Skip to content

Commit 188633e

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] Implement 'modifier-list' lowering
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 llvm#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 5a194c1 commit 188633e

File tree

4 files changed

+160
-38
lines changed

4 files changed

+160
-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: 33 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -286,16 +286,28 @@ class OpenACCClauseCIREmitter final
286286
std::move(bounds)};
287287
}
288288

289+
mlir::acc::DataClauseModifier
290+
convertModifiers(OpenACCModifierKind modifiers) {
291+
using namespace mlir::acc;
292+
DataClauseModifier mlirModifiers{};
293+
294+
// The MLIR representation of this represents `always` as `alwaysin` +
295+
// `alwaysout`. So do a small fixup here.
296+
if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) {
297+
mlirModifiers = mlirModifiers | DataClauseModifier::always;
298+
modifiers &= ~OpenACCModifierKind::Always;
299+
}
300+
301+
mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
302+
return mlirModifiers;
303+
}
304+
289305
template <typename BeforeOpTy, typename AfterOpTy>
290306
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
291-
bool structured, bool implicit) {
307+
OpenACCModifierKind modifiers, bool structured,
308+
bool implicit) {
292309
DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
293310

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-
299311
auto beforeOp =
300312
builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
301313
implicit, opInfo.name, opInfo.bounds);
@@ -323,6 +335,8 @@ class OpenACCClauseCIREmitter final
323335
// Set the 'rest' of the info for both operations.
324336
beforeOp.setDataClause(dataClause);
325337
afterOp.setDataClause(dataClause);
338+
beforeOp.setModifiers(convertModifiers(modifiers));
339+
afterOp.setModifiers(convertModifiers(modifiers));
326340

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

332346
template <typename BeforeOpTy>
333347
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
334-
bool structured, bool implicit) {
348+
OpenACCModifierKind modifiers, bool structured,
349+
bool implicit) {
335350
DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
336351
auto beforeOp =
337352
builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
@@ -340,6 +355,8 @@ class OpenACCClauseCIREmitter final
340355

341356
// Set the 'rest' of the info for the operation.
342357
beforeOp.setDataClause(dataClause);
358+
beforeOp.setModifiers(convertModifiers(modifiers));
359+
343360
// Make sure we record these, so 'async' values can be updated later.
344361
dataOperands.push_back(beforeOp.getOperation());
345362
}
@@ -818,7 +835,8 @@ class OpenACCClauseCIREmitter final
818835
mlir::acc::KernelsOp>) {
819836
for (auto var : clause.getVarList())
820837
addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
821-
var, mlir::acc::DataClause::acc_copy, /*structured=*/true,
838+
var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
839+
/*structured=*/true,
822840
/*implicit=*/false);
823841
} else if constexpr (isCombinedType<OpTy>) {
824842
applyToComputeOp(clause);
@@ -833,8 +851,8 @@ class OpenACCClauseCIREmitter final
833851
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
834852
for (auto var : clause.getVarList())
835853
addDataOperand<mlir::acc::UseDeviceOp>(
836-
var, mlir::acc::DataClause::acc_use_device,
837-
/*structured=*/true, /*implicit=*/false);
854+
var, mlir::acc::DataClause::acc_use_device, {}, /*structured=*/true,
855+
/*implicit=*/false);
838856
} else {
839857
llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
840858
}
@@ -845,7 +863,8 @@ class OpenACCClauseCIREmitter final
845863
mlir::acc::KernelsOp>) {
846864
for (auto var : clause.getVarList())
847865
addDataOperand<mlir::acc::DevicePtrOp>(
848-
var, mlir::acc::DataClause::acc_deviceptr, /*structured=*/true,
866+
var, mlir::acc::DataClause::acc_deviceptr, {},
867+
/*structured=*/true,
849868
/*implicit=*/false);
850869
} else if constexpr (isCombinedType<OpTy>) {
851870
applyToComputeOp(clause);
@@ -861,7 +880,7 @@ class OpenACCClauseCIREmitter final
861880
mlir::acc::KernelsOp>) {
862881
for (auto var : clause.getVarList())
863882
addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
864-
var, mlir::acc::DataClause::acc_no_create, /*structured=*/true,
883+
var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
865884
/*implicit=*/false);
866885
} else if constexpr (isCombinedType<OpTy>) {
867886
applyToComputeOp(clause);
@@ -877,7 +896,7 @@ class OpenACCClauseCIREmitter final
877896
mlir::acc::KernelsOp>) {
878897
for (auto var : clause.getVarList())
879898
addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
880-
var, mlir::acc::DataClause::acc_present, /*structured=*/true,
899+
var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
881900
/*implicit=*/false);
882901
} else if constexpr (isCombinedType<OpTy>) {
883902
applyToComputeOp(clause);
@@ -893,7 +912,7 @@ class OpenACCClauseCIREmitter final
893912
mlir::acc::KernelsOp>) {
894913
for (auto var : clause.getVarList())
895914
addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
896-
var, mlir::acc::DataClause::acc_attach, /*structured=*/true,
915+
var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
897916
/*implicit=*/false);
898917
} else if constexpr (isCombinedType<OpTy>) {
899918
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)