[clang] 2403865 - [OpenACC][CIR] implement 'collapse' lowering for combined constructs
via cfe-commits
cfe-commits at lists.llvm.org
Fri May 9 10:52:33 PDT 2025
Author: erichkeane
Date: 2025-05-09T10:52:23-07:00
New Revision: 24038650d9ca5d66b07d3075afdebe81012ab1f2
URL: https://github.com/llvm/llvm-project/commit/24038650d9ca5d66b07d3075afdebe81012ab1f2
DIFF: https://github.com/llvm/llvm-project/commit/24038650d9ca5d66b07d3075afdebe81012ab1f2.diff
LOG: [OpenACC][CIR] implement 'collapse' lowering for combined constructs
Another trivial implementation. It has a constant value that doesn't
require any insertion of instructions, so this just works with minimal
effort.
Added:
Modified:
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
clang/test/CIR/CodeGenOpenACC/combined.cpp
Removed:
################################################################################
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
index 8652a0fee0994..8892c49e41202 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
@@ -414,10 +414,10 @@ class OpenACCClauseCIREmitter final
value = value.sextOrTrunc(64);
operation.setCollapseForDeviceTypes(builder.getContext(),
lastDeviceTypeValues, value);
+ } else if constexpr (isCombinedType<OpTy>) {
+ applyToLoopOp(clause);
} else {
- // TODO: When we've implemented this for everything, switch this to an
- // unreachable. Combined constructs remain.
- return clauseNotImplemented(clause);
+ llvm_unreachable("Unknown construct kind in VisitCollapseClause");
}
}
diff --git a/clang/test/CIR/CodeGenOpenACC/combined.cpp b/clang/test/CIR/CodeGenOpenACC/combined.cpp
index 38ca45daf048f..3b2ae8a97d8c5 100644
--- a/clang/test/CIR/CodeGenOpenACC/combined.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/combined.cpp
@@ -134,4 +134,46 @@ extern "C" void acc_combined(int N) {
// CHECK: acc.terminator
// CHECK-NEXT: } loc
+ #pragma acc parallel loop collapse(1) device_type(radeon)
+ for(unsigned I = 0; I < N; ++I)
+ for(unsigned J = 0; J < N; ++J)
+ for(unsigned K = 0; K < N; ++K);
+ // CHECK: acc.parallel combined(loop) {
+ // CHECK: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {collapse = [1], collapseDeviceType = [#acc.device_type<none>]}
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+
+ #pragma acc serial loop collapse(1) device_type(radeon) collapse (2)
+ for(unsigned I = 0; I < N; ++I)
+ for(unsigned J = 0; J < N; ++J)
+ for(unsigned K = 0; K < N; ++K);
+ // CHECK: acc.serial combined(loop) {
+ // CHECK: acc.loop combined(serial) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {collapse = [1, 2], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>]}
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+
+ #pragma acc kernels loop collapse(1) device_type(radeon, nvidia) collapse (2)
+ for(unsigned I = 0; I < N; ++I)
+ for(unsigned J = 0; J < N; ++J)
+ for(unsigned K = 0; K < N; ++K);
+ // CHECK: acc.kernels combined(loop) {
+ // CHECK: acc.loop combined(kernels) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {collapse = [1, 2, 2], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>, #acc.device_type<nvidia>]}
+ // CHECK: acc.terminator
+ // CHECK-NEXT: } loc
+ #pragma acc parallel loop collapse(1) device_type(radeon, nvidia) collapse(2) device_type(host) collapse(3)
+ for(unsigned I = 0; I < N; ++I)
+ for(unsigned J = 0; J < N; ++J)
+ for(unsigned K = 0; K < N; ++K);
+ // CHECK: acc.parallel combined(loop) {
+ // CHECK: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {collapse = [1, 2, 2, 3], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>, #acc.device_type<nvidia>, #acc.device_type<host>]}
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
}
More information about the cfe-commits
mailing list