llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: darkbuck <details> <summary>Changes</summary> …ormat. NFC Assisted-by: DeepSeek V4 Flash --- Patch is 955.09 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/217729.diff 39 Files Affected: - (modified) clang/test/CIR/CodeGenOpenACC/cache.c (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/combined-copy.c (+169-169) - (modified) clang/test/CIR/CodeGenOpenACC/combined-copy.cpp (+48-48) - (modified) clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c (+61-61) - (modified) clang/test/CIR/CodeGenOpenACC/combined-firstprivate-clause.cpp (+57-57) - (modified) clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp (+57-57) - (modified) clang/test/CIR/CodeGenOpenACC/combined.cpp (+141-141) - (modified) clang/test/CIR/CodeGenOpenACC/compute-copy.c (+162-162) - (modified) clang/test/CIR/CodeGenOpenACC/compute-copy.cpp (+48-48) - (modified) clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c (+50-50) - (modified) clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause-templates.cpp (+4-4) - (modified) clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.c (+18-18) - (modified) clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/compute-private-clause-templates.cpp (+4-4) - (modified) clang/test/CIR/CodeGenOpenACC/compute-private-clause.c (+18-18) - (modified) clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c (+74-74) - (modified) clang/test/CIR/CodeGenOpenACC/data.c (+32-32) - (modified) clang/test/CIR/CodeGenOpenACC/declare-copy.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp (+72-72) - (modified) clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/declare-create.cpp (+72-72) - (modified) clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp (+72-72) - (modified) clang/test/CIR/CodeGenOpenACC/declare-link.cpp (+19-19) - (modified) clang/test/CIR/CodeGenOpenACC/declare-present.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/enter-data.c (+20-20) - (modified) clang/test/CIR/CodeGenOpenACC/exit-data.c (+40-40) - (modified) clang/test/CIR/CodeGenOpenACC/host_data.c (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/init.c (+5-5) - (modified) clang/test/CIR/CodeGenOpenACC/kernels.c (+27-27) - (modified) clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp (+57-57) - (modified) clang/test/CIR/CodeGenOpenACC/loop.cpp (+53-53) - (modified) clang/test/CIR/CodeGenOpenACC/parallel.c (+27-27) - (modified) clang/test/CIR/CodeGenOpenACC/serial.c (+27-27) - (modified) clang/test/CIR/CodeGenOpenACC/set.c (+3-3) - (modified) clang/test/CIR/CodeGenOpenACC/shutdown.c (+5-5) - (modified) clang/test/CIR/CodeGenOpenACC/update.c (+61-61) - (modified) clang/test/CIR/IR/openacc.cir (+3-3) ``````````diff diff --git a/clang/test/CIR/CodeGenOpenACC/cache.c b/clang/test/CIR/CodeGenOpenACC/cache.c index c67f0dd4d2f3a..7d65c4deba9e2 100644 --- a/clang/test/CIR/CodeGenOpenACC/cache.c +++ b/clang/test/CIR/CodeGenOpenACC/cache.c @@ -32,7 +32,7 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) structured(false) name("iArr[1]") -> !cir.ptr<!cir.array<!s32i x 10>> // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 @@ -41,11 +41,11 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) structured(false) name("fArr[1:5]") -> !cir.ptr<!cir.array<!cir.float x 10>> // // CHECK-NEXT: acc.loop cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) { // CHECK: acc.yield - // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} + // CHECK-NEXT: } independent #pragma acc loop for(int i = 0; i < 5; ++i) { @@ -57,7 +57,7 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) structured(false) name("iArr[1]") -> !cir.ptr<!cir.array<!s32i x 10>> // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 @@ -66,11 +66,11 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) structured(false) name("fArr[1:5]") -> !cir.ptr<!cir.array<!cir.float x 10>> // // CHECK-NEXT: acc.loop cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) { // CHECK: acc.yield - // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} + // CHECK-NEXT: } independent #pragma acc parallel loop for(int i = 0; i < 5; ++i) { @@ -83,7 +83,7 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) structured(false) name("iArr[1]") -> !cir.ptr<!cir.array<!s32i x 10>> // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 @@ -92,11 +92,11 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) structured(false) name("fArr[1:5]") -> !cir.ptr<!cir.array<!cir.float x 10>> // // CHECK-NEXT: acc.loop combined(parallel) cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) { // CHECK: acc.yield - // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} + // CHECK-NEXT: } independent // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc @@ -113,7 +113,7 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) structured(false) name("iArr[1]") -> !cir.ptr<!cir.array<!s32i x 10>> // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 @@ -122,11 +122,11 @@ void acc_cache() { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) - // 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} + // CHECK-NEXT: %[[CACHE2:.*]] = acc.cache varPtr(%[[FARR]] : !cir.ptr<!cir.array<!cir.float x 10>>) bounds(%[[BOUNDS]]) structured(false) name("fArr[1:5]") -> !cir.ptr<!cir.array<!cir.float x 10>> // // CHECK-NEXT: acc.loop combined(parallel) cache(%[[CACHE1]], %[[CACHE2]] : !cir.ptr<!cir.array<!s32i x 10>>, !cir.ptr<!cir.array<!cir.float x 10>>) { // CHECK: acc.yield - // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} + // CHECK-NEXT: } independent // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc } diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c b/clang/test/CIR/CodeGenOpenACC/combined-copy.c index a9277667af2b4..7af7dbd3faf64 100644 --- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c @@ -18,192 +18,192 @@ void acc_compute(int parmVar) { #pragma acc parallel loop copy(localVar1, parmVar) copy(localVar2) copy(localVar3, parmVar) for(int i = 0; i < 5; ++i); - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc - // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: %[[COPYIN4:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: %[[COPYIN5:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") -> !cir.ptr<!s32i> loc + // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") -> !cir.ptr<!s32i> loc + // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") -> !cir.ptr<!s16i> loc + // CHECK-NEXT: %[[COPYIN4:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") -> !cir.ptr<!cir.float> loc + // CHECK-NEXT: %[[COPYIN5:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") -> !cir.ptr<!s32i> loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]], %[[COPYIN4]], %[[COPYIN5]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>, !cir.ptr<!s32i>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN5]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN4]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN5]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN4]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") loc #pragma acc serial loop copy(localVar1, parmVar) copy(localVar2) copy(localVar3, parmVar) for(int i = 0; i < 5; ++i); - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc - // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: %[[COPYIN4:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: %[[COPYIN5:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") -> !cir.ptr<!s32i> loc + // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") -> !cir.ptr<!s32i> loc + // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") -> !cir.ptr<!s16i> loc + // CHECK-NEXT: %[[COPYIN4:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") -> !cir.ptr<!cir.float> loc + // CHECK-NEXT: %[[COPYIN5:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") -> !cir.ptr<!s32i> loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]], %[[COPYIN4]], %[[COPYIN5]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>, !cir.ptr<!s32i>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN5]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN4]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN5]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN4]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") loc #pragma acc kernels loop copy(localVar1, parmVar) copy(localVar2) copy(localVar3, parmVar) for(int i = 0; i < 5; ++i); - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc - // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: %[[COPYIN4:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: %[[COPYIN5:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") -> !cir.ptr<!s32i> loc + // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") -> !cir.ptr<!s32i> loc + // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") -> !cir.ptr<!s16i> loc + // CHECK-NEXT: %[[COPYIN4:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") -> !cir.ptr<!cir.float> loc + // CHECK-NEXT: %[[COPYIN5:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") -> !cir.ptr<!s32i> loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]], %[[COPYIN4]], %[[COPYIN5]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>, !cir.ptr<!s32i>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN5]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN4]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN5]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN4]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") loc #pragma acc parallel loop copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3) for(int i = 0; i < 5; ++i); - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc - // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc - // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i> loc + // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/217729 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
