133 lines · c
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s2 3void acc_cache() {4 // CHECK: cir.func{{.*}} @acc_cache()5 6 int iArr[10];7 // CHECK-NEXT: %[[IARR:.*]] = cir.alloca !cir.array<!s32i x 10>, !cir.ptr<!cir.array<!s32i x 10>>, ["iArr"]8 float fArr[10];9 // CHECK-NEXT: %[[FARR:.*]] = cir.alloca !cir.array<!cir.float x 10>, !cir.ptr<!cir.array<!cir.float x 10>>, ["fArr"]10 11#pragma acc cache(iArr[1], fArr[1:5])12 // This does nothing, as it is not in a loop.13 14#pragma acc parallel15 {16#pragma acc cache(iArr[1], fArr[1:5])17 // This does nothing, as it is not in a loop.18 }19 // CHECK-NEXT: acc.parallel {20 // CHECK-NEXT: acc.yield21 // CHECK-NEXT: } loc22 23#pragma acc loop24 for(int i = 0; i < 5; ++i) {25 for(int j = 0; j < 5; ++j) {26#pragma acc cache(iArr[1], fArr[1:5])27 }28 }29 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>30 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3231 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6432 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6433 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6434 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)35 // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false}36 //37 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>38 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3239 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5>40 // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si3241 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6442 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6443 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64)44 // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!cir.float x 10>> {name = "fArr[1:5]", structured = false}45 //46 // CHECK-NEXT: acc.loop cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) {47 // CHECK: acc.yield48 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]}49 50#pragma acc loop51 for(int i = 0; i < 5; ++i) {52#pragma acc cache(iArr[1], fArr[1:5])53 }54 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>55 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3256 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6457 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6458 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6459 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)60 // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false}61 //62 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>63 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3264 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5>65 // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si3266 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6467 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6468 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64)69 // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!cir.float x 10>> {name = "fArr[1:5]", structured = false}70 //71 // CHECK-NEXT: acc.loop cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) {72 // CHECK: acc.yield73 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]}74 75#pragma acc parallel loop76 for(int i = 0; i < 5; ++i) {77#pragma acc cache(iArr[1], fArr[1:5])78 }79 // CHECK-NEXT: acc.parallel combined(loop) {80 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>81 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3282 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6483 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6484 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i6485 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)86 // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false}87 //88 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>89 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si3290 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5>91 // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si3292 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i6493 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i6494 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64)95 // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!cir.float x 10>> {name = "fArr[1:5]", structured = false}96 //97 // CHECK-NEXT: acc.loop combined(parallel) cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) {98 // CHECK: acc.yield99 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]}100 // CHECK-NEXT: acc.yield101 // CHECK-NEXT: } loc102 103#pragma acc parallel loop104 for(int i = 0; i < 5; ++i) {105 int localArr[5];106 // The first term here isn't lowered, because it references data inside of the 'loop'.107#pragma acc cache(localArr[i], iArr[1], fArr[1:5])108 }109 // CHECK-NEXT: acc.parallel combined(loop) {110 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>111 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32112 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64113 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64114 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64115 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64)116 // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false}117 //118 // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1>119 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32120 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5>121 // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si32122 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64123 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64124 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64)125 // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!cir.float x 10>> {name = "fArr[1:5]", structured = false}126 //127 // CHECK-NEXT: acc.loop combined(parallel) cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) {128 // CHECK: acc.yield129 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]}130 // CHECK-NEXT: acc.yield131 // CHECK-NEXT: } loc132}133