Skip to content

Commit 69bbf21

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] Implement 'modifier-list' lowering (#145770)
Some of the 'data' clauses can have a 'modifier-list' which specifies one of a few keywords from a list. This patch adds support for lowering them following #144806. We have to keep a separate enum from MLIR, since we have to keep 'always' around for semantic reasons, whereas the dialect doesn't differentiate these. This patch ensures we get these right for the only applicable clause so far, which is 'copy'.
1 parent 2e39959 commit 69bbf21

File tree

4 files changed

+171
-38
lines changed

4 files changed

+171
-38
lines changed

‎clang/include/clang/Basic/OpenACCKinds.h‎

Lines changed: 11 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -634,16 +634,19 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
634634
}
635635

636636
// Represents the 'modifier' of a 'modifier-list', as applied to copy, copyin,
637-
// copyout, and create. Implemented as a 'bitmask'
637+
// copyout, and create. Implemented as a 'bitmask'.
638+
// Note: This attempts to synchronize with mlir::acc::DataClauseModifier,
639+
// however has to store `Always` separately(whereas MLIR has it as AlwaysIn &
640+
// AlwaysOut). However, we keep them in sync so that we can cast between them.
638641
enum class OpenACCModifierKind : uint8_t {
639642
Invalid = 0,
640-
Always = 1 << 0,
641-
AlwaysIn = 1 << 1,
642-
AlwaysOut = 1 << 2,
643-
Readonly = 1 << 3,
644-
Zero = 1 << 4,
645-
Capture = 1 << 5,
646-
LLVM_MARK_AS_BITMASK_ENUM(Capture)
643+
Zero = 1 << 0,
644+
Readonly = 1 << 1,
645+
AlwaysIn = 1 << 2,
646+
AlwaysOut = 1 << 3,
647+
Capture = 1 << 4,
648+
Always = 1 << 5,
649+
LLVM_MARK_AS_BITMASK_ENUM(Always)
647650
};
648651

649652
inline bool isOpenACCModifierBitSet(OpenACCModifierKind List,

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

Lines changed: 44 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -286,16 +286,39 @@ class OpenACCClauseCIREmitter final
286286
std::move(bounds)};
287287
}
288288

289+
mlir::acc::DataClauseModifier
290+
convertModifiers(OpenACCModifierKind modifiers) {
291+
using namespace mlir::acc;
292+
static_assert(static_cast<int>(OpenACCModifierKind::Zero) ==
293+
static_cast<int>(DataClauseModifier::zero) &&
294+
static_cast<int>(OpenACCModifierKind::Readonly) ==
295+
static_cast<int>(DataClauseModifier::readonly) &&
296+
static_cast<int>(OpenACCModifierKind::AlwaysIn) ==
297+
static_cast<int>(DataClauseModifier::alwaysin) &&
298+
static_cast<int>(OpenACCModifierKind::AlwaysOut) ==
299+
static_cast<int>(DataClauseModifier::alwaysout) &&
300+
static_cast<int>(OpenACCModifierKind::Capture) ==
301+
static_cast<int>(DataClauseModifier::capture));
302+
303+
DataClauseModifier mlirModifiers{};
304+
305+
// The MLIR representation of this represents `always` as `alwaysin` +
306+
// `alwaysout`. So do a small fixup here.
307+
if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) {
308+
mlirModifiers = mlirModifiers | DataClauseModifier::always;
309+
modifiers &= ~OpenACCModifierKind::Always;
310+
}
311+
312+
mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
313+
return mlirModifiers;
314+
}
315+
289316
template <typename BeforeOpTy, typename AfterOpTy>
290317
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
291-
bool structured, bool implicit) {
318+
OpenACCModifierKind modifiers, bool structured,
319+
bool implicit) {
292320
DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
293321

294-
// TODO: OpenACC: we should comprehend the 'modifier-list' here for the data
295-
// operand. At the moment, we don't have a uniform way to assign these
296-
// properly, and the dialect cannot represent anything other than 'readonly'
297-
// and 'zero' on copyin/copyout/create, so for now, we skip it.
298-
299322
auto beforeOp =
300323
builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
301324
implicit, opInfo.name, opInfo.bounds);
@@ -323,6 +346,8 @@ class OpenACCClauseCIREmitter final
323346
// Set the 'rest' of the info for both operations.
324347
beforeOp.setDataClause(dataClause);
325348
afterOp.setDataClause(dataClause);
349+
beforeOp.setModifiers(convertModifiers(modifiers));
350+
afterOp.setModifiers(convertModifiers(modifiers));
326351

