Author: Erich Keane Date: 2025-11-25T17:23:50Z New Revision: 53e5cfdf8b13e2427797ca6eeda1860f8aa190ef
URL: https://github.com/llvm/llvm-project/commit/53e5cfdf8b13e2427797ca6eeda1860f8aa190ef DIFF: https://github.com/llvm/llvm-project/commit/53e5cfdf8b13e2427797ca6eeda1860f8aa190ef.diff LOG: [OpenACC][CIR] link clause lowering for global declare (#169524) The 'link' clause is like the rest of the global clauses (copyin, create, device_resident), except it only has an entry op(thus no dtor). This patch also removes a bunch of now stales TODOs from the tests. Added: Modified: clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp clang/test/CIR/CodeGenOpenACC/combined-copy.c clang/test/CIR/CodeGenOpenACC/compute-copy.c clang/test/CIR/CodeGenOpenACC/declare-copy.cpp clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp clang/test/CIR/CodeGenOpenACC/declare-link.cpp clang/test/CIR/CodeGenOpenACC/declare-present.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index c1a1f8a83f5cd..405c1aad2f159 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -231,16 +231,12 @@ namespace { class OpenACCGlobalDeclareClauseEmitter final : public OpenACCClauseVisitor<OpenACCGlobalDeclareClauseEmitter> { CIRGenModule &cgm; - void clauseNotImplemented(const OpenACCClause &c) { - cgm.errorNYI(c.getSourceRange(), "OpenACC Global Declare Clause", - c.getClauseKind()); - } public: OpenACCGlobalDeclareClauseEmitter(CIRGenModule &cgm) : cgm(cgm) {} void VisitClause(const OpenACCClause &clause) { - clauseNotImplemented(clause); + llvm_unreachable("Invalid OpenACC clause on global Declare"); } void emitClauses(ArrayRef<const OpenACCClause *> clauses) { @@ -271,6 +267,14 @@ class OpenACCGlobalDeclareClauseEmitter final /*structured=*/true, /*implicit=*/false, /*requiresDtor=*/true); } + + void VisitLinkClause(const OpenACCLinkClause &clause) { + for (const Expr *var : clause.getVarList()) + cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::DeclareLinkOp>( + var, mlir::acc::DataClause::acc_declare_link, {}, + /*structured=*/true, + /*implicit=*/false, /*requiresDtor=*/false); + } }; } // namespace diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c b/clang/test/CIR/CodeGenOpenACC/combined-copy.c index 31956b383df02..e1b4e593a86fd 100644 --- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c @@ -73,8 +73,6 @@ void acc_compute(int parmVar) { // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc - // TODO: OpenACC: Represent alwaysin/alwaysout/always correctly. For now, - // these do nothing to the IR. #pragma acc parallel loop copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3) for(int i = 0; i < 5; ++i); // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copy.c b/clang/test/CIR/CodeGenOpenACC/compute-copy.c index 41e594ec3551b..fd8b5ee3761c4 100644 --- a/clang/test/CIR/CodeGenOpenACC/compute-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/compute-copy.c @@ -65,8 +65,6 @@ void acc_compute(int parmVar) { // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc - // TODO: OpenACC: Represent alwaysin/alwaysout/always correctly. For now, - // these do nothing to the IR. #pragma acc parallel copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3) ; // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp index a8a9115a21b29..1dd66826da96b 100644 --- a/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp +++ b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp @@ -5,15 +5,11 @@ struct 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(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) { // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}) // CHECK-NEXT: cir.alloca{{.*}}["this" diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp index 1d79cef894d5e..33e76a3b93e9c 100644 --- a/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp +++ b/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp @@ -5,15 +5,11 @@ struct 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(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) { // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}) // CHECK-NEXT: cir.alloca{{.*}}["this" diff --git a/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp b/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp index d8021ef9a9dc5..f6591f78aa225 100644 --- a/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp +++ b/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp @@ -5,15 +5,11 @@ struct 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(HasSideEffects *ArgHSE, int *ArgInt) { // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}}) // CHECK-NEXT: cir.alloca{{.*}}["this" diff --git a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp index 8494a2354c7db..5fc78167ce991 100644 --- a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp +++ b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp @@ -5,14 +5,116 @@ struct HasSideEffects { ~HasSideEffects(); }; -// TODO: OpenACC: Implement 'global', NS lowering. +HasSideEffects GlobalHSE1; +HasSideEffects GlobalHSEArr[5]; +int GlobalInt1; -struct Struct { - static const HasSideEffects StaticMemHSE; - static const HasSideEffects StaticMemHSEArr[5]; - static const int StaticMemInt; +#pragma acc declare link(GlobalHSE1, GlobalInt1, GlobalHSEArr[1:1]) +// CHECK: acc.global_ctor @GlobalHSE1_acc_ctor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr<!rec_HasSideEffects> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "GlobalHSE1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// CHECK-NOT: acc.global_dtor +// +// CHECK: acc.global_ctor @GlobalInt1_acc_ctor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr<!s32i> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "GlobalInt1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// +// CHECK: acc.global_ctor @GlobalHSEArr_acc_ctor { +// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i +// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] +// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i +// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] +// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 +// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 +// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "GlobalHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } + +namespace NS { + +HasSideEffects NSHSE1; +HasSideEffects NSHSEArr[5]; +int NSInt1; + +#pragma acc declare link(NSHSE1, NSInt1, NSHSEArr[1:1]) +// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "NSHSE1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// +// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "NSInt1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// +// CHECK: acc.global_ctor @{{.*}}NSHSEArr{{.*}}_acc_ctor { +// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i +// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] +// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i +// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] +// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 +// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 +// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "NSHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } + +} // namespace NS + +namespace { - // TODO: OpenACC: Implement static-local lowering. +HasSideEffects AnonNSHSE1; +HasSideEffects AnonNSHSEArr[5]; +int AnonNSInt1; + +#pragma acc declare link(AnonNSHSE1, AnonNSInt1, AnonNSHSEArr[1:1]) +// CHECK: acc.global_ctor @{{.*}}AnonNSHSE1{{.*}}_acc_ctor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "AnonNSHSE1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// +// CHECK: acc.global_ctor @{{.*}}AnonNSInt1{{.*}}_acc_ctor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr<!s32i> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "AnonNSInt1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// +// CHECK: acc.global_ctor @{{.*}}AnonNSHSEArr{{.*}}_acc_ctor { +// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i +// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]] +// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i +// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]] +// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 +// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 +// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> +// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "AnonNSHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } + +} // namespace NS + + +struct Struct { void MemFunc1() { // CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) { diff --git a/clang/test/CIR/CodeGenOpenACC/declare-present.cpp b/clang/test/CIR/CodeGenOpenACC/declare-present.cpp index c17b9597adf12..9c646d62a4f3c 100644 --- a/clang/test/CIR/CodeGenOpenACC/declare-present.cpp +++ b/clang/test/CIR/CodeGenOpenACC/declare-present.cpp @@ -5,15 +5,11 @@ struct 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(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) { // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}) // CHECK-NEXT: cir.alloca{{.*}}["this" _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
