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 //===----------------------------------------------------------------------===// _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
