Skip to content

Commit ea5449d

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] Implement 'async'/'if' lowering for 'data' construct
These two are trivial, and work the same as the compute construct versions of these, so this adds tests to do so, and adds them to the implementation.
1 parent a83b4a2 commit ea5449d

File tree

3 files changed

+119
-5
lines changed

3 files changed

+119
-5
lines changed

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

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -204,7 +204,8 @@ class OpenACCClauseCIREmitter final
204204
if (!clause.getArchitectures().empty())
205205
operation.setDeviceType(
206206
decodeDeviceType(clause.getArchitectures()[0].getIdentifierInfo()));
207-
} else if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
207+
} else if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp,
208+
DataOp>) {
208209
// Nothing to do here, these constructs don't have any IR for these, as
209210
// they just modify the other clauses IR. So setting of `lastDeviceType`
210211
// (done above) is all we need.
@@ -243,7 +244,7 @@ class OpenACCClauseCIREmitter final
243244
}
244245

245246
void VisitAsyncClause(const OpenACCAsyncClause &clause) {
246-
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
247+
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, DataOp>) {
247248
if (!clause.hasIntExpr()) {
248249
operation.setAsyncOnlyAttr(
249250
handleDeviceTypeAffectedClause(operation.getAsyncOnlyAttr()));
@@ -278,7 +279,7 @@ class OpenACCClauseCIREmitter final
278279

279280
void VisitIfClause(const OpenACCIfClause &clause) {
280281
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, InitOp,
281-
ShutdownOp, SetOp>) {
282+
ShutdownOp, SetOp, DataOp>) {
282283
operation.getIfCondMutable().append(
283284
createCondition(clause.getConditionExpr()));
284285
} else {

‎clang/test/CIR/CodeGenOpenACC/data.c

Lines changed: 78 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,9 @@
11
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
22

3-
void acc_data(void) {
4-
// CHECK: cir.func @acc_data() {
3+
void acc_data(int cond) {
4+
// CHECK: cir.func @acc_data(%[[ARG:.*]]: !s32i{{.*}}) {
5+
// CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
6+
// CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
57

68
#pragma acc data default(none)
79
{
@@ -33,5 +35,79 @@ void acc_data(void) {
3335
// CHECK-NEXT: acc.terminator
3436
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}
3537

38+
#pragma acc data default(none) async
39+
{}
40+
// CHECK-NEXT: acc.data {
41+
// CHECK-NEXT: acc.terminator
42+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>], defaultAttr = #acc<defaultvalue none>}
43+
44+
#pragma acc data default(none) async(cond)
45+
{}
46+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
47+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
48+
// CHECK-NEXT: acc.data async(%[[CONV_CAST]] : si32) {
49+
// CHECK-NEXT: acc.terminator
50+
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
51+
52+
#pragma acc data default(none) async device_type(nvidia, radeon) async
53+
{}
54+
// CHECK-NEXT: acc.data {
55+
// CHECK-NEXT: acc.terminator
56+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>], defaultAttr = #acc<defaultvalue none>}
57+
58+
#pragma acc data default(none) async(3) device_type(nvidia, radeon) async(cond)
59+
{}
60+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
61+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
62+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
63+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
64+
// CHECK-NEXT: acc.data async(%[[THREE_CAST]] : si32, %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
65+
// CHECK-NEXT: acc.terminator
66+
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
67+
68+
#pragma acc data default(none) async device_type(nvidia, radeon) async(cond)
69+
{}
70+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
71+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
72+
// CHECK-NEXT: acc.data async(%[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
73+
// CHECK-NEXT: acc.terminator
74+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>], defaultAttr = #acc<defaultvalue none>}
75+
76+
#pragma acc data default(none) async(3) device_type(nvidia, radeon) async
77+
{}
78+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
79+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
80+
// CHECK-NEXT: acc.data async(%[[THREE_CAST]] : si32) {
81+
// CHECK-NEXT: acc.terminator
82+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<nvidia>, #acc.device_type<radeon>], defaultAttr = #acc<defaultvalue none>}
83+
84+
#pragma acc data default(none) if(cond)
85+
{}
86+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
87+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
88+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
89+
// CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {
90+
// CHECK-NEXT: acc.terminator
91+
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
92+
93+
#pragma acc data default(none) if(1)
94+
{}
95+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
96+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ONE_LITERAL]] : !s32i), !cir.bool
97+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
98+
// CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {
99+
// CHECK-NEXT: acc.terminator
100+
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
101+
102+
#pragma acc data default(none) if(cond == 1)
103+
{}
104+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
105+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
106+
// CHECK-NEXT: %[[EQ_RES:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
107+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES]] : !cir.bool to i1
108+
// CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {
109+
// CHECK-NEXT: acc.terminator
110+
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
111+
36112
// CHECK-NEXT: cir.return
37113
}

‎clang/test/SemaOpenACC/data-construct-async-clause.c

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,4 +31,41 @@ void Test() {
3131
// expected-error@+2{{expected ')'}}
3232
// expected-note@+1{{to match this '('}}
3333
#pragma acc enter data copyin(I) async(I, I)
34+
//
35+
// expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'data' directive}}
36+
// expected-note@+1{{previous clause is here}}
37+
#pragma acc data default(none) async async
38+
while(1);
39+
40+
// expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'data' directive}}
41+
// expected-note@+1{{previous clause is here}}
42+
#pragma acc data default(none) async(1) async(2)
43+
while(1);
44+
45+
// expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'data' directive}}
46+
// expected-note@+1{{previous clause is here}}
47+
#pragma acc data default(none) async(1) async(2)
48+
while(1);
49+
50+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
51+
// expected-note@+2{{previous clause is here}}
52+
// expected-note@+1{{previous clause is here}}
53+
#pragma acc data default(none) async(1) device_type(*) async(1) async(2)
54+
while(1);
55+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
56+
// expected-note@+2{{previous clause is here}}
57+
// expected-note@+1{{previous clause is here}}
58+
#pragma acc data default(none) async device_type(*) async async
59+
while(1);
60+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
61+
// expected-note@+2{{previous clause is here}}
62+
// expected-note@+1{{previous clause is here}}
63+
#pragma acc data default(none) async(1) device_type(*) async async(2)
64+
while(1);
65+
66+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
67+
// expected-note@+2{{previous clause is here}}
68+
// expected-note@+1{{previous clause is here}}
69+
#pragma acc data default(none) device_type(*) async async
70+
while(1);
3471
}

0 commit comments

Comments
 (0)