Skip to content

Commit 4eb7231

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] Implement 'present' lowering for compute/combined
'present' has a data-in operation acc.present, and uses 'delete' for its data out. Otherwise it is identical to no_create.
1 parent b09b1d6 commit 4eb7231

File tree

5 files changed

+104
-0
lines changed

5 files changed

+104
-0
lines changed

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

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -871,6 +871,22 @@ class OpenACCClauseCIREmitter final
871871
return clauseNotImplemented(clause);
872872
}
873873
}
874+
875+
void VisitPresentClause(const OpenACCPresentClause &clause) {
876+
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
877+
mlir::acc::KernelsOp>) {
878+
for (auto var : clause.getVarList())
879+
addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
880+
var, mlir::acc::DataClause::acc_present, /*structured=*/true,
881+
/*implicit=*/false);
882+
} else if constexpr (isCombinedType<OpTy>) {
883+
applyToComputeOp(clause);
884+
} else {
885+
// TODO: When we've implemented this for everything, switch this to an
886+
// unreachable. data & declare remain.
887+
return clauseNotImplemented(clause);
888+
}
889+
}
874890
};
875891

876892
template <typename OpTy>

‎clang/test/CIR/CodeGenOpenACC/combined.cpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1114,4 +1114,39 @@ extern "C" void acc_combined_data_clauses(int *arg1, int *arg2) {
11141114
// CHECK-NEXT: } loc
11151115
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg2"}
11161116
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg1"}
1117+
1118+
#pragma acc parallel loop present(arg1)
1119+
for(unsigned I = 0; I < 5; ++I);
1120+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
1121+
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {
1122+
// CHECK-NEXT: acc.loop combined(parallel) {
1123+
// CHECK: acc.yield
1124+
// CHECK-NEXT: } loc
1125+
// CHECK-NEXT: acc.yield
1126+
// CHECK-NEXT: } loc
1127+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
1128+
1129+
#pragma acc serial loop present(arg2)
1130+
for(unsigned I = 0; I < 5; ++I);
1131+
// CHECK-NEXT: %[[NOCREATE2:.*]] = acc.present varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
1132+
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) {
1133+
// CHECK-NEXT: acc.loop combined(serial) {
1134+
// CHECK: acc.yield
1135+
// CHECK-NEXT: } loc
1136+
// CHECK-NEXT: acc.yield
1137+
// CHECK-NEXT: } loc
1138+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_present>, name = "arg2"}
1139+
1140+
#pragma acc kernels loop present(arg1, arg2) device_type(host) async
1141+
for(unsigned I = 0; I < 5; ++I);
1142+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
1143+
// CHECK-NEXT: %[[NOCREATE2:.*]] = acc.present varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
1144+
// CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {
1145+
// CHECK-NEXT: acc.loop combined(kernels) {
1146+
// CHECK: acc.yield
1147+
// CHECK-NEXT: } loc
1148+
// CHECK-NEXT: acc.terminator
1149+
// CHECK-NEXT: } loc
1150+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_present>, name = "arg2"}
1151+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
11171152
}

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

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -469,4 +469,22 @@ void acc_kernels_data_clauses(int *arg1, int *arg2) {
469469
// CHECK-NEXT: } loc
470470
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg2"}
471471
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg1"}
472+
473+
#pragma acc kernels present(arg1)
474+
;
475+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
476+
// CHECK-NEXT: acc.kernels dataOperands(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {
477+
// CHECK-NEXT: acc.terminator
478+
// CHECK-NEXT: } loc
479+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
480+
481+
#pragma acc kernels present(arg1, arg2) device_type(nvidia) async
482+
;
483+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
484+
// CHECK-NEXT: %[[NOCREATE2:.*]] = acc.present varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
485+
// CHECK-NEXT: acc.kernels dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {
486+
// CHECK-NEXT: acc.terminator
487+
// CHECK-NEXT: } loc
488+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_present>, name = "arg2"}
489+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
472490
}

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

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -496,5 +496,23 @@ void acc_parallel_data_clauses(int *arg1, int *arg2) {
496496
// CHECK-NEXT: } loc
497497
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg2"}
498498
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg1"}
499+
500+
#pragma acc parallel present(arg1)
501+
;
502+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
503+
// CHECK-NEXT: acc.parallel dataOperands(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {
504+
// CHECK-NEXT: acc.yield
505+
// CHECK-NEXT: } loc
506+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
507+
508+
#pragma acc parallel present(arg1, arg2) device_type(radeon) async
509+
;
510+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
511+
// CHECK-NEXT: %[[NOCREATE2:.*]] = acc.present varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
512+
// CHECK-NEXT: acc.parallel dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {
513+
// CHECK-NEXT: acc.yield
514+
// CHECK-NEXT: } loc
515+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_present>, name = "arg2"}
516+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
499517
}
500518

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

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -319,4 +319,21 @@ void acc_serial_data_clauses(int *arg1, int *arg2) {
319319
// CHECK-NEXT: } loc
320320
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg2"}
321321
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg1"}
322+
323+
#pragma acc serial present(arg1)
324+
;
325+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
326+
// CHECK-NEXT: acc.serial dataOperands(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {
327+
// CHECK-NEXT: acc.yield
328+
// CHECK-NEXT: } loc
329+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
330+
#pragma acc serial present(arg1, arg2) device_type(nvidia) async
331+
;
332+
// CHECK-NEXT: %[[NOCREATE1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
333+
// CHECK-NEXT: %[[NOCREATE2:.*]] = acc.present varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
334+
// CHECK-NEXT: acc.serial dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {
335+
// CHECK-NEXT: acc.yield
336+
// CHECK-NEXT: } loc
337+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_present>, name = "arg2"}
338+
// CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {dataClause = #acc<data_clause acc_present>, name = "arg1"}
322339
}

0 commit comments

Comments
 (0)