Skip to content

Commit 554256a

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] Implement 'async' lowering.
Async acts just like num_workers/vector_length in that it gets a new variant per device_type and is lowered as an operand. However, it has one additional complication, in that it can have a variant that has no argument, which produces an attribute with the correct devicetype. Additionally, this syncronizes us with the implementation of flang, which prohibits multiple 'async' clauses per-device_type.
1 parent de0153d commit 554256a

File tree

6 files changed

+216
-9
lines changed

6 files changed

+216
-9
lines changed

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

Lines changed: 33 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -97,11 +97,13 @@ class OpenACCClauseCIREmitter final
9797

9898
// Handle a clause affected by the 'device-type' to the point that they need
9999
// to have the attributes added in the correct/corresponding order, such as
100-
// 'num_workers' or 'vector_length' on a compute construct.
101-
mlir::ArrayAttr
102-
handleDeviceTypeAffectedClause(mlir::ArrayAttr existingDeviceTypes,
103-
mlir::Value argument,
104-
mlir::MutableOperandRange &argCollection) {
100+
// 'num_workers' or 'vector_length' on a compute construct. For cases where we
101+
// don't have an argument that needs to be added to an additional one (such as
102+
// asyncOnly) we can use this with 'argument' as std::nullopt.
103+
mlir::ArrayAttr handleDeviceTypeAffectedClause(
104+
mlir::ArrayAttr existingDeviceTypes,
105+
std::optional<mlir::Value> argument = std::nullopt,
106+
mlir::MutableOperandRange *argCollection = nullptr) {
105107
llvm::SmallVector<mlir::Attribute> deviceTypes;
106108

107109
// Collect the 'existing' device-type attributes so we can re-create them
@@ -120,13 +122,19 @@ class OpenACCClauseCIREmitter final
120122
lastDeviceTypeClause->getArchitectures()) {
121123
deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
122124
builder.getContext(), decodeDeviceType(arch.getIdentifierInfo())));
123-
argCollection.append(argument);
125+
if (argument) {
126+
assert(argCollection);
127+
argCollection->append(*argument);
128+
}
124129
}
125130
} else {
126131
// Else, we just add a single for 'none'.
127132
deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
128133
builder.getContext(), mlir::acc::DeviceType::None));
129-
argCollection.append(argument);
134+
if (argument) {
135+
assert(argCollection);
136+
argCollection->append(*argument);
137+
}
130138
}
131139

132140
return mlir::ArrayAttr::get(builder.getContext(), deviceTypes);
@@ -205,7 +213,7 @@ class OpenACCClauseCIREmitter final
205213
mlir::MutableOperandRange range = operation.getNumWorkersMutable();
206214
operation.setNumWorkersDeviceTypeAttr(handleDeviceTypeAffectedClause(
207215
operation.getNumWorkersDeviceTypeAttr(),
208-
createIntExpr(clause.getIntExpr()), range));
216+
createIntExpr(clause.getIntExpr()), &range));
209217
} else if constexpr (isOneOfTypes<OpTy, SerialOp>) {
210218
llvm_unreachable("num_workers not valid on serial");
211219
} else {
@@ -218,14 +226,30 @@ class OpenACCClauseCIREmitter final
218226
mlir::MutableOperandRange range = operation.getVectorLengthMutable();
219227
operation.setVectorLengthDeviceTypeAttr(handleDeviceTypeAffectedClause(
220228
operation.getVectorLengthDeviceTypeAttr(),
221-
createIntExpr(clause.getIntExpr()), range));
229+
createIntExpr(clause.getIntExpr()), &range));
222230
} else if constexpr (isOneOfTypes<OpTy, SerialOp>) {
223231
llvm_unreachable("vector_length not valid on serial");
224232
} else {
225233
return clauseNotImplemented(clause);
226234
}
227235
}
228236

237+
void VisitAsyncClause(const OpenACCAsyncClause &clause) {
238+
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
239+
if (!clause.hasIntExpr()) {
240+
operation.setAsyncOnlyAttr(
241+
handleDeviceTypeAffectedClause(operation.getAsyncOnlyAttr()));
242+
} else {
243+
mlir::MutableOperandRange range = operation.getAsyncOperandsMutable();
244+
operation.setAsyncOperandsDeviceTypeAttr(handleDeviceTypeAffectedClause(
245+
operation.getAsyncOperandsDeviceTypeAttr(),
246+
createIntExpr(clause.getIntExpr()), &range));
247+
}
248+
} else {
249+
return clauseNotImplemented(clause);
250+
}
251+
}
252+
229253
void VisitSelfClause(const OpenACCSelfClause &clause) {
230254
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
231255
if (clause.isEmptySelfClause()) {

‎clang/lib/Sema/SemaOpenACCClause.cpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -639,6 +639,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
639639

640640
OpenACCClause *SemaOpenACCClauseVisitor::VisitAsyncClause(
641641
SemaOpenACC::OpenACCParsedClause &Clause) {
642+
if (DisallowSinceLastDeviceType<OpenACCAsyncClause>(Clause))
643+
return nullptr;
644+
642645
assert(Clause.getNumIntExprs() < 2 &&
643646
"Invalid number of expressions for Async");
644647
return OpenACCAsyncClause::Create(

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

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -210,5 +210,51 @@ void acc_kernels(int cond) {
210210
// CHECK-NEXT: acc.terminator
211211
// CHECK-NEXT: } loc
212212

213+
#pragma acc kernels async
214+
{}
215+
// CHECK-NEXT: acc.kernels {
216+
// CHECK-NEXT: acc.terminator
217+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>]}
218+
219+
#pragma acc kernels async(cond)
220+
{}
221+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
222+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
223+
// CHECK-NEXT: acc.kernels async(%[[CONV_CAST]] : si32) {
224+
// CHECK-NEXT: acc.terminator
225+
// CHECK-NEXT: } loc
226+
227+
#pragma acc kernels async device_type(nvidia, radeon) async
228+
{}
229+
// CHECK-NEXT: acc.kernels {
230+
// CHECK-NEXT: acc.terminator
231+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]}
232+
233+
#pragma acc kernels async(3) device_type(nvidia, radeon) async(cond)
234+
{}
235+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
236+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
237+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
238+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
239+
// CHECK-NEXT: acc.kernels async(%[[THREE_CAST]] : si32, %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
240+
// CHECK-NEXT: acc.terminator
241+
// CHECK-NEXT: } loc
242+
243+
#pragma acc kernels async device_type(nvidia, radeon) async(cond)
244+
{}
245+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
246+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
247+
// CHECK-NEXT: acc.kernels async(%[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
248+
// CHECK-NEXT: acc.terminator
249+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>]}
250+
251+
#pragma acc kernels async(3) device_type(nvidia, radeon) async
252+
{}
253+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
254+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
255+
// CHECK-NEXT: acc.kernels async(%[[THREE_CAST]] : si32) {
256+
// CHECK-NEXT: acc.terminator
257+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<nvidia>, #acc.device_type<radeon>]}
258+
213259
// CHECK-NEXT: cir.return
214260
}

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

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -209,5 +209,51 @@ void acc_parallel(int cond) {
209209
// CHECK-NEXT: acc.yield
210210
// CHECK-NEXT: } loc
211211

212+
#pragma acc parallel async
213+
{}
214+
// CHECK-NEXT: acc.parallel {
215+
// CHECK-NEXT: acc.yield
216+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>]}
217+
218+
#pragma acc parallel async(cond)
219+
{}
220+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
221+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
222+
// CHECK-NEXT: acc.parallel async(%[[CONV_CAST]] : si32) {
223+
// CHECK-NEXT: acc.yield
224+
// CHECK-NEXT: } loc
225+
226+
#pragma acc parallel async device_type(nvidia, radeon) async
227+
{}
228+
// CHECK-NEXT: acc.parallel {
229+
// CHECK-NEXT: acc.yield
230+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]}
231+
232+
#pragma acc parallel async(3) device_type(nvidia, radeon) async(cond)
233+
{}
234+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
235+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
236+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
237+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
238+
// CHECK-NEXT: acc.parallel async(%[[THREE_CAST]] : si32, %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
239+
// CHECK-NEXT: acc.yield
240+
// CHECK-NEXT: } loc
241+
242+
#pragma acc parallel async device_type(nvidia, radeon) async(cond)
243+
{}
244+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
245+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
246+
// CHECK-NEXT: acc.parallel async(%[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
247+
// CHECK-NEXT: acc.yield
248+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>]}
249+
250+
#pragma acc parallel async(3) device_type(nvidia, radeon) async
251+
{}
252+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
253+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
254+
// CHECK-NEXT: acc.parallel async(%[[THREE_CAST]] : si32) {
255+
// CHECK-NEXT: acc.yield
256+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<nvidia>, #acc.device_type<radeon>]}
257+
212258
// CHECK-NEXT: cir.return
213259
}

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

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -106,5 +106,51 @@ void acc_serial(int cond) {
106106
// CHECK-NEXT: acc.yield
107107
// CHECK-NEXT: } loc
108108

109+
#pragma acc serial async
110+
{}
111+
// CHECK-NEXT: acc.serial {
112+
// CHECK-NEXT: acc.yield
113+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>]}
114+
115+
#pragma acc serial async(cond)
116+
{}
117+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
118+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
119+
// CHECK-NEXT: acc.serial async(%[[CONV_CAST]] : si32) {
120+
// CHECK-NEXT: acc.yield
121+
// CHECK-NEXT: } loc
122+
123+
#pragma acc serial async device_type(nvidia, radeon) async
124+
{}
125+
// CHECK-NEXT: acc.serial {
126+
// CHECK-NEXT: acc.yield
127+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]}
128+
129+
#pragma acc serial async(3) device_type(nvidia, radeon) async(cond)
130+
{}
131+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
132+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
133+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
134+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
135+
// CHECK-NEXT: acc.serial async(%[[THREE_CAST]] : si32, %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
136+
// CHECK-NEXT: acc.yield
137+
// CHECK-NEXT: } loc
138+
139+
#pragma acc serial async device_type(nvidia, radeon) async(cond)
140+
{}
141+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
142+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
143+
// CHECK-NEXT: acc.serial async(%[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
144+
// CHECK-NEXT: acc.yield
145+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>]}
146+
147+
#pragma acc serial async(3) device_type(nvidia, radeon) async
148+
{}
149+
// CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
150+
// CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
151+
// CHECK-NEXT: acc.serial async(%[[THREE_CAST]] : si32) {
152+
// CHECK-NEXT: acc.yield
153+
// CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<nvidia>, #acc.device_type<radeon>]}
154+
109155
// CHECK-NEXT: cir.return
110156
}

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

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,48 @@ void Test() {
2020
#pragma acc serial async(1, 2)
2121
while(1);
2222

23+
// expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'kernels' directive}}
24+
// expected-note@+1{{previous clause is here}}
25+
#pragma acc kernels async async
26+
while(1);
27+
28+
// expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'kernels' directive}}
29+
// expected-note@+1{{previous clause is here}}
30+
#pragma acc kernels async(1) async(2)
31+
while(1);
32+
33+
// expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'parallel' directive}}
34+
// expected-note@+1{{previous clause is here}}
35+
#pragma acc parallel async(1) async(2)
36+
while(1);
37+
38+
// expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'serial' directive}}
39+
// expected-note@+1{{previous clause is here}}
40+
#pragma acc serial async(1) async(2)
41+
while(1);
42+
43+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'kernels' directive}}
44+
// expected-note@+2{{previous clause is here}}
45+
// expected-note@+1{{previous clause is here}}
46+
#pragma acc kernels async(1) device_type(*) async(1) async(2)
47+
while(1);
48+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}}
49+
// expected-note@+2{{previous clause is here}}
50+
// expected-note@+1{{previous clause is here}}
51+
#pragma acc parallel async device_type(*) async async
52+
while(1);
53+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'serial' directive}}
54+
// expected-note@+2{{previous clause is here}}
55+
// expected-note@+1{{previous clause is here}}
56+
#pragma acc serial async(1) device_type(*) async async(2)
57+
while(1);
58+
59+
// expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}}
60+
// expected-note@+2{{previous clause is here}}
61+
// expected-note@+1{{previous clause is here}}
62+
#pragma acc parallel device_type(*) async async
63+
while(1);
64+
2365
struct NotConvertible{} NC;
2466
// expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('struct NotConvertible' invalid)}}
2567
#pragma acc parallel async(NC)

0 commit comments

Comments
 (0)