105 lines · cpp
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s2 3struct HasSideEffects {4 HasSideEffects();5 ~HasSideEffects();6};7 8struct Struct {9 static const HasSideEffects StaticMemHSE;10 static const HasSideEffects StaticMemHSEArr[5];11 static const int StaticMemInt;12 13 void MemFunc1(HasSideEffects *ArgHSE, int *ArgInt) {14 // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})15 // CHECK-NEXT: cir.alloca{{.*}}["this"16 // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE17 // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt18 // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE19 // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt20 // CHECK-NEXT: cir.store21 // CHECK-NEXT: cir.store22 // CHECK-NEXT: cir.store23 // CHECK-NEXT: cir.load24 HasSideEffects *LocalHSE;25 int *LocalInt;26#pragma acc declare deviceptr(ArgHSE, ArgInt, LocalHSE, LocalInt)27 // CHECK-NEXT: %[[DEV_PTR_ARG_HSE:.*]] = acc.deviceptr varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSE"}28 // CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}29 // CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}30 // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}31 // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]], %[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)32 33 // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])34 }35 void MemFunc2(HasSideEffects *ArgHSE, int *ArgInt);36};37 38void use() {39 Struct s;40 s.MemFunc1(nullptr, nullptr);41}42 43void Struct::MemFunc2(HasSideEffects *ArgHSE, int *ArgInt) {44 // CHECK: cir.func {{.*}}MemFunc2{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})45 // CHECK-NEXT: cir.alloca{{.*}}["this"46 // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE47 // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt48 // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE49 // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt50 // CHECK-NEXT: cir.store51 // CHECK-NEXT: cir.store52 // CHECK-NEXT: cir.store53 // CHECK-NEXT: cir.load54 HasSideEffects *LocalHSE;55 int *LocalInt;56#pragma acc declare deviceptr(ArgHSE, ArgInt)57 // CHECK-NEXT: %[[DEV_PTR_ARG_HSE:.*]] = acc.deviceptr varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSE"}58 // CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}59 // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)60 61#pragma acc declare deviceptr(LocalHSE, LocalInt)62 // CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}63 // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}64 // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)65 //66 // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])67 // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])68}69 70extern "C" void do_thing();71 72void NormalFunc(HasSideEffects *ArgHSE, int *ArgInt) {73 // CHECK: cir.func {{.*}}NormalFunc{{.*}}(%[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})74 // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE75 // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt76 // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE77 // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt78 // CHECK-NEXT: cir.store79 // CHECK-NEXT: cir.store80 HasSideEffects *LocalHSE;81 int *LocalInt;82#pragma acc declare deviceptr(ArgHSE, ArgInt)83 // CHECK-NEXT: %[[DEV_PTR_ARG_HSE:.*]] = acc.deviceptr varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSE"}84 // CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}85 // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)86 {87 // CHECK-NEXT: cir.scope {88#pragma acc declare deviceptr(LocalHSE, LocalInt)89 // CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}90 // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}91 // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)92 do_thing();93 // CHECK-NEXT: cir.call @do_thing94 // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])95 96 }97 // CHECK-NEXT: }98 99 // Make sure that cleanup gets put in the right scope.100 do_thing();101 // CHECK-NEXT: cir.call @do_thing102 // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])103}104 105