Skip to content

Commit 33d2082

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] Implement enter-data + clause lowering (#146146)
'enter data' is a new construct type that requires one of the data clauses, so we had to wait for all clauses to be ready before we could commit this. Most of the clauses are simple, but there is a little bit of work to get 'async' and 'wait' to have similar interfaces in the ACC dialect, where helpers were added.
1 parent 8d2034c commit 33d2082

File tree

5 files changed

+247
-18
lines changed

5 files changed

+247
-18
lines changed

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

Lines changed: 52 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -376,10 +376,19 @@ class OpenACCClauseCIREmitter final
376376
// on all operation types.
377377
mlir::ArrayAttr getAsyncOnlyAttr() {
378378
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
379-
mlir::acc::KernelsOp, mlir::acc::DataOp>)
379+
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
380380
return operation.getAsyncOnlyAttr();
381-
else if constexpr (isCombinedType<OpTy>)
381+
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
382+
if (!operation.getAsyncAttr())
383+
return mlir::ArrayAttr{};
384+
385+
llvm::SmallVector<mlir::Attribute> devTysTemp;
386+
devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
387+
builder.getContext(), mlir::acc::DeviceType::None));
388+
return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
389+
} else if constexpr (isCombinedType<OpTy>) {
382390
return operation.computeOp.getAsyncOnlyAttr();
391+
}
383392

384393
// Note: 'wait' has async as well, but it cannot have data clauses, so we
385394
// don't have to handle them here.
@@ -391,10 +400,19 @@ class OpenACCClauseCIREmitter final
391400
// on all operation types.
392401
mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
393402
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
394-
mlir::acc::KernelsOp, mlir::acc::DataOp>)
403+
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
395404
return operation.getAsyncOperandsDeviceTypeAttr();
396-
else if constexpr (isCombinedType<OpTy>)
405+
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
406+
if (!operation.getAsyncOperand())
407+
return mlir::ArrayAttr{};
408+
409+
llvm::SmallVector<mlir::Attribute> devTysTemp;
410+
devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
411+
builder.getContext(), mlir::acc::DeviceType::None));
412+
return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
413+
} else if constexpr (isCombinedType<OpTy>) {
397414
return operation.computeOp.getAsyncOperandsDeviceTypeAttr();
415+
}
398416

399417
// Note: 'wait' has async as well, but it cannot have data clauses, so we
400418
// don't have to handle them here.
@@ -409,6 +427,8 @@ class OpenACCClauseCIREmitter final
409427
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
410428
mlir::acc::KernelsOp, mlir::acc::DataOp>)
411429
return operation.getAsyncOperands();
430+
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>)
431+
return operation.getAsyncOperandMutable();
412432
else if constexpr (isCombinedType<OpTy>)
413433
return operation.computeOp.getAsyncOperands();
414434