327352
// Make sure we record these, so 'async' values can be updated later.
328353
dataOperands.push_back(beforeOp.getOperation());
@@ -331,7 +356,8 @@ class OpenACCClauseCIREmitter final
331356

332357
template <typename BeforeOpTy>
333358
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
334-
bool structured, bool implicit) {
359+
OpenACCModifierKind modifiers, bool structured,
360+
bool implicit) {
335361
DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
336362
auto beforeOp =
337363
builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
@@ -340,6 +366,8 @@ class OpenACCClauseCIREmitter final
340366

341367
// Set the 'rest' of the info for the operation.
342368
beforeOp.setDataClause(dataClause);
369+
beforeOp.setModifiers(convertModifiers(modifiers));
370+
343371
// Make sure we record these, so 'async' values can be updated later.
344372
dataOperands.push_back(beforeOp.getOperation());
345373
}
@@ -818,7 +846,8 @@ class OpenACCClauseCIREmitter final
818846
mlir::acc::KernelsOp>) {
819847
for (auto var : clause.getVarList())
820848
addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
821-
var, mlir::acc::DataClause::acc_copy, /*structured=*/true,
849+
var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
850+
/*structured=*/true,
822851
/*implicit=*/false);
823852
} else if constexpr (isCombinedType<OpTy>) {
824853
applyToComputeOp(clause);
@@ -833,8 +862,8 @@ class OpenACCClauseCIREmitter final
833862
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
834863
for (auto var : clause.getVarList())
835864
addDataOperand<mlir::acc::UseDeviceOp>(
836-
var, mlir::acc::DataClause::acc_use_device,
837-
/*structured=*/true, /*implicit=*/false);
865+
var, mlir::acc::DataClause::acc_use_device, {}, /*structured=*/true,
866+
/*implicit=*/false);
838867
} else {
839868
llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
840869
}
@@ -845,7 +874,8 @@ class OpenACCClauseCIREmitter final
845874
mlir::acc::KernelsOp>) {
846875
for (auto var : clause.getVarList())
847876
addDataOperand<mlir::acc::DevicePtrOp>(
848-
var, mlir::acc::DataClause::acc_deviceptr, /*structured=*/true,
877+
var, mlir::acc::DataClause::acc_deviceptr, {},
878+
/*structured=*/true,
849879
/*implicit=*/false);
850880
} else if constexpr (isCombinedType<OpTy>) {
851881
applyToComputeOp(clause);
@@ -861,7 +891,7 @@ class OpenACCClauseCIREmitter final
861891
mlir::acc::KernelsOp>) {
862892
for (auto var : clause.getVarList())
863893
addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
864-
var, mlir::acc::DataClause::acc_no_create, /*structured=*/true,
894+
var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
865895
/*implicit=*/false);
866896
} else if constexpr (isCombinedType<OpTy>) {
867897
applyToComputeOp(clause);
@@ -877,7 +907,7 @@ class OpenACCClauseCIREmitter final
877907
mlir::acc::KernelsOp>) {
878908
for (auto var : clause.getVarList())
879909
addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
880-
var, mlir::acc::DataClause::acc_present, /*structured=*/true,
910+
var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
881911
/*implicit=*/false);
882912
} else if constexpr (isCombinedType<OpTy>) {
883913
applyToComputeOp(clause);
@@ -893,7 +923,7 @@ class OpenACCClauseCIREmitter final
893923
mlir::acc::KernelsOp>) {
894924
for (auto var : clause.getVarList())
895925
addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
896-
var, mlir::acc::DataClause::acc_attach, /*structured=*/true,
926+
var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
897927
/*implicit=*/false);
898928
} else if constexpr (isCombinedType<OpTy>) {
899929
applyToComputeOp(clause);

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

Lines changed: 65 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -77,29 +77,29 @@ void acc_compute(int parmVar) {
7777
// these do nothing to the IR.
7878
#pragma acc parallel loop copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3)
7979
for(int i = 0; i < 5; ++i);
80-
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc
81-
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc
82-
// CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc
80+
// 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
81+
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc
82+
// CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"} loc
8383
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>) {
8484
// CHECK-NEXT: acc.loop combined(parallel) {
8585
// CHECK: acc.yield
8686
// CHECK-NEXT: }
8787
// CHECK-NEXT: acc.yield
8888
// CHECK-NEXT: } loc
89-
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc
90-
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc
91-
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc
89+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"} loc
90+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc
91+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc
9292

9393
#pragma acc serial loop copy(always, alwaysin, alwaysout: localVar1)
9494
for(int i = 0; i < 5; ++i);
95-
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc
95+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc
9696
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
9797
// CHECK-NEXT: acc.loop combined(serial) {
9898
// CHECK: acc.yield
9999
// CHECK-NEXT: }
100100
// CHECK-NEXT: acc.yield
101101
// CHECK-NEXT: } loc
102-
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc
102+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc
103103

104104
short *localPointer;
105105
float localArray[5];
@@ -1102,3 +1102,60 @@ void copy_member_of_array_element_member() {
11021102
// CHECK-NEXT: } loc
11031103
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETB]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "outer.inner[2].b"}
11041104
}
1105+
1106+
void modifier_list() {
1107+
// CHECK: cir.func @modifier_list() {
1108+
int localVar;
1109+
// CHECK-NEXT: %[[LOCALVAR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["localVar"]
1110+
1111+
#pragma acc parallel loop copy(always:localVar)
1112+
for(int i = 0; i < 5; ++i);
1113+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"}
1114+
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
1115+
// CHECK-NEXT: acc.loop combined(parallel) {
1116+
// CHECK: acc.yield
1117+
// CHECK-NEXT: } loc
1118+
// CHECK-NEXT: acc.yield
1119+
// CHECK-NEXT: } loc
1120+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"}
1121+
#pragma acc serial loop copy(alwaysin:localVar)
1122+
for(int i = 0; i < 5; ++i);
1123+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"}
1124+
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
1125+
// CHECK-NEXT: acc.loop combined(serial) {
1126+
// CHECK: acc.yield
1127+
// CHECK-NEXT: } loc
1128+
// CHECK-NEXT: acc.yield
1129+
// CHECK-NEXT: } loc
1130+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"}
1131+
#pragma acc kernels loop copy(alwaysout:localVar)
1132+
for(int i = 0; i < 5; ++i);
1133+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"}
1134+
// CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
1135+
// CHECK-NEXT: acc.loop combined(kernels) {
1136+
// CHECK: acc.yield
1137+
// CHECK-NEXT: } loc
1138+
// CHECK-NEXT: acc.terminator
1139+
// CHECK-NEXT: } loc
1140+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"}
1141+
#pragma acc parallel loop copy(capture:localVar)
1142+
for(int i = 0; i < 5; ++i);
1143+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"}
1144+
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
1145+
// CHECK-NEXT: acc.loop combined(parallel) {
1146+
// CHECK: acc.yield
1147+
// CHECK-NEXT: } loc
1148+
// CHECK-NEXT: acc.yield
1149+
// CHECK-NEXT: } loc
1150+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"}
1151+
#pragma acc serial loop copy(capture, always, alwaysin, alwaysout:localVar)
1152+
for(int i = 0; i < 5; ++i);
1153+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"}
1154+
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
1155+
// CHECK-NEXT: acc.loop combined(serial) {
1156+
// CHECK: acc.yield
1157+
// CHECK-NEXT: } loc
1158+
// CHECK-NEXT: acc.yield
1159+
// CHECK-NEXT: } loc
1160+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"}
1161+
}

0 commit comments

Comments
 (0)