https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/170358
These two are both incredibly similar and simple, basically identical to 'seq'. This patch adds them both together. >From ed8e2ed741907566f2c1828c44dc5fdf8ab1f8d0 Mon Sep 17 00:00:00 2001 From: erichkeane <[email protected]> Date: Tue, 2 Dec 2025 08:37:25 -0800 Subject: [PATCH] [OpenACC][CIR] Add worker/vector clause lowering for Routine These two are both incredibly similar and simple, basically identical to 'seq'. This patch adds them both together. --- clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 6 +++ .../CIR/CodeGenOpenACC/routine-clauses.cpp | 38 +++++++++++++++++++ .../mlir/Dialect/OpenACC/OpenACCOps.td | 4 ++ mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 12 ++++++ 4 files changed, 60 insertions(+) create mode 100644 clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index 0b3a877202fb1..19def59db14ba 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -323,6 +323,12 @@ class OpenACCRoutineClauseEmitter final void VisitSeqClause(const OpenACCSeqClause &clause) { routineOp.addSeq(builder.getContext(), lastDeviceTypeValues); } + void VisitWorkerClause(const OpenACCWorkerClause &clause) { + routineOp.addWorker(builder.getContext(), lastDeviceTypeValues); + } + void VisitVectorClause(const OpenACCVectorClause &clause) { + routineOp.addVector(builder.getContext(), lastDeviceTypeValues); + } }; } // namespace diff --git a/clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp b/clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp new file mode 100644 index 0000000000000..7ad64f2d05158 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp @@ -0,0 +1,38 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +#pragma acc routine seq +void Func1() {} + +void Func2() {} +#pragma acc routine(Func2) seq + +#pragma acc routine worker +void Func3() {} + +void Func4() {} +#pragma acc routine(Func4) worker + +#pragma acc routine vector +void Func5() {} + +void Func6() {} +#pragma acc routine(Func6) vector + +// CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq + +// CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>} + +// CHECK: cir.func{{.*}} @[[F3_NAME:.*Func3[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F3_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F3_R_NAME]] func(@[[F3_NAME]]) worker + +// CHECK: cir.func{{.*}} @[[F4_NAME:.*Func4[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F4_R_NAME:.*]]]>} + +// CHECK: cir.func{{.*}} @[[F5_NAME:.*Func5[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F5_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F5_R_NAME]] func(@[[F5_NAME]]) vector + +// CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>} + +// CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) seq +// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) worker +// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) vector diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td index be05b9d6fbddc..77d1a6f8d53b5 100644 --- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td +++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td @@ -3282,6 +3282,10 @@ def OpenACC_RoutineOp : OpenACC_Op<"routine", [IsolatedFromAbove]> { // Add an entry to the 'seq' attribute for each additional device types. void addSeq(MLIRContext *, llvm::ArrayRef<DeviceType>); + // Add an entry to the 'vector' attribute for each additional device types. + void addVector(MLIRContext *, llvm::ArrayRef<DeviceType>); + // Add an entry to the 'worker' attribute for each additional device types. + void addWorker(MLIRContext *, llvm::ArrayRef<DeviceType>); }]; let assemblyFormat = [{ diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp index 565af9b38cdf4..828f329e48d4a 100644 --- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp +++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp @@ -4299,6 +4299,18 @@ void RoutineOp::addSeq(MLIRContext *context, effectiveDeviceTypes)); } +void RoutineOp::addVector(MLIRContext *context, + llvm::ArrayRef<DeviceType> effectiveDeviceTypes) { + setVectorAttr(addDeviceTypeAffectedOperandHelper(context, getVectorAttr(), + effectiveDeviceTypes)); +} + +void RoutineOp::addWorker(MLIRContext *context, + llvm::ArrayRef<DeviceType> effectiveDeviceTypes) { + setWorkerAttr(addDeviceTypeAffectedOperandHelper(context, getWorkerAttr(), + effectiveDeviceTypes)); +} + //===----------------------------------------------------------------------===// // InitOp //===----------------------------------------------------------------------===// _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
