Skip to content

Commit 39bb267

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR][NFC] Add device_ptr async clause tests
Add a test to ensure that device_ptr properly respects the 'async' functionality we added for copy/etc.
1 parent 8f7e574 commit 39bb267

File tree

4 files changed

+74
-0
lines changed

4 files changed

+74
-0
lines changed

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

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1048,4 +1048,35 @@ extern "C" void acc_combined_deviceptr(int *arg1, int *arg2) {
10481048
// CHECK-NEXT: } loc
10491049
// CHECK-NEXT: acc.terminator
10501050
// CHECK-NEXT: } loc
1051+
1052+
#pragma acc parallel loop deviceptr(arg1) async
1053+
for(unsigned I = 0; I < 5; ++I);
1054+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
1055+
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) async {
1056+
// CHECK-NEXT: acc.loop combined(parallel) {
1057+
// CHECK: acc.yield
1058+
// CHECK-NEXT: } loc
1059+
// CHECK-NEXT: acc.yield
1060+
// CHECK-NEXT: } loc
1061+
1062+
#pragma acc serial loop deviceptr(arg2) async device_type(nvidia)
1063+
for(unsigned I = 0; I < 5; ++I);
1064+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
1065+
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>) async {
1066+
// CHECK-NEXT: acc.loop combined(serial) {
1067+
// CHECK: acc.yield
1068+
// CHECK-NEXT: } loc
1069+
// CHECK-NEXT: acc.yield
1070+
// CHECK-NEXT: } loc
1071+
1072+
#pragma acc kernels loop deviceptr(arg1, arg2) device_type(nvidia) async
1073+
for(unsigned I = 0; I < 5; ++I);
1074+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
1075+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
1076+
// CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {
1077+
// CHECK-NEXT: acc.loop combined(kernels) {
1078+
// CHECK: acc.yield
1079+
// CHECK-NEXT: } loc
1080+
// CHECK-NEXT: acc.terminator
1081+
// CHECK-NEXT: } loc
10511082
}

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

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -437,4 +437,18 @@ void acc_kernels_deviceptr(int *arg1, int *arg2) {
437437
// CHECK-NEXT: acc.kernels dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {
438438
// CHECK-NEXT: acc.terminator
439439
// CHECK-NEXT: } loc
440+
441+
#pragma acc kernels deviceptr(arg1) async
442+
;
443+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
444+
// CHECK-NEXT: acc.kernels dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) async {
445+
// CHECK-NEXT: acc.terminator
446+
// CHECK-NEXT: } loc
447+
#pragma acc kernels deviceptr(arg1, arg2) device_type(radeon) async
448+
;
449+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
450+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
451+
// CHECK-NEXT: acc.kernels dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {
452+
// CHECK-NEXT: acc.terminator
453+
// CHECK-NEXT: } loc
440454
}

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

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -464,4 +464,18 @@ void acc_parallel_deviceptr(int *arg1, int *arg2) {
464464
// CHECK-NEXT: acc.parallel dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {
465465
// CHECK-NEXT: acc.yield
466466
// CHECK-NEXT: } loc
467+
468+
#pragma acc parallel deviceptr(arg1) async
469+
;
470+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
471+
// CHECK-NEXT: acc.parallel dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) async {
472+
// CHECK-NEXT: acc.yield
473+
// CHECK-NEXT: } loc
474+
#pragma acc parallel deviceptr(arg1, arg2) device_type(radeon) async
475+
;
476+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
477+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
478+
// CHECK-NEXT: acc.parallel dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {
479+
// CHECK-NEXT: acc.yield
480+
// CHECK-NEXT: } loc
467481
}

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

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -287,4 +287,19 @@ void acc_serial_deviceptr(int *arg1, int *arg2) {
287287
// CHECK-NEXT: acc.serial dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {
288288
// CHECK-NEXT: acc.yield
289289
// CHECK-NEXT: } loc
290+
291+
#pragma acc serial deviceptr(arg1) async
292+
;
293+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
294+
// CHECK-NEXT: acc.serial dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) async {
295+
// CHECK-NEXT: acc.yield
296+
// CHECK-NEXT: } loc
297+
298+
#pragma acc serial deviceptr(arg1, arg2) device_type(radeon) async
299+
;
300+
// CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}
301+
// CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}
302+
// CHECK-NEXT: acc.serial dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {
303+
// CHECK-NEXT: acc.yield
304+
// CHECK-NEXT: } loc
290305
}

0 commit comments

Comments
 (0)