[clang] 4efcc52 - [OpenACC][CIR] Implement Loop lowering of seq/auto/independent (#138164)
via cfe-commits
cfe-commits at lists.llvm.org
Thu May 1 14:30:14 PDT 2025
Author: Erich Keane
Date: 2025-05-01T14:30:11-07:00
New Revision: 4efcc52ed839c4348c69a01538c7ecd399e4b113
URL: https://github.com/llvm/llvm-project/commit/4efcc52ed839c4348c69a01538c7ecd399e4b113
DIFF: https://github.com/llvm/llvm-project/commit/4efcc52ed839c4348c69a01538c7ecd399e4b113.diff
LOG: [OpenACC][CIR] Implement Loop lowering of seq/auto/independent (#138164)
These just add a standard 'device_type' flag to the acc.loop, so
implement that lowering. This also modifies the dialect to add helpers
for these as well, to be consistent with the previous ones.
Added:
Modified:
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
clang/test/CIR/CodeGenOpenACC/loop.cpp
mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
Removed:
################################################################################
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
index b54682402d961..ff0bf6e7f55dd 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
@@ -147,13 +147,13 @@ class OpenACCClauseCIREmitter final
decodeDeviceType(clause.getArchitectures()[0].getIdentifierInfo()));
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp, mlir::acc::KernelsOp,
- mlir::acc::DataOp>) {
+ mlir::acc::DataOp, mlir::acc::LoopOp>) {
// Nothing to do here, these constructs don't have any IR for these, as
// they just modify the other clauses IR. So setting of
// `lastDeviceTypeValues` (done above) is all we need.
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. update, data, loop, routine, combined constructs remain.
+ // unreachable. update, data, routine, combined constructs remain.
return clauseNotImplemented(clause);
}
}
@@ -306,6 +306,36 @@ class OpenACCClauseCIREmitter final
llvm_unreachable("set, is only valid device_num constructs");
}
}
+
+ void VisitSeqClause(const OpenACCSeqClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+ operation.addSeq(builder.getContext(), lastDeviceTypeValues);
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. Routine, Combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
+
+ void VisitAutoClause(const OpenACCAutoClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+ operation.addAuto(builder.getContext(), lastDeviceTypeValues);
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. Routine, Combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
+
+ void VisitIndependentClause(const OpenACCIndependentClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+ operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. Routine, Combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
};
template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/loop.cpp b/clang/test/CIR/CodeGenOpenACC/loop.cpp
index 792edfedaacc6..2757d935e1f76 100644
--- a/clang/test/CIR/CodeGenOpenACC/loop.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/loop.cpp
@@ -30,4 +30,83 @@ extern "C" void acc_loop(int *A, int *B, int *C, int N) {
// CHECK-NEXT: } loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
+
+
+#pragma acc loop seq
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+#pragma acc loop device_type(nvidia, radeon) seq
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {seq = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc
+#pragma acc loop device_type(radeon) seq
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {seq = [#acc.device_type<radeon>]} loc
+#pragma acc loop seq device_type(nvidia, radeon)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+#pragma acc loop seq device_type(radeon)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+
+#pragma acc loop independent
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc
+#pragma acc loop device_type(nvidia, radeon) independent
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {independent = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc
+#pragma acc loop device_type(radeon) independent
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {independent = [#acc.device_type<radeon>]} loc
+#pragma acc loop independent device_type(nvidia, radeon)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc
+#pragma acc loop independent device_type(radeon)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc
+
+#pragma acc loop auto
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc
+#pragma acc loop device_type(nvidia, radeon) auto
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc
+#pragma acc loop device_type(radeon) auto
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<radeon>]} loc
+#pragma acc loop auto device_type(nvidia, radeon)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc
+#pragma acc loop auto device_type(radeon)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK: acc.loop {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc
}
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 3ad8e4f9ccbeb..c3df064cf0ead 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2198,6 +2198,14 @@ def OpenACC_LoopOp : OpenACC_Op<"loop",
/// Return the value of the worker clause for the given device_type
/// if present.
mlir::Value getGangValue(mlir::acc::GangArgType gangArgType, mlir::acc::DeviceType deviceType);
+
+ // Add an entry to the 'seq' attribute for each additional device types.
+ void addSeq(MLIRContext *, llvm::ArrayRef<DeviceType>);
+ // Add an entry to the 'independent' attribute for each additional device
+ // types.
+ void addIndependent(MLIRContext *, llvm::ArrayRef<DeviceType>);
+ // Add an entry to the 'auto' attribute for each additional device types.
+ void addAuto(MLIRContext *, llvm::ArrayRef<DeviceType>);
}];
let hasCustomAssemblyFormat = 1;
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index d23563f1f0fb0..39dbb0c92a309 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -2651,6 +2651,24 @@ void printLoopControl(OpAsmPrinter &p, Operation *op, Region ®ion,
p.printRegion(region, /*printEntryBlockArgs=*/false);
}
+void acc::LoopOp::addSeq(MLIRContext *context,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ setSeqAttr(addDeviceTypeAffectedOperandHelper(context, getSeqAttr(),
+ effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addIndependent(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ setIndependentAttr(addDeviceTypeAffectedOperandHelper(
+ context, getIndependentAttr(), effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addAuto(MLIRContext *context,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ setAuto_Attr(addDeviceTypeAffectedOperandHelper(context, getAuto_Attr(),
+ effectiveDeviceTypes));
+}
+
//===----------------------------------------------------------------------===//
// DataOp
//===----------------------------------------------------------------------===//
More information about the cfe-commits
mailing list