https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/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.
>From 47b0f5d4387aec83ef2f0d9affe822e12b4213ee Mon Sep 17 00:00:00 2001
From: erichkeane <[email protected]>
Date: Mon, 24 Nov 2025 14:22:15 -0800
Subject: [PATCH] [OpenACC][CIR] link clause lowering for global declare
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.
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 14 ++-
clang/test/CIR/CodeGenOpenACC/combined-copy.c | 2 -
clang/test/CIR/CodeGenOpenACC/compute-copy.c | 2 -
.../test/CIR/CodeGenOpenACC/declare-copy.cpp | 4 -
.../CIR/CodeGenOpenACC/declare-copyout.cpp | 4 -
.../CIR/CodeGenOpenACC/declare-deviceptr.cpp | 4 -
.../test/CIR/CodeGenOpenACC/declare-link.cpp | 114 +++++++++++++++++-
.../CIR/CodeGenOpenACC/declare-present.cpp | 4 -
8 files changed, 117 insertions(+), 31 deletions(-)
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
