Skip to content

Commit bdaba7b

Browse files
erichkeanerorth
erichkeane
authored and
rorth
committed
[OpenACC][CIR] Lowering for 'deviceptr' for compute/combined constructs
This ends up being a simple clause that only adds 'acc.deviceptr' to the dataOperands list on the compute construct operation.
1 parent e21ca1a commit bdaba7b

File tree

5 files changed

+120
-0
lines changed

5 files changed

+120
-0
lines changed

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

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -829,6 +829,22 @@ class OpenACCClauseCIREmitter final
829829
llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
830830
}
831831
}
832+
833+
void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) {
834+
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
835+
mlir::acc::KernelsOp>) {
836+
for (auto var : clause.getVarList())
837+
addDataOperand<mlir::acc::DevicePtrOp>(
838+
var, mlir::acc::DataClause::acc_deviceptr, /*structured=*/true,
839+
/*implicit=*/false);
840+
} else if constexpr (isCombinedType<OpTy>) {
841+
applyToComputeOp(clause);
842+
} else {
843+
// TODO: When we've implemented this for everything, switch this to an
844+
// unreachable. data, declare remain.
845+
return clauseNotImplemented(clause);
846+
}
847+
}
832848
};
833849

834850
template <typename OpTy>

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

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1011,3 +1011,41 @@ extern "C" void acc_combined(int N, int cond) {
10111011
// CHECK-NEXT: acc.terminator
10121012
// CHECK-NEXT: } loc
10131013
}
1014+
extern "C" void acc_combined_deviceptr(int *arg1, int *arg2) {
1015+
// CHECK: cir.func @acc_combined_deviceptr(%[[ARG1_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}, %[[ARG2_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}) {
1016+
// CHECK-NEXT: %[[ARG1:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg1", init]
1017+
// CHECK-NEXT: %[[ARG2:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg2", init]
1018+
// CHECK-NEXT: cir.store %[[ARG1_PARAM]], %[[ARG1]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
1019+
// CHECK-NEXT: cir.store %[[ARG2_PARAM]], %[[ARG2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
1020+
1021+
#pragma acc parallel loop deviceptr(arg1)
1022+
for(unsigned I = 0; I < 5; ++I);
1023+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
1024+
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) {
1025+
// CHECK-NEXT: acc.loop combined(parallel) {
1026+
// CHECK: acc.yield
1027+
// CHECK-NEXT: } loc
1028+
// CHECK-NEXT: acc.yield
1029+
// CHECK-NEXT: } loc
1030+
1031+
#pragma acc serial loop deviceptr(arg2)
1032+
for(unsigned I = 0; I < 5; ++I);
1033+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
1034+
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>) {
1035+
// CHECK-NEXT: acc.loop combined(serial) {
1036+
// CHECK: acc.yield
1037+
// CHECK-NEXT: } loc
1038+
// CHECK-NEXT: acc.yield
1039+
// CHECK-NEXT: } loc
1040+
1041+
#pragma acc kernels loop deviceptr(arg1, arg2)
1042+
for(unsigned I = 0; I < 5; ++I);
1043+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
1044+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
1045+
// CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {
1046+
// CHECK-NEXT: acc.loop combined(kernels) {
1047+
// CHECK: acc.yield
1048+
// CHECK-NEXT: } loc
1049+
// CHECK-NEXT: acc.terminator
1050+
// CHECK-NEXT: } loc
1051+
}

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -416,3 +416,25 @@ void acc_kernels(int cond) {
416416

417417
// CHECK-NEXT: cir.return
418418
}
419+
420+
void acc_kernels_deviceptr(int *arg1, int *arg2) {
421+
// CHECK: cir.func @acc_kernels_deviceptr(%[[ARG1_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}, %[[ARG2_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}) {
422+
// CHECK-NEXT: %[[ARG1:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg1", init]
423+
// CHECK-NEXT: %[[ARG2:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg2", init]
424+
// CHECK-NEXT: cir.store %[[ARG1_PARAM]], %[[ARG1]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
425+
// CHECK-NEXT: cir.store %[[ARG2_PARAM]], %[[ARG2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
426+
427+
#pragma acc kernels deviceptr(arg1)
428+
;
429+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
430+
// CHECK-NEXT: acc.kernels dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) {
431+
// CHECK-NEXT: acc.terminator
432+
// CHECK-NEXT: } loc
433+
#pragma acc kernels deviceptr(arg1, arg2)
434+
;
435+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
436+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
437+
// CHECK-NEXT: acc.kernels dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {
438+
// CHECK-NEXT: acc.terminator
439+
// CHECK-NEXT: } loc
440+
}

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -443,3 +443,25 @@ void acc_parallel(int cond) {
443443

444444
// CHECK-NEXT: cir.return
445445
}
446+
447+
void acc_parallel_deviceptr(int *arg1, int *arg2) {
448+
// CHECK: cir.func @acc_parallel_deviceptr(%[[ARG1_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}, %[[ARG2_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}) {
449+
// CHECK-NEXT: %[[ARG1:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg1", init]
450+
// CHECK-NEXT: %[[ARG2:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg2", init]
451+
// CHECK-NEXT: cir.store %[[ARG1_PARAM]], %[[ARG1]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
452+
// CHECK-NEXT: cir.store %[[ARG2_PARAM]], %[[ARG2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
453+
454+
#pragma acc parallel deviceptr(arg1)
455+
;
456+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
457+
// CHECK-NEXT: acc.parallel dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) {
458+
// CHECK-NEXT: acc.yield
459+
// CHECK-NEXT: } loc
460+
#pragma acc parallel deviceptr(arg1, arg2)
461+
;
462+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
463+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
464+
// CHECK-NEXT: acc.parallel dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {
465+
// CHECK-NEXT: acc.yield
466+
// CHECK-NEXT: } loc
467+
}

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -266,3 +266,25 @@ void acc_serial(int cond) {
266266

267267
// CHECK-NEXT: cir.return
268268
}
269+
270+
void acc_serial_deviceptr(int *arg1, int *arg2) {
271+
// CHECK: cir.func @acc_serial_deviceptr(%[[ARG1_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}, %[[ARG2_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}) {
272+
// CHECK-NEXT: %[[ARG1:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg1", init]
273+
// CHECK-NEXT: %[[ARG2:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg2", init]
274+
// CHECK-NEXT: cir.store %[[ARG1_PARAM]], %[[ARG1]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
275+
// CHECK-NEXT: cir.store %[[ARG2_PARAM]], %[[ARG2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
276+
277+
#pragma acc serial deviceptr(arg1)
278+
;
279+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
280+
// CHECK-NEXT: acc.serial dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) {
281+
// CHECK-NEXT: acc.yield
282+
// CHECK-NEXT: } loc
283+
#pragma acc serial deviceptr(arg1, arg2)
284+
;
285+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
286+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
287+
// CHECK-NEXT: acc.serial dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {
288+
// CHECK-NEXT: acc.yield
289+
// CHECK-NEXT: } loc
290+
}

0 commit comments

Comments
 (0)