161 lines · c
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s2 3void acc_combined(int parmVar) {4 // CHECK: cir.func{{.*}} @acc_combined(%[[ARG:.*]]: !s32i{{.*}}) {5 // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]6 7 int localVar1;8 // CHECK-NEXT: %[[LV1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["localVar1"]9 float localVar2;10 // CHECK-NEXT: %[[LV2:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["localVar2"]11 // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]]12#pragma acc parallel loop copyin(parmVar) copyout(localVar1) create(localVar2)13 for(int i = 0; i < 5; ++i);14 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}15 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "localVar1"}16 // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "localVar2"}17 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]], %[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {18 // CHECK-NEXT: acc.loop combined(parallel) {19 // CHECK: acc.yield20 // CHECK-NEXT: } loc21 // CHECK-NEXT: acc.yield22 // CHECK-NEXT: } loc23 // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_create>, name = "localVar2"}24 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {name = "localVar1"}25 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "parmVar"}26 27#pragma acc serial loop copyin(parmVar, localVar1)28 for(int i = 0; i < 5; ++i);29 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}30 // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}31 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {32 // CHECK-NEXT: acc.loop combined(serial) {33 // CHECK: acc.yield34 // CHECK-NEXT: } loc35 // CHECK-NEXT: acc.yield36 // CHECK-NEXT: } loc37 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "localVar1"}38 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "parmVar"}39 40#pragma acc kernels loop copyout(parmVar, localVar1)41 for(int i = 0; i < 5; ++i);42 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar"}43 // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "localVar1"}44 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {45 // CHECK-NEXT: acc.loop combined(kernels) {46 // CHECK: acc.yield47 // CHECK-NEXT: } loc48 // CHECK-NEXT: acc.terminator49 // CHECK-NEXT: } loc50 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {name = "localVar1"}51 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar"}52 53#pragma acc parallel loop create (parmVar, localVar2)54 for(int i = 0; i < 5; ++i);55 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}56 // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "localVar2"}57 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {58 // CHECK-NEXT: acc.loop combined(parallel) {59 // CHECK: acc.yield60 // CHECK-NEXT: } loc61 // CHECK-NEXT: acc.yield62 // CHECK-NEXT: } loc63 // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_create>, name = "localVar2"}64 // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "parmVar"}65 66#pragma acc serial loop copyin(capture: parmVar) copyin(always: localVar1)67 for(int i = 0; i < 5; ++i);68 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}69 // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier always>, name = "localVar1"}70 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {71 // CHECK-NEXT: acc.loop combined(serial) {72 // CHECK: acc.yield73 // CHECK-NEXT: } loc74 // CHECK-NEXT: acc.yield75 // CHECK-NEXT: } loc76 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"}77 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}78 79#pragma acc kernels loop copyout(capture: parmVar) copyout(always: localVar1)80 for(int i = 0; i < 5; ++i);81 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}82 // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"}83 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {84 // CHECK-NEXT: acc.loop combined(kernels) {85 // CHECK: acc.yield86 // CHECK-NEXT: } loc87 // CHECK-NEXT: acc.terminator88 // CHECK-NEXT: } loc89 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "localVar1"}90 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}91 92#pragma acc parallel loop create(capture: parmVar)93 for(int i = 0; i < 5; ++i);94 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}95 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {96 // CHECK-NEXT: acc.loop combined(parallel) {97 // CHECK: acc.yield98 // CHECK-NEXT: } loc99 // CHECK-NEXT: acc.yield100 // CHECK-NEXT: } loc101 // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}102 103#pragma acc serial loop copyin(capture, always: parmVar, localVar1)104 for(int i = 0; i < 5; ++i);105 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier always,capture>, name = "parmVar"}106 // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier always,capture>, name = "localVar1"}107 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {108 // CHECK-NEXT: acc.loop combined(serial) {109 // CHECK: acc.yield110 // CHECK-NEXT: } loc111 // CHECK-NEXT: acc.yield112 // CHECK-NEXT: } loc113 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always,capture>, name = "localVar1"}114 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always,capture>, name = "parmVar"}115 116#pragma acc kernels loop copyin(readonly, always, alwaysin, capture: parmVar, localVar1, localVar2)117 for(int i = 0; i < 5; ++i);118 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "parmVar"}119 // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "localVar1"}120 // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LV2]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "localVar2"}121 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {122 // CHECK-NEXT: acc.loop combined(kernels) {123 // CHECK: acc.yield124 // CHECK-NEXT: } loc125 // CHECK-NEXT: acc.terminator126 // CHECK-NEXT: } loc127 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "localVar2"}128 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "localVar1"}129 // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "parmVar"}130 131#pragma acc parallel loop copyout(zero, always, alwaysout, capture: parmVar, localVar1, localVar2)132 for(int i = 0; i < 5; ++i);133 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name = "parmVar"}134 // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name = "localVar1"}135 // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name = "localVar2"}136 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]], %[[CREATE3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {137 // CHECK-NEXT: acc.loop combined(parallel) {138 // CHECK: acc.yield139 // CHECK-NEXT: } loc140 // CHECK-NEXT: acc.yield141 // CHECK-NEXT: } loc142 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) to varPtr(%[[LV2]] : !cir.ptr<!cir.float>) {modifiers = #acc<data_clause_modifier always,zero,capture>, name = "localVar2"}143 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always,zero,capture>, name = "localVar1"}144 // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always,zero,capture>, name = "parmVar"}145 146#pragma acc serial loop create(zero, capture: parmVar, localVar1, localVar2)147 for(int i = 0; i < 5; ++i);148 // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero,capture>, name = "parmVar"}149 // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero,capture>, name = "localVar1"}150 // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {modifiers = #acc<data_clause_modifier zero,capture>, name = "localVar2"}151 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]], %[[CREATE3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {152 // CHECK-NEXT: acc.loop combined(serial) {153 // CHECK: acc.yield154 // CHECK-NEXT: } loc155 // CHECK-NEXT: acc.yield156 // CHECK-NEXT: } loc157 // CHECK-NEXT: acc.delete accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero,capture>, name = "localVar2"}158 // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero,capture>, name = "localVar1"}159 // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero,capture>, name = "parmVar"}160}161