@@ -542,10 +562,11 @@ class OpenACCClauseCIREmitter final
542562
void VisitAsyncClause(const OpenACCAsyncClause &clause) {
543563
hasAsyncClause = true;
544564
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
545-
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
546-
if (!clause.hasIntExpr())
565+
mlir::acc::KernelsOp, mlir::acc::DataOp,
566+
mlir::acc::EnterDataOp>) {
567+
if (!clause.hasIntExpr()) {
547568
operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
548-
else {
569+
} else {
549570

550571
mlir::Value intExpr;
551572
{
@@ -572,8 +593,8 @@ class OpenACCClauseCIREmitter final
572593
applyToComputeOp(clause);
573594
} else {
574595
// TODO: When we've implemented this for everything, switch this to an
575-
// unreachable. Combined constructs remain. Data, enter data, exit data,
576-
// update constructs remain.
596+
// unreachable. Combined constructs remain. Exit data, update constructs
597+
// remain.
577598
return clauseNotImplemented(clause);
578599
}
579600
}
@@ -604,7 +625,7 @@ class OpenACCClauseCIREmitter final
604625
mlir::acc::KernelsOp, mlir::acc::InitOp,
605626
mlir::acc::ShutdownOp, mlir::acc::SetOp,
606627
mlir::acc::DataOp, mlir::acc::WaitOp,
607-
mlir::acc::HostDataOp>) {
628+
mlir::acc::HostDataOp, mlir::acc::EnterDataOp>) {
608629
operation.getIfCondMutable().append(
609630
createCondition(clause.getConditionExpr()));
610631
} else if constexpr (isCombinedType<OpTy>) {
@@ -659,7 +680,8 @@ class OpenACCClauseCIREmitter final
659680

660681
void VisitWaitClause(const OpenACCWaitClause &clause) {
661682
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
662-
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
683+
mlir::acc::KernelsOp, mlir::acc::DataOp,
684+
mlir::acc::EnterDataOp>) {
663685
if (!clause.hasExprs()) {
664686
operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
665687
} else {
@@ -866,11 +888,16 @@ class OpenACCClauseCIREmitter final
866888
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
867889
/*structured=*/true,
868890
/*implicit=*/false);
891+
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
892+
for (const Expr *var : clause.getVarList())
893+
addDataOperand<mlir::acc::CopyinOp>(
894+
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
895+
/*structured=*/false, /*implicit=*/false);
869896
} else if constexpr (isCombinedType<OpTy>) {
870897
applyToComputeOp(clause);
871898
} else {
872899
// TODO: When we've implemented this for everything, switch this to an
873-
// unreachable. enter-data, declare constructs remain.
900+
// unreachable. declare construct remains.
874901
return clauseNotImplemented(clause);
875902
}
876903
}
@@ -900,11 +927,16 @@ class OpenACCClauseCIREmitter final
900927
var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
901928
/*structured=*/true,
902929
/*implicit=*/false);
930+
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
931+
for (const Expr *var : clause.getVarList())
932+
addDataOperand<mlir::acc::CreateOp>(
933+
var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
934+
/*structured=*/false, /*implicit=*/false);
903935
} else if constexpr (isCombinedType<OpTy>) {
904936
applyToComputeOp(clause);
905937
} else {
906938
// TODO: When we've implemented this for everything, switch this to an
907-
// unreachable. enter-data, declare constructs remain.
939+
// unreachable. declare construct remains.
908940
return clauseNotImplemented(clause);
909941
}
910942
}
@@ -974,12 +1006,15 @@ class OpenACCClauseCIREmitter final
9741006
addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
9751007
var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
9761008
/*implicit=*/false);
1009+
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
1010+
for (const Expr *var : clause.getVarList())
1011+
addDataOperand<mlir::acc::AttachOp>(
1012+
var, mlir::acc::DataClause::acc_attach, {},
1013+
/*structured=*/false, /*implicit=*/false);
9771014
} else if constexpr (isCombinedType<OpTy>) {
9781015
applyToComputeOp(clause);
9791016
} else {
980-
// TODO: When we've implemented this for everything, switch this to an
981-
// unreachable. enter data remains.
982-
return clauseNotImplemented(clause);
1017+
llvm_unreachable("Unknown construct kind in VisitAttachClause");
9831018
}
9841019
}
9851020
};
@@ -1018,6 +1053,7 @@ EXPL_SPEC(mlir::acc::ShutdownOp)
10181053
EXPL_SPEC(mlir::acc::SetOp)
10191054
EXPL_SPEC(mlir::acc::WaitOp)
10201055
EXPL_SPEC(mlir::acc::HostDataOp)
1056+
EXPL_SPEC(mlir::acc::EnterDataOp)
10211057
#undef EXPL_SPEC
10221058

10231059
template <typename ComputeOp, typename LoopOp>

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -250,8 +250,10 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
250250

251251
mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
252252
const OpenACCEnterDataConstruct &s) {
253-
cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
254-
return mlir::failure();
253+
mlir::Location start = getLoc(s.getSourceRange().getBegin());
254+
emitOpenACCOp<EnterDataOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
255+
s.clauses());
256+
return mlir::success();
255257
}
256258
mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
257259
const OpenACCExitDataConstruct &s) {

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

Lines changed: 125 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
2+
void acc_data(int parmVar, int *ptrParmVar) {
3+
// CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) {
4+
// CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]
5+
// CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]
6+
// CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>
7+
// CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
8+
9+
#pragma acc enter data copyin(parmVar)
10+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
11+
// CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
12+
13+
#pragma acc enter data copyin(readonly, alwaysin: parmVar)
14+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
15+
// CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
16+
17+
#pragma acc enter data copyin(readonly, alwaysin: parmVar) async
18+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
19+
// CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
20+
21+
#pragma acc enter data async copyin(readonly, alwaysin: parmVar)
22+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
23+
// CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
24+
25+
#pragma acc enter data copyin(readonly, alwaysin: parmVar) async(parmVar)
26+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
27+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
28+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
29+
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
30+
31+
#pragma acc enter data async(parmVar) copyin(readonly, alwaysin: parmVar)
32+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
33+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
34+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
35+
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
36+
37+
#pragma acc enter data create(parmVar)
38+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
39+
// CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
40+
41+
#pragma acc enter data create(zero: parmVar)
42+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}
43+
// CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
44+
45+
#pragma acc enter data create(zero: parmVar) async
46+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}
47+
// CHECK-NEXT: acc.enter_data async dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
48+
49+
#pragma acc enter data create(zero: parmVar) async(parmVar)
50+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
51+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
52+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}
53+
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
54+
55+
#pragma acc enter data attach(ptrParmVar)
56+
// CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
57+
// CHECK-NEXT: acc.enter_data dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)
58+
59+
#pragma acc enter data attach(ptrParmVar) async
60+
// CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
61+
// CHECK-NEXT: acc.enter_data async dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)
62+
63+
#pragma acc enter data attach(ptrParmVar) async(parmVar)
64+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
65+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
66+
// CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
67+
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)
68+
69+
#pragma acc enter data if (parmVar == 1) copyin(parmVar)
70+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
71+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
72+
// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
73+
// CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
74+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
75+
// CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
76+
77+
#pragma acc enter data async if (parmVar == 1) copyin(parmVar)
78+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
79+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
80+
// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
81+
// CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
82+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
83+
// CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
84+
85+
#pragma acc enter data if (parmVar == 1) async(parmVar) copyin(parmVar)
86+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
87+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
88+
// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
89+
// CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
90+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
91+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
92+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
93+
// CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
94+
95+
#pragma acc enter data wait create(parmVar)
96+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
97+
// CHECK-NEXT: acc.enter_data wait dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
98+
99+
#pragma acc enter data wait(1) create(parmVar)
100+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
101+
// CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
102+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
103+
// CHECK-NEXT: acc.enter_data wait(%[[ONE_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
104+
105+
#pragma acc enter data wait(parmVar, 1, 2) create(parmVar)
106+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
107+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
108+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
109+
// CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
110+
// CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
111+
// CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]
112+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
113+
// CHECK-NEXT: acc.enter_data wait(%[[PARM_CAST]], %[[ONE_CAST]], %[[TWO_CAST]] : si32, si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
114+
115+
#pragma acc enter data wait(devnum: parmVar: 1, 2) create(parmVar)
116+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
117+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
118+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
119+
// CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
120+
// CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
121+
// CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]
122+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
123+
// CHECK-NEXT: acc.enter_data wait_devnum(%[[PARM_CAST]] : si32) wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
124+
125+
}

‎mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2010,6 +2010,25 @@ def OpenACC_EnterDataOp : OpenACC_Op<"enter_data",
20102010

20112011
/// The i-th data operand passed.
20122012
Value getDataOperand(unsigned i);
2013+
2014+
/// Add an entry to the 'async-only' attribute (clause spelled without
2015+
/// arguments). DeviceType array is supplied even though it should always be
2016+
/// empty, so this can mirror other versions of this function.
2017+
void addAsyncOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
2018+
/// Add a value to the 'async'. DeviceType array is supplied even though it
2019+
/// should always be empty, so this can mirror other versions of this
2020+
/// function.
2021+
void addAsyncOperand(MLIRContext *, mlir::Value,
2022+
llvm::ArrayRef<DeviceType>);
2023+
/// Add an entry to the 'wait-only' attribute (clause spelled without
2024+
/// arguments). DeviceType array is supplied even though it should always be
2025+
/// empty, so this can mirror other versions of this function.
2026+
void addWaitOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
2027+
/// Add an array-like entry to the 'wait'. DeviceType array is supplied
2028+
/// even though it should always be empty, so this can mirror other versions
2029+
/// of this function.
2030+
void addWaitOperands(MLIRContext *, bool hasDevnum, mlir::ValueRange,
2031+
llvm::ArrayRef<DeviceType>);
20132032
}];
20142033

20152034
let assemblyFormat = [{

0 commit comments

Comments
 (0)