Skip to content

Commit 53e5cfd

Browse files
erichkeaneerichkeane
authored
[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.
1 parent a860c83 commit 53e5cfd

File tree

8 files changed

+117
-31
lines changed

8 files changed

+117
-31
lines changed

‎clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp‎

Lines changed: 9 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -231,16 +231,12 @@ namespace {
231231
class OpenACCGlobalDeclareClauseEmitter final
232232
: public OpenACCClauseVisitor<OpenACCGlobalDeclareClauseEmitter> {
233233
CIRGenModule &cgm;
234-
void clauseNotImplemented(const OpenACCClause &c) {
235-
cgm.errorNYI(c.getSourceRange(), "OpenACC Global Declare Clause",
236-
c.getClauseKind());
237-
}
238234

239235
public:
240236
OpenACCGlobalDeclareClauseEmitter(CIRGenModule &cgm) : cgm(cgm) {}
241237

242238
void VisitClause(const OpenACCClause &clause) {
243-
clauseNotImplemented(clause);
239+
llvm_unreachable("Invalid OpenACC clause on global Declare");
244240
}
245241

246242
void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
@@ -271,6 +267,14 @@ class OpenACCGlobalDeclareClauseEmitter final
271267
/*structured=*/true,
272268
/*implicit=*/false, /*requiresDtor=*/true);
273269
}
270+
271+
void VisitLinkClause(const OpenACCLinkClause &clause) {
272+
for (const Expr *var : clause.getVarList())
273+
cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::DeclareLinkOp>(
274+
var, mlir::acc::DataClause::acc_declare_link, {},
275+
/*structured=*/true,
276+
/*implicit=*/false, /*requiresDtor=*/false);
277+
}
274278
};
275279
} // namespace
276280

‎clang/test/CIR/CodeGenOpenACC/combined-copy.c‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -73,8 +73,6 @@ void acc_compute(int parmVar) {
7373
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc
7474
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc
7575

76-
// TODO: OpenACC: Represent alwaysin/alwaysout/always correctly. For now,
77-
// these do nothing to the IR.
7876
#pragma acc parallel loop copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3)
7977
for(int i = 0; i < 5; ++i);
8078
// 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

‎clang/test/CIR/CodeGenOpenACC/compute-copy.c‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -65,8 +65,6 @@ void acc_compute(int parmVar) {
6565
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc
6666
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc
6767

68-
// TODO: OpenACC: Represent alwaysin/alwaysout/always correctly. For now,
69-
// these do nothing to the IR.
7068
#pragma acc parallel copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3)
7169
;
7270
// 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

‎clang/test/CIR/CodeGenOpenACC/declare-copy.cpp‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,15 +5,11 @@ struct HasSideEffects {
55
~HasSideEffects();
66
};
77

8-
// TODO: OpenACC: Implement 'global', NS lowering.
9-
108
struct Struct {
119
static const HasSideEffects StaticMemHSE;
1210
static const HasSideEffects StaticMemHSEArr[5];
1311
static const int StaticMemInt;
1412

15-
// TODO: OpenACC: Implement static-local lowering.
16-
1713
void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
1814
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})
1915
// CHECK-NEXT: cir.alloca{{.*}}["this"

‎clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,15 +5,11 @@ struct HasSideEffects {
55
~HasSideEffects();
66
};
77

8-
// TODO: OpenACC: Implement 'global', NS lowering.
9-
108
struct Struct {
119
static const HasSideEffects StaticMemHSE;
1210
static const HasSideEffects StaticMemHSEArr[5];
1311
static const int StaticMemInt;
1412

15-
// TODO: OpenACC: Implement static-local lowering.
16-
1713
void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
1814
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})
1915
// CHECK-NEXT: cir.alloca{{.*}}["this"

‎clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,15 +5,11 @@ struct HasSideEffects {
55
~HasSideEffects();
66
};
77

8-
// TODO: OpenACC: Implement 'global', NS lowering.
9-
108
struct Struct {
119
static const HasSideEffects StaticMemHSE;
1210
static const HasSideEffects StaticMemHSEArr[5];
1311
static const int StaticMemInt;
1412

15-
// TODO: OpenACC: Implement static-local lowering.
16-
1713
void MemFunc1(HasSideEffects *ArgHSE, int *ArgInt) {
1814
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
1915
// CHECK-NEXT: cir.alloca{{.*}}["this"

‎clang/test/CIR/CodeGenOpenACC/declare-link.cpp‎

Lines changed: 108 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -5,14 +5,116 @@ struct HasSideEffects {
55
~HasSideEffects();
66
};
77

8-
// TODO: OpenACC: Implement 'global', NS lowering.
8+
HasSideEffects GlobalHSE1;
9+
HasSideEffects GlobalHSEArr[5];
10+
int GlobalInt1;
911

10-
struct Struct {
11-
static const HasSideEffects StaticMemHSE;
12-
static const HasSideEffects StaticMemHSEArr[5];
13-
static const int StaticMemInt;
12+
#pragma acc declare link(GlobalHSE1, GlobalInt1, GlobalHSEArr[1:1])
13+
// CHECK: acc.global_ctor @GlobalHSE1_acc_ctor {
14+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr<!rec_HasSideEffects>
15+
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "GlobalHSE1"}
16+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
17+
// CHECK-NEXT: acc.terminator
18+
// CHECK-NEXT: }
19+
// CHECK-NOT: acc.global_dtor
20+
//
21+
// CHECK: acc.global_ctor @GlobalInt1_acc_ctor {
22+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr<!s32i>
23+
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "GlobalInt1"}
24+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
25+
// CHECK-NEXT: acc.terminator
26+
// CHECK-NEXT: }
27+
//
28+
// CHECK: acc.global_ctor @GlobalHSEArr_acc_ctor {
29+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
30+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
31+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
32+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
33+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
34+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
35+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
36+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
37+
// 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]"}
38+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
39+
// CHECK-NEXT: acc.terminator
40+
// CHECK-NEXT: }
41+
42+
namespace NS {
43+
44+
HasSideEffects NSHSE1;
45+
HasSideEffects NSHSEArr[5];
46+
int NSInt1;
47+
48+
#pragma acc declare link(NSHSE1, NSInt1, NSHSEArr[1:1])
49+
// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor {
50+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
51+
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "NSHSE1"}
52+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
53+
// CHECK-NEXT: acc.terminator
54+
// CHECK-NEXT: }
55+
//
56+
// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor {
57+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
58+
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "NSInt1"}
59+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
60+
// CHECK-NEXT: acc.terminator
61+
// CHECK-NEXT: }
62+
//
63+
// CHECK: acc.global_ctor @{{.*}}NSHSEArr{{.*}}_acc_ctor {
64+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
65+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
66+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
67+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
68+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
69+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
70+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
71+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
72+
// 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]"}
73+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
74+
// CHECK-NEXT: acc.terminator
75+
// CHECK-NEXT: }
76+
77+
} // namespace NS
78+
79+
namespace {
1480

15-
// TODO: OpenACC: Implement static-local lowering.
81+
HasSideEffects AnonNSHSE1;
82+
HasSideEffects AnonNSHSEArr[5];
83+
int AnonNSInt1;
84+
85+
#pragma acc declare link(AnonNSHSE1, AnonNSInt1, AnonNSHSEArr[1:1])
86+
// CHECK: acc.global_ctor @{{.*}}AnonNSHSE1{{.*}}_acc_ctor {
87+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
88+
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "AnonNSHSE1"}
89+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
90+
// CHECK-NEXT: acc.terminator
91+
// CHECK-NEXT: }
92+
//
93+
// CHECK: acc.global_ctor @{{.*}}AnonNSInt1{{.*}}_acc_ctor {
94+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr<!s32i>
95+
// CHECK-NEXT: %[[CREATE:.*]] = acc.declare_link varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "AnonNSInt1"}
96+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
97+
// CHECK-NEXT: acc.terminator
98+
// CHECK-NEXT: }
99+
//
100+
// CHECK: acc.global_ctor @{{.*}}AnonNSHSEArr{{.*}}_acc_ctor {
101+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
102+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
103+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
104+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
105+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
106+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
107+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
108+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
109+
// 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]"}
110+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
111+
// CHECK-NEXT: acc.terminator
112+
// CHECK-NEXT: }
113+
114+
} // namespace NS
115+
116+
117+
struct Struct {
16118

17119
void MemFunc1() {
18120
// CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) {

‎clang/test/CIR/CodeGenOpenACC/declare-present.cpp‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,15 +5,11 @@ struct HasSideEffects {
55
~HasSideEffects();
66
};
77

8-
// TODO: OpenACC: Implement 'global', NS lowering.
9-
108
struct Struct {
119
static const HasSideEffects StaticMemHSE;
1210
static const HasSideEffects StaticMemHSEArr[5];
1311
static const int StaticMemInt;
1412

15-
// TODO: OpenACC: Implement static-local lowering.
16-
1713
void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
1814
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})
1915
// CHECK-NEXT: cir.alloca{{.*}}["this"

0 commit comments

Comments
 (0)