[OpenACC][CIR] link clause lowering for global declare #169524
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
erichkeane
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clangir @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThe 'link' clause is like the rest of the global clauses (copyin, This patch also removes a bunch of now stales TODOs from the tests. Full diff: https://github.com/llvm/llvm-project/pull/169524.diff 8 Files Affected:
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"
|
zwuis
added
the
skip-precommit-approval
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.