// RUN: not %clang_cc1 -fopenacc -triple x86_64-linux-gnu -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple x86_64-linux-pc %s -o - | FileCheck %s struct NoCopyConstruct {}; // int // CHECK: acc.private.recipe @privatization__ZTSi : !cir.ptr init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr {{.*}}): // CHECK-NEXT: cir.alloca !s32i, !cir.ptr, ["openacc.private.init"] // CHECK-NEXT: acc.yield // CHECK-NEXT: } // // float // CHECK-NEXT: acc.private.recipe @privatization__ZTSf : !cir.ptr init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr {{.*}}): // CHECK-NEXT: cir.alloca !cir.float, !cir.ptr, ["openacc.private.init"] // CHECK-NEXT: acc.yield // CHECK-NEXT: } // // NoCopyConstruct // CHECK-NEXT: acc.private.recipe @privatization__ZTS15NoCopyConstruct : !cir.ptr init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr {{.*}}): // CHECK-NEXT: cir.alloca !rec_NoCopyConstruct, !cir.ptr, ["openacc.private.init"] // CHECK-NEXT: acc.yield // CHECK-NEXT: } // // int[5] with 1 'bound' // CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_i : !cir.ptr> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } // // float[5] with 1 'bound' // CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_f : !cir.ptr> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } // // NoCopyConstruct[5] with 1 'bound' // CHECK-NEXT: acc.private.recipe @privatization__Bcnt1__ZTSA5_15NoCopyConstruct : !cir.ptr> init { // CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr> {{.*}}, %[[BOUND1:.*]]: !acc.data_bounds_ty {{.*}}): // TODO: Add Init here. // CHECK-NEXT: acc.yield // CHECK-NEXT: } void acc_compute() { // CHECK: cir.func{{.*}} @acc_compute() { int someInt; // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !s32i, !cir.ptr, ["someInt"] float someFloat; // CHECK-NEXT: %[[SOMEFLOAT:.*]] = cir.alloca !cir.float, !cir.ptr, ["someFloat"] struct NoCopyConstruct noCopy; // CHECK-NEXT: %[[NOCOPY:.*]] = cir.alloca !rec_NoCopyConstruct, !cir.ptr, ["noCopy"] int someIntArr[5]; // CHECK-NEXT: %[[INTARR:.*]] = cir.alloca !cir.array, !cir.ptr>, ["someIntArr"] float someFloatArr[5]; // CHECK-NEXT: %[[FLOATARR:.*]] = cir.alloca !cir.array, !cir.ptr>, ["someFloatArr"] struct NoCopyConstruct noCopyArr[5]; // CHECK-NEXT: %[[NOCOPYARR:.*]] = cir.alloca !cir.array, !cir.ptr>, ["noCopyArr"] #pragma acc parallel private(someInt) ; // CHECK: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr) -> !cir.ptr {name = "someInt"} // CHECK-NEXT: acc.parallel private(@privatization__ZTSi -> %[[PRIVATE]] : !cir.ptr) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc serial private(someFloat) ; // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr) -> !cir.ptr {name = "someFloat"} // CHECK-NEXT: acc.serial private(@privatization__ZTSf -> %[[PRIVATE]] : !cir.ptr) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc parallel private(noCopy) ; // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr) -> !cir.ptr {name = "noCopy"} // CHECK-NEXT: acc.parallel private(@privatization__ZTS15NoCopyConstruct -> %[[PRIVATE]] : !cir.ptr // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc parallel private(someInt, someFloat, noCopy) ; // CHECK: %[[PRIVATE1:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr) -> !cir.ptr {name = "someInt"} // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr) -> !cir.ptr {name = "someFloat"} // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr) -> !cir.ptr {name = "noCopy"} // CHECK-NEXT: acc.parallel private(@privatization__ZTSi -> %[[PRIVATE1]] : !cir.ptr, // CHECK-SAME: @privatization__ZTSf -> %[[PRIVATE2]] : !cir.ptr, // CHECK-SAME: @privatization__ZTS15NoCopyConstruct -> %[[PRIVATE3]] : !cir.ptr) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc serial private(someIntArr[1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someIntArr[1]"} // CHECK-NEXT: acc.serial private(@privatization__Bcnt1__ZTSA5_i -> %[[PRIVATE]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc parallel private(someFloatArr[1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someFloatArr[1]"} // CHECK-NEXT: acc.parallel private(@privatization__Bcnt1__ZTSA5_f -> %[[PRIVATE]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc serial private(noCopyArr[1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "noCopyArr[1]"} // CHECK-NEXT: acc.serial private(@privatization__Bcnt1__ZTSA5_15NoCopyConstruct -> %[[PRIVATE]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc serial private(someIntArr[1], someFloatArr[1], noCopyArr[1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someIntArr[1]"} // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someFloatArr[1]"} // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "noCopyArr[1]"} // CHECK-NEXT: acc.serial private(@privatization__Bcnt1__ZTSA5_i -> %[[PRIVATE1]] : !cir.ptr>, // CHECK-SAME: @privatization__Bcnt1__ZTSA5_f -> %[[PRIVATE2]] : !cir.ptr>, // CHECK-SAME: @privatization__Bcnt1__ZTSA5_15NoCopyConstruct -> %[[PRIVATE3]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc parallel private(someIntArr[1:1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someIntArr[1:1]"} // CHECK-NEXT: acc.parallel private(@privatization__Bcnt1__ZTSA5_i -> %[[PRIVATE]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc serial private(someFloatArr[1:1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someFloatArr[1:1]"} // CHECK-NEXT: acc.serial private(@privatization__Bcnt1__ZTSA5_f -> %[[PRIVATE]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc parallel private(noCopyArr[1:1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "noCopyArr[1:1]"} // CHECK-NEXT: acc.parallel private(@privatization__Bcnt1__ZTSA5_15NoCopyConstruct -> %[[PRIVATE]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc #pragma acc parallel private(someIntArr[1:1], someFloatArr[1:1], noCopyArr[1:1]) ; // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.private varPtr(%[[INTARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someIntArr[1:1]"} // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[FLOATARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "someFloatArr[1:1]"} // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[NOCOPYARR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "noCopyArr[1:1]"} // CHECK-NEXT: acc.parallel private(@privatization__Bcnt1__ZTSA5_i -> %[[PRIVATE1]] : !cir.ptr>, // CHECK-SAME: @privatization__Bcnt1__ZTSA5_f -> %[[PRIVATE2]] : !cir.ptr>, // CHECK-SAME: @privatization__Bcnt1__ZTSA5_15NoCopyConstruct -> %[[PRIVATE3]] : !cir.ptr>) // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc }