126 lines · c
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s2void acc_data(int parmVar, int *ptrParmVar) {3 // CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) {4 // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]5 // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]6 // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>7 // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>8 9#pragma acc enter data copyin(parmVar)10 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}11 // CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)12 13#pragma acc enter data copyin(readonly, alwaysin: parmVar)14 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}15 // CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)16 17#pragma acc enter data copyin(readonly, alwaysin: parmVar) async18 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}19 // CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)20 21#pragma acc enter data async copyin(readonly, alwaysin: parmVar)22 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}23 // CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)24 25#pragma acc enter data copyin(readonly, alwaysin: parmVar) async(parmVar)26 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]27 // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]28 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}29 // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)30 31#pragma acc enter data async(parmVar) copyin(readonly, alwaysin: parmVar)32 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]33 // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]34 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}35 // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)36 37#pragma acc enter data create(parmVar)38 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}39 // CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)40 41#pragma acc enter data create(zero: parmVar)42 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}43 // CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)44 45#pragma acc enter data create(zero: parmVar) async46 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}47 // CHECK-NEXT: acc.enter_data async dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)48 49#pragma acc enter data create(zero: parmVar) async(parmVar)50 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]51 // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]52 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}53 // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)54 55#pragma acc enter data attach(ptrParmVar)56 // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}57 // CHECK-NEXT: acc.enter_data dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)58 59#pragma acc enter data attach(ptrParmVar) async60 // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}61 // CHECK-NEXT: acc.enter_data async dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)62 63#pragma acc enter data attach(ptrParmVar) async(parmVar)64 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]65 // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]66 // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}67 // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)68 69#pragma acc enter data if (parmVar == 1) copyin(parmVar)70 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]71 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i72 // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])73 // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]74 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}75 // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)76 77#pragma acc enter data async if (parmVar == 1) copyin(parmVar)78 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]79 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i80 // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])81 // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]82 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {name = "parmVar", structured = false}83 // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)84 85#pragma acc enter data if (parmVar == 1) async(parmVar) copyin(parmVar)86 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]87 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i88 // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])89 // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]90 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]91 // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]92 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}93 // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)94 95#pragma acc enter data wait create(parmVar)96 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}97 // CHECK-NEXT: acc.enter_data wait dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)98 99#pragma acc enter data wait(1) create(parmVar)100 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i101 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]102 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}103 // CHECK-NEXT: acc.enter_data wait(%[[ONE_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)104 105#pragma acc enter data wait(parmVar, 1, 2) create(parmVar)106 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]107 // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]108 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i109 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]110 // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i111 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]112 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}113 // CHECK-NEXT: acc.enter_data wait(%[[PARM_CAST]], %[[ONE_CAST]], %[[TWO_CAST]] : si32, si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)114 115#pragma acc enter data wait(devnum: parmVar: 1, 2) create(parmVar)116 // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]117 // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]118 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i119 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]120 // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i121 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]122 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}123 // CHECK-NEXT: acc.enter_data wait_devnum(%[[PARM_CAST]] : si32) wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)124 125}126