https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/168793
>From 22caecc11fc0d5d6113bfd6ba24f3644316e5350 Mon Sep 17 00:00:00 2001 From: erichkeane <[email protected]> Date: Wed, 19 Nov 2025 13:28:18 -0800 Subject: [PATCH 1/2] [OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) 'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'. A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical. 'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses. --- clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 32 ++++- clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 13 ++ .../test/CIR/CodeGenOpenACC/declare-link.cpp | 130 ++++++++++++++++++ .../openacc-not-implemented.cpp | 5 +- 4 files changed, 177 insertions(+), 3 deletions(-) create mode 100644 clang/test/CIR/CodeGenOpenACC/declare-link.cpp diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index b588a50aa0404..f6680cbaa8c78 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -12,12 +12,42 @@ #include "CIRGenFunction.h" #include "clang/AST/DeclOpenACC.h" +#include "mlir/Dialect/OpenACC/OpenACC.h" using namespace clang; using namespace clang::CIRGen; +namespace { + struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { + mlir::acc::DeclareEnterOp enterOp; + + OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) + : enterOp(enterOp) {} + + void emit(CIRGenFunction &cgf) override { + mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(), + enterOp, {}); + + // TODO(OpenACC): Some clauses require that we add info about them to the + // DeclareExitOp. However, we don't have any of those implemented yet, so + // we should add infrastructure here to do that once we have one + // implemented. + } + + }; +} // namespace + void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) { - getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Declare Construct"); + mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc()); + auto enterOp = mlir::acc::DeclareEnterOp::create( + builder, exprLoc, + mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()), {}); + + emitOpenACCClauses(enterOp, OpenACCDirectiveKind::Declare, d.getBeginLoc(), + d.clauses()); + + ehStack.pushCleanup<OpenACCDeclareCleanup>(CleanupKind::NormalCleanup, + enterOp); } void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) { diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index 527dfd21db8a5..c7e6a256c3868 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -876,6 +876,18 @@ class OpenACCClauseCIREmitter final } } + void VisitLinkClause(const OpenACCLinkClause &clause) { + if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) { + for (const Expr *var : clause.getVarList()) + addDataOperand<mlir::acc::DeclareLinkOp>( + var, mlir::acc::DataClause::acc_declare_link, {}, + /*structured=*/true, + /*implicit=*/false); + } else { + llvm_unreachable("Unknown construct kind in VisitLinkClause"); + } + } + void VisitDeleteClause(const OpenACCDeleteClause &clause) { if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) { for (const Expr *var : clause.getVarList()) @@ -1151,6 +1163,7 @@ EXPL_SPEC(mlir::acc::AtomicReadOp) EXPL_SPEC(mlir::acc::AtomicWriteOp) EXPL_SPEC(mlir::acc::AtomicCaptureOp) EXPL_SPEC(mlir::acc::AtomicUpdateOp) +EXPL_SPEC(mlir::acc::DeclareEnterOp) #undef EXPL_SPEC template <typename ComputeOp, typename LoopOp> diff --git a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp new file mode 100644 index 0000000000000..8494a2354c7db --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp @@ -0,0 +1,130 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +struct HasSideEffects { + HasSideEffects(); + ~HasSideEffects(); +}; + +// TODO: OpenACC: Implement 'global', NS lowering. + +struct Struct { + static const HasSideEffects StaticMemHSE; + static const HasSideEffects StaticMemHSEArr[5]; + static const int StaticMemInt; + + // TODO: OpenACC: Implement static-local lowering. + + void MemFunc1() { + // CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) { + // CHECK-NEXT: cir.alloca{{.*}}["this" + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.load + extern HasSideEffects LocalHSE; + extern HasSideEffects LocalHSEArr[5]; + extern int LocalInt; +#pragma acc declare link(LocalHSE, LocalInt, LocalHSEArr[1:1]) + + // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE : !cir.ptr<!rec_HasSideEffects> + // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"} + // + // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt : !cir.ptr<!s32i> + // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"} + // + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64) + // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> + // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"} + // + // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) + // + // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) + } + + void MemFunc2(); +}; +void use() { + Struct s; + s.MemFunc1(); +} + +void Struct::MemFunc2() { + // CHECK: cir.func {{.*}}MemFunc2{{.*}}({{.*}}) { + // CHECK-NEXT: cir.alloca{{.*}}["this" + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.load + extern HasSideEffects LocalHSE2; + extern HasSideEffects LocalHSEArr2[5]; + extern int LocalInt2; + +#pragma acc declare link(LocalHSE2, LocalInt2, LocalHSEArr2[1:1]) + // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE2 : !cir.ptr<!rec_HasSideEffects> + // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE2"} + // + // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt2 : !cir.ptr<!s32i> + // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt2"} + // + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64) + // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr2 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> + // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr2[1:1]"} + // + // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) + // + // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) +} + +extern "C" void do_thing(); + +void NormalFunc() { + // CHECK: cir.func {{.*}}NormalFunc{{.*}}() + extern HasSideEffects LocalHSE3; + extern HasSideEffects LocalHSEArr3[5]; + extern int LocalInt3; + // CHECK-NEXT: cir.scope + { + extern HasSideEffects InnerHSE; +#pragma acc declare link(LocalHSE3, LocalInt3, LocalHSEArr3[1:1], InnerHSE) + // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE3 : !cir.ptr<!rec_HasSideEffects> + // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE3"} + // + // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt3 : !cir.ptr<!s32i> + // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt3"} + // + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64) + // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr3 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> + // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr3[1:1]"} + // + // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @InnerHSE : !cir.ptr<!rec_HasSideEffects> + // CHECK-NEXT: %[[INNERHSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "InnerHSE"} + // + // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]], %[[INNERHSE_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>, !cir.ptr<!rec_HasSideEffects>) + // + // CHECK + + do_thing(); + // CHECK-NEXT: cir.call @do_thing + + // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) + } + // CHECK-NEXT: } + + do_thing(); + // CHECK-NEXT: cir.call @do_thing +} + diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp index e85c26718acb8..c8b85a12f84e7 100644 --- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp +++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp @@ -1,7 +1,8 @@ // RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify void HelloWorld(int *A) { + extern int *E; -// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}} -#pragma acc declare create(A) +// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: create}} +#pragma acc declare link(E) create(A) } >From 126313abdaaf4e3756e494d4b4ca313ee7a6abaa Mon Sep 17 00:00:00 2001 From: erichkeane <[email protected]> Date: Wed, 19 Nov 2025 15:46:59 -0800 Subject: [PATCH 2/2] clang-format --- clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 32 ++++++++++----------- 1 file changed, 15 insertions(+), 17 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index f6680cbaa8c78..551027bb1c8eb 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -11,37 +11,35 @@ //===----------------------------------------------------------------------===// #include "CIRGenFunction.h" -#include "clang/AST/DeclOpenACC.h" #include "mlir/Dialect/OpenACC/OpenACC.h" +#include "clang/AST/DeclOpenACC.h" using namespace clang; using namespace clang::CIRGen; namespace { - struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { - mlir::acc::DeclareEnterOp enterOp; - - OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) - : enterOp(enterOp) {} +struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { + mlir::acc::DeclareEnterOp enterOp; - void emit(CIRGenFunction &cgf) override { - mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(), - enterOp, {}); + OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) : enterOp(enterOp) {} - // TODO(OpenACC): Some clauses require that we add info about them to the - // DeclareExitOp. However, we don't have any of those implemented yet, so - // we should add infrastructure here to do that once we have one - // implemented. - } + void emit(CIRGenFunction &cgf) override { + mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(), + enterOp, {}); - }; + // TODO(OpenACC): Some clauses require that we add info about them to the + // DeclareExitOp. However, we don't have any of those implemented yet, so + // we should add infrastructure here to do that once we have one + // implemented. + } +}; } // namespace void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) { mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc()); auto enterOp = mlir::acc::DeclareEnterOp::create( - builder, exprLoc, - mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()), {}); + builder, exprLoc, mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()), + {}); emitOpenACCClauses(enterOp, OpenACCDirectiveKind::Declare, d.getBeginLoc(), d.clauses()); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
