[clang] [cir] Fix tests post 3e3927f (Commit) / #217589 (PR) (PR #218925)
Andrzej Warzyński via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 26 06:58:35 PDT 2026
https://github.com/banach-space created https://github.com/llvm/llvm-project/pull/218925
None
>From aa00ef9e581f91bb31fe348a9fdb9c2c2f41ef20 Mon Sep 17 00:00:00 2001
From: Andrzej Warzynski <andrzej.warzynski at arm.com>
Date: Wed, 26 Aug 2026 13:45:43 +0000
Subject: [PATCH] [cir] Fix tests post 3e3927f (Commit) / #217589 (PR)
---
clang/test/CIR/CodeGenOpenACC/combined-copy.c | 36 ++++-----
.../combined-copyin-copyout-create.c | 64 +++++++--------
clang/test/CIR/CodeGenOpenACC/compute-copy.c | 36 ++++-----
.../compute-copyin-copyout-create.c | 64 +++++++--------
.../data-copy-copyin-copyout-create.c | 48 +++++------
.../test/CIR/CodeGenOpenACC/declare-copy.cpp | 61 +++++++-------
.../CIR/CodeGenOpenACC/declare-copyin.cpp | 79 +++++++++----------
.../CIR/CodeGenOpenACC/declare-copyout.cpp | 61 +++++++-------
.../CIR/CodeGenOpenACC/declare-create.cpp | 78 +++++++++---------
clang/test/CIR/CodeGenOpenACC/enter-data.c | 16 ++--
clang/test/CIR/CodeGenOpenACC/exit-data.c | 20 ++---
11 files changed, 280 insertions(+), 283 deletions(-)
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c b/clang/test/CIR/CodeGenOpenACC/combined-copy.c
index be07aa96fd1be..a3037349bbd09 100644
--- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c
+++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c
@@ -75,29 +75,29 @@ void acc_compute(int parmVar) {
#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>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i> loc
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!s16i> loc
- // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.float> loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = alwaysin> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <modifiers = alwaysout> -> !cir.ptr<!s16i> loc
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <modifiers = "always"> -> !cir.ptr<!cir.float> loc
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.loop combined(parallel) {
// CHECK: acc.yield
// CHECK-NEXT: }
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <{modifiers = #acc<data_clause_modifier always>}> loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <{modifiers = #acc<data_clause_modifier alwaysout>}> loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier alwaysin>}> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <modifiers = "always"> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <modifiers = alwaysout> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = alwaysin> loc
#pragma acc serial loop copy(always, alwaysin, alwaysout: localVar1)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = "always"> -> !cir.ptr<!s32i> loc
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !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(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = "always"> loc
short *localPointer;
float localArray[5];
@@ -1107,52 +1107,52 @@ void modifier_list() {
#pragma acc parallel loop copy(always:localVar)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(parallel) {
// CHECK: acc.yield
// CHECK-NEXT: } independent loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always">
#pragma acc serial loop copy(alwaysin:localVar)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysin> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(serial) {
// CHECK: acc.yield
// CHECK-NEXT: } seq loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysin>
#pragma acc kernels loop copy(alwaysout:localVar)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysout> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(kernels) {
// CHECK: acc.yield
// CHECK-NEXT: } auto_ loc
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysout>
#pragma acc parallel loop copy(capture:localVar)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = capture> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(parallel) {
// CHECK: acc.yield
// CHECK-NEXT: } independent loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = capture>
#pragma acc serial loop copy(capture, always, alwaysin, alwaysout:localVar)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always,capture"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(serial) {
// CHECK: acc.yield
// CHECK-NEXT: } seq loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always,capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always,capture">
}
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c b/clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c
index d219414e1601a..2f75c9f0a6670 100644
--- a/clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c
+++ b/clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c
@@ -65,96 +65,96 @@ void acc_combined(int parmVar) {
#pragma acc serial loop copyin(capture: parmVar) copyin(always: localVar1)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(serial) {
// CHECK: acc.yield
// CHECK-NEXT: } seq loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = capture>
#pragma acc kernels loop copyout(capture: parmVar) copyout(always: localVar1)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(kernels) {
// CHECK: acc.yield
// CHECK-NEXT: } auto_ loc
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture>
#pragma acc parallel loop create(capture: parmVar)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(parallel) {
// CHECK: acc.yield
// CHECK-NEXT: } independent loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <modifiers = capture>
#pragma acc serial loop copyin(capture, always: parmVar, localVar1)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always,capture"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.loop combined(serial) {
// CHECK: acc.yield
// CHECK-NEXT: } seq loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <{modifiers = #acc<data_clause_modifier always,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier always,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <modifiers = "always,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = "always,capture">
#pragma acc kernels loop copyin(readonly, always, alwaysin, capture: parmVar, localVar1, localVar2)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}> -> !cir.ptr<!cir.float>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,readonly,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always,readonly,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <modifiers = "always,readonly,capture"> -> !cir.ptr<!cir.float>
// CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.loop combined(kernels) {
// CHECK: acc.yield
// CHECK-NEXT: } auto_ loc
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) dataClause(acc_copyin) name("localVar2") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) dataClause(acc_copyin) name("localVar2") <modifiers = "always,readonly,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <modifiers = "always,readonly,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = "always,readonly,capture">
#pragma acc parallel loop copyout(zero, always, alwaysout, capture: parmVar, localVar1, localVar2)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <{modifiers = #acc<data_clause_modifier always,zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <{modifiers = #acc<data_clause_modifier always,zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) dataClause(acc_copyout) name("localVar2") <{modifiers = #acc<data_clause_modifier always,zero,capture>}> -> !cir.ptr<!cir.float>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <modifiers = "always,zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <modifiers = "always,zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) dataClause(acc_copyout) name("localVar2") <modifiers = "always,zero,capture"> -> !cir.ptr<!cir.float>
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]], %[[CREATE3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.loop combined(parallel) {
// CHECK: acc.yield
// CHECK-NEXT: } independent loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) to varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <{modifiers = #acc<data_clause_modifier always,zero,capture>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always,zero,capture>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,zero,capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) to varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <modifiers = "always,zero,capture">
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always,zero,capture">
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,zero,capture">
#pragma acc serial loop create(zero, capture: parmVar, localVar1, localVar2)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <{modifiers = #acc<data_clause_modifier zero,capture>}> -> !cir.ptr<!cir.float>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <modifiers = "zero,capture"> -> !cir.ptr<!cir.float>
// CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[CREATE1]], %[[CREATE2]], %[[CREATE3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.loop combined(serial) {
// CHECK: acc.yield
// CHECK-NEXT: } seq loc
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) dataClause(acc_create) name("localVar2") <{modifiers = #acc<data_clause_modifier zero,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) dataClause(acc_create) name("localVar1") <{modifiers = #acc<data_clause_modifier zero,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) dataClause(acc_create) name("localVar2") <modifiers = "zero,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) dataClause(acc_create) name("localVar1") <modifiers = "zero,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <modifiers = "zero,capture">
}
diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copy.c b/clang/test/CIR/CodeGenOpenACC/compute-copy.c
index f1f99e5d7805e..8771b0c8b692c 100644
--- a/clang/test/CIR/CodeGenOpenACC/compute-copy.c
+++ b/clang/test/CIR/CodeGenOpenACC/compute-copy.c
@@ -67,23 +67,23 @@ void acc_compute(int parmVar) {
#pragma acc parallel copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3)
;
- // 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]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!s16i> loc
- // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.float> loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = alwaysin> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <modifiers = alwaysout> -> !cir.ptr<!s16i> loc
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <modifiers = "always"> -> !cir.ptr<!cir.float> loc
// CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <{modifiers = #acc<data_clause_modifier always>}> loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <{modifiers = #acc<data_clause_modifier alwaysout>}> loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier alwaysin>}> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) dataClause(acc_copy) name("localVar3") <modifiers = "always"> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) dataClause(acc_copy) name("localVar2") <modifiers = alwaysout> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = alwaysin> loc
#pragma acc serial copy(always, alwaysin, alwaysout: localVar1)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = "always"> -> !cir.ptr<!s32i> loc
// CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar1") <modifiers = "always"> loc
short *localPointer;
float localArray[5];
@@ -903,38 +903,38 @@ void modifier_list() {
#pragma acc parallel copy(always:localVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always">
#pragma acc serial copy(alwaysin:localVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysin> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysin>
#pragma acc kernels copy(alwaysout:localVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysout> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = alwaysout>
#pragma acc parallel copy(capture:localVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = capture> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = capture>
#pragma acc serial copy(capture, always, alwaysin, alwaysout:localVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always,capture"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <{modifiers = #acc<data_clause_modifier always,capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("localVar") <modifiers = "always,capture">
}
diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c b/clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c
index 459fbf540aff3..59fdc7a92a1ff 100644
--- a/clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c
+++ b/clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c
@@ -54,75 +54,75 @@ void acc_compute(int parmVar) {
#pragma acc serial copyin(capture: parmVar) copyin(always: localVar1)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always"> -> !cir.ptr<!s32i> loc
// CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = capture>
#pragma acc kernels copyout(capture: parmVar) copyout(always: localVar1)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.kernels dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture>
#pragma acc parallel create(capture: parmVar)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.parallel dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <modifiers = capture>
#pragma acc serial copyin(capture, always: parmVar, localVar1)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,capture"> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always,capture"> -> !cir.ptr<!s32i> loc
// CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <{modifiers = #acc<data_clause_modifier always,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier always,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <modifiers = "always,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = "always,capture">
#pragma acc kernels copyin(readonly, always, alwaysin, capture: parmVar, localVar1, localVar2)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}> -> !cir.ptr<!cir.float>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,readonly,capture"> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always,readonly,capture"> -> !cir.ptr<!s32i> loc
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <modifiers = "always,readonly,capture"> -> !cir.ptr<!cir.float> loc
// CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) dataClause(acc_copyin) name("localVar2") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) dataClause(acc_copyin) name("localVar2") <modifiers = "always,readonly,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("localVar1") <modifiers = "always,readonly,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = "always,readonly,capture">
#pragma acc parallel copyout(zero, always, alwaysout, capture: parmVar, localVar1, localVar2)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <{modifiers = #acc<data_clause_modifier always,zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <{modifiers = #acc<data_clause_modifier always,zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) dataClause(acc_copyout) name("localVar2") <{modifiers = #acc<data_clause_modifier always,zero,capture>}> -> !cir.ptr<!cir.float>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <modifiers = "always,zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("localVar1") <modifiers = "always,zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) dataClause(acc_copyout) name("localVar2") <modifiers = "always,zero,capture"> -> !cir.ptr<!cir.float>
// CHECK-NEXT: acc.parallel dataOperands(%[[CREATE1]], %[[CREATE2]], %[[CREATE3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) to varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <{modifiers = #acc<data_clause_modifier always,zero,capture>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier always,zero,capture>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,zero,capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) to varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <modifiers = "always,zero,capture">
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "always,zero,capture">
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,zero,capture">
#pragma acc serial create(zero, capture: parmVar, localVar1, localVar2)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <{modifiers = #acc<data_clause_modifier zero,capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <{modifiers = #acc<data_clause_modifier zero,capture>}> -> !cir.ptr<!cir.float>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) name("localVar1") <modifiers = "zero,capture"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] : !cir.ptr<!cir.float>) name("localVar2") <modifiers = "zero,capture"> -> !cir.ptr<!cir.float>
// CHECK-NEXT: acc.serial dataOperands(%[[CREATE1]], %[[CREATE2]], %[[CREATE3]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) dataClause(acc_create) name("localVar2") <{modifiers = #acc<data_clause_modifier zero,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) dataClause(acc_create) name("localVar1") <{modifiers = #acc<data_clause_modifier zero,capture>}>
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>) dataClause(acc_create) name("localVar2") <modifiers = "zero,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) dataClause(acc_create) name("localVar1") <modifiers = "zero,capture">
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <modifiers = "zero,capture">
}
diff --git a/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c b/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c
index 7e9d03e3ed603..62811fad64b7a 100644
--- a/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c
+++ b/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c
@@ -35,21 +35,21 @@ void acc_data(int parmVar) {
#pragma acc data copy(capture, always, alwaysin, alwaysout: parmVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <{modifiers = #acc<data_clause_modifier always,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <modifiers = "always,capture"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <{modifiers = #acc<data_clause_modifier always,capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <modifiers = "always,capture">
#pragma acc data copy(capture: parmVar) copy(always, alwaysin, alwaysout: parmVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("parmVar") <modifiers = capture>
#pragma acc data copyin(parmVar)
;
@@ -79,21 +79,21 @@ void acc_data(int parmVar) {
#pragma acc data copyin(capture, always, alwaysin, readonly: parmVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,readonly,capture"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = "always,readonly,capture">
#pragma acc data copyin(capture: parmVar) copyin(readonly, always, alwaysin: parmVar)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always,readonly"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier always,readonly>}>
- // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = "always,readonly">
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("parmVar") <modifiers = capture>
#pragma acc data copyout(parmVar)
;
@@ -125,21 +125,21 @@ void acc_data(int parmVar) {
#pragma acc data copyout(capture, zero, alwaysout: parmVar)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <modifiers = "zero,alwaysout,capture"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout,capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "zero,alwaysout,capture">
#pragma acc data copyout(capture: parmVar) copyout(always, alwaysout: parmVar)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("parmVar") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture>
#pragma acc data create(parmVar)
;
@@ -172,21 +172,21 @@ void acc_data(int parmVar) {
#pragma acc data create(capture, zero: parmVar)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,capture>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = "zero,capture"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <modifiers = "zero,capture">
#pragma acc data create(capture: parmVar) create(zero: parmVar)
;
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = capture> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) name("parmVar") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } loc
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) dataClause(acc_create) name("parmVar") <modifiers = capture>
#pragma acc data no_create(parmVar)
;
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp
index 6acf95a42f43a..2af67028e829d 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp
@@ -33,10 +33,10 @@ struct Struct {
int LocalInt;
#pragma acc declare copy(always:ArgHSE, ArgInt, LocalHSE, LocalInt, ArgHSEPtr[1:1], LocalHSEArr[1:1])
- // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <modifiers = "always"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -44,7 +44,7 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -52,19 +52,19 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copy) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copy) name("LocalHSEArr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
//
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_copy) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_copy) name("LocalHSEArr[1:1]") <modifiers = "always">
// CHECK-NEXT: cir.yield
// CHRCK-NEXT: }
}
@@ -98,8 +98,8 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK: }
int LocalInt;
#pragma acc declare copy(alwaysin:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <modifiers = alwaysin> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <modifiers = alwaysin> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -107,13 +107,13 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <modifiers = alwaysin> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
#pragma acc declare copy(alwaysout:LocalHSE, LocalInt, LocalHSEArr[1:1])
// CHECK-NEXT: cir.cleanup.scope {
- // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <modifiers = alwaysout> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <modifiers = alwaysout> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -121,16 +121,16 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copy) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copy) name("LocalHSEArr[1:1]") <modifiers = alwaysout> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <{modifiers = #acc<data_clause_modifier alwaysout>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <{modifiers = #acc<data_clause_modifier alwaysout>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_copy) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("LocalHSE") <modifiers = alwaysout>
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("LocalInt") <modifiers = alwaysout>
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_copy) name("LocalHSEArr[1:1]") <modifiers = alwaysout>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
@@ -138,9 +138,9 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <{modifiers = #acc<data_clause_modifier alwaysin>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <{modifiers = #acc<data_clause_modifier alwaysin>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <modifiers = alwaysin>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <modifiers = alwaysin>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <modifiers = alwaysin>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
@@ -166,8 +166,8 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK: }
int LocalInt;
#pragma acc declare copy(capture:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <modifiers = capture> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <modifiers = capture> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -175,7 +175,7 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier capture>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <modifiers = capture> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
{
@@ -214,9 +214,8 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: cir.yield
// CHECK-NEXT: cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <{modifiers = #acc<data_clause_modifier capture>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <{modifiers = #acc<data_clause_modifier capture>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier capture>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copy) name("ArgHSE") <modifiers = capture>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("ArgInt") <modifiers = capture>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) dataClause(acc_copy) name("ArgHSEPtr[1:1]") <modifiers = capture>
// CHECK-NEXT: cir.yield
}
-
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
index a3b382569d55e..b5c8bc90e56d7 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
@@ -77,29 +77,29 @@ int NSInt1;
#pragma acc declare copyin(alwaysin: NSHSE1, NSInt1, NSHSEArr[1:1])
// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
-// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) name("NSHSE1") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) name("NSHSE1") <modifiers = alwaysin> -> !cir.ptr<!rec_HasSideEffects>
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!rec_HasSideEffects>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
// CHECK: acc.global_dtor @{{.*}}NSHSE1{{.*}}_acc_dtor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("NSHSE1") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("NSHSE1") <modifiers = alwaysin> -> !cir.ptr<!rec_HasSideEffects>
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("NSHSE1") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("NSHSE1") <modifiers = alwaysin>
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
-// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) name("NSInt1") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) name("NSInt1") <modifiers = alwaysin> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!s32i>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
// CHECK: acc.global_dtor @{{.*}}NSInt1{{.*}}_acc_dtor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("NSInt1") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("NSInt1") <modifiers = alwaysin> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("NSInt1") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("NSInt1") <modifiers = alwaysin>
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
@@ -112,7 +112,7 @@ int NSInt1;
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
-// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) name("NSHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) name("NSHSEArr[1:1]") <modifiers = alwaysin> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
@@ -125,9 +125,9 @@ int NSInt1;
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) dataClause(acc_copyin) name("NSHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) dataClause(acc_copyin) name("NSHSEArr[1:1]") <modifiers = alwaysin> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_copyin) name("NSHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_copyin) name("NSHSEArr[1:1]") <modifiers = alwaysin>
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
@@ -281,10 +281,10 @@ struct Struct {
int LocalInt;
#pragma acc declare copyin(always:ArgHSE, ArgInt, LocalHSE, LocalInt, ArgHSEPtr[1:1], LocalHSEArr[1:1])
- // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = "always"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -292,7 +292,7 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -300,19 +300,19 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
//
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("LocalHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("LocalInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyin) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyin) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("ArgHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("ArgInt") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("LocalHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("LocalInt") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyin) name("ArgHSEPtr[1:1]") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyin) name("LocalHSEArr[1:1]") <modifiers = "always">
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
@@ -346,8 +346,8 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK: }
int LocalInt;
#pragma acc declare copyin(alwaysin:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = alwaysin> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = alwaysin> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -355,13 +355,13 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <modifiers = alwaysin> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
#pragma acc declare copyin(alwaysin:LocalHSE, LocalInt, LocalHSEArr[1:1])
// CHECK-NEXT: cir.cleanup.scope {
- // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <modifiers = alwaysin> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <modifiers = alwaysin> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -369,16 +369,16 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <modifiers = alwaysin> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("LocalHSE") <{modifiers = #acc<data_clause_modifier alwaysin>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("LocalInt") <{modifiers = #acc<data_clause_modifier alwaysin>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyin) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("LocalHSE") <modifiers = alwaysin>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("LocalInt") <modifiers = alwaysin>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyin) name("LocalHSEArr[1:1]") <modifiers = alwaysin>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
@@ -386,9 +386,9 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("ArgHSE") <{modifiers = #acc<data_clause_modifier alwaysin>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("ArgInt") <{modifiers = #acc<data_clause_modifier alwaysin>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyin) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysin>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("ArgHSE") <modifiers = alwaysin>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("ArgInt") <modifiers = alwaysin>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyin) name("ArgHSEPtr[1:1]") <modifiers = alwaysin>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
@@ -414,8 +414,8 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK: }
int LocalInt;
#pragma acc declare copyin(always:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -423,7 +423,7 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
{
// CHECK-NEXT: cir.cleanup.scope {
@@ -461,8 +461,7 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: cir.yield
// CHECK-NEXT: cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyin) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyin) name("ArgHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr<!s32i>) dataClause(acc_copyin) name("ArgInt") <modifiers = "always">
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyin) name("ArgHSEPtr[1:1]") <modifiers = "always">
}
-
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp
index 47a97cf947f0d..09399cb65dc48 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp
@@ -33,10 +33,10 @@ struct Struct {
int LocalInt;
#pragma acc declare copyout(always:ArgHSE, ArgInt, LocalHSE, LocalInt, ArgHSEPtr[1:1], LocalHSEArr[1:1])
- // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("LocalHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("LocalInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("ArgHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("ArgInt") <modifiers = "always"> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("LocalHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("LocalInt") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -44,7 +44,7 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyout) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyout) name("ArgHSEPtr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -52,19 +52,19 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyout) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyout) name("LocalHSEArr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) name("ArgHSEPtr[1:1]") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) name("LocalHSEArr[1:1]") <modifiers = "always">
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
@@ -98,8 +98,8 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK: }
int LocalInt;
#pragma acc declare copyout(alwaysout:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("ArgHSE") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("ArgInt") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("ArgHSE") <modifiers = alwaysout> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("ArgInt") <modifiers = alwaysout> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -107,13 +107,13 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyout) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyout) name("ArgHSEPtr[1:1]") <modifiers = alwaysout> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
#pragma acc declare copyout(alwaysout:LocalHSE, LocalInt, LocalHSEArr[1:1])
// CHECK-NEXT: cir.cleanup.scope {
- // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("LocalHSE") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("LocalInt") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("LocalHSE") <modifiers = alwaysout> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("LocalInt") <modifiers = alwaysout> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -121,16 +121,16 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyout) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysout>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_copyout) name("LocalHSEArr[1:1]") <modifiers = alwaysout> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <{modifiers = #acc<data_clause_modifier alwaysout>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <{modifiers = #acc<data_clause_modifier alwaysout>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <modifiers = alwaysout>
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <modifiers = alwaysout>
+ // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) name("LocalHSEArr[1:1]") <modifiers = alwaysout>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
@@ -138,9 +138,9 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier alwaysout>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier alwaysout>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = alwaysout>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = alwaysout>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) name("ArgHSEPtr[1:1]") <modifiers = alwaysout>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
@@ -166,8 +166,8 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK: }
int LocalInt;
#pragma acc declare copyout(always:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_copyout) name("ArgHSE") <modifiers = "always"> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) dataClause(acc_copyout) name("ArgInt") <modifiers = "always"> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -175,7 +175,7 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyout) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_copyout) name("ArgHSEPtr[1:1]") <modifiers = "always"> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
{
@@ -216,8 +216,7 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier always>}>
- // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier always>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = "always">
+ // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) name("ArgHSEPtr[1:1]") <modifiers = "always">
}
-
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
index 4dc32b59c8c82..d0a6ace2fe4bf 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
@@ -77,29 +77,29 @@ int NSInt1;
#pragma acc declare create(zero: NSHSE1, NSInt1, NSHSEArr[1:1])
// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
-// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) name("NSHSE1") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) name("NSHSE1") <modifiers = zero> -> !cir.ptr<!rec_HasSideEffects>
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
// CHECK: acc.global_dtor @{{.*}}NSHSE1{{.*}}_acc_dtor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("NSHSE1") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("NSHSE1") <modifiers = zero> -> !cir.ptr<!rec_HasSideEffects>
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("NSHSE1") <{modifiers = #acc<data_clause_modifier zero>}>
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("NSHSE1") <modifiers = zero>
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
-// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) name("NSInt1") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) name("NSInt1") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
// CHECK: acc.global_dtor @{{.*}}NSInt1{{.*}}_acc_dtor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) dataClause(acc_create) name("NSInt1") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) dataClause(acc_create) name("NSInt1") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) dataClause(acc_create) name("NSInt1") <{modifiers = #acc<data_clause_modifier zero>}>
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) dataClause(acc_create) name("NSInt1") <modifiers = zero>
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
@@ -112,7 +112,7 @@ int NSInt1;
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
-// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) name("NSHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) name("NSHSEArr[1:1]") <modifiers = zero> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
@@ -125,9 +125,9 @@ int NSInt1;
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) dataClause(acc_create) name("NSHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) dataClause(acc_create) name("NSHSEArr[1:1]") <modifiers = zero> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_create) name("NSHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}>
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) dataClause(acc_create) name("NSHSEArr[1:1]") <modifiers = zero>
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
@@ -282,10 +282,10 @@ struct Struct {
int LocalInt;
#pragma acc declare create(zero:ArgHSE, ArgInt, LocalHSE, LocalInt, ArgHSEPtr[1:1], LocalHSEArr[1:1])
- // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
- // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = zero> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = zero> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <modifiers = zero> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -293,7 +293,7 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <modifiers = zero> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -301,19 +301,19 @@ struct Struct {
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <modifiers = zero> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: } cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("ArgHSE") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("ArgInt") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("LocalHSE") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("LocalInt") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_create) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_create) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("ArgHSE") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("ArgInt") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("LocalHSE") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("LocalInt") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_create) name("ArgHSEPtr[1:1]") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_create) name("LocalHSEArr[1:1]") <modifiers = zero>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
@@ -347,8 +347,8 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK: }
int LocalInt;
#pragma acc declare create(zero:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = zero> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -356,13 +356,13 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <modifiers = zero> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
#pragma acc declare create(zero:LocalHSE, LocalInt, LocalHSEArr[1:1])
// CHECK-NEXT: cir.cleanup.scope {
- // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[LOC_HSE_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("LocalHSE") <modifiers = zero> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[LOC_INT_CREATE:.*]] = acc.create varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!s32i>) name("LocalInt") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -370,16 +370,16 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[LOC_HSE_ARR_CREATE:.*]] = acc.create varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) name("LocalHSEArr[1:1]") <modifiers = zero> -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
// CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: cir.cleanup.scope {
// CHECK-NEXT: cir.yield
// CHECK-NEXT: cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_CREATE]], %[[LOC_INT_CREATE]], %[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("LocalHSE") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("LocalInt") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_create) name("LocalHSEArr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("LocalHSE") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("LocalInt") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUND2]]) dataClause(acc_create) name("LocalHSEArr[1:1]") <modifiers = zero>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
@@ -387,9 +387,9 @@ void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEP
// CHECK-NEXT: cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("ArgHSE") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("ArgInt") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_create) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("ArgHSE") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("ArgInt") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_create) name("ArgHSEPtr[1:1]") <modifiers = zero>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
@@ -415,8 +415,8 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK: }
int LocalInt;
#pragma acc declare create(zero:ArgHSE, ArgInt, ArgHSEPtr[1:1])
- // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!rec_HasSideEffects>
- // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK: %[[ARG_HSE_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!rec_HasSideEffects>) name("ArgHSE") <modifiers = zero> -> !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[ARG_INT_CREATE:.*]] = acc.create varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!s32i>) name("ArgInt") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[LB:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
@@ -424,7 +424,7 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
- // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
+ // CHECK-NEXT: %[[ARG_HSE_PTR_CREATE:.*]] = acc.create varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) name("ArgHSEPtr[1:1]") <modifiers = zero> -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>>
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
{
// CHECK-NEXT: cir.cleanup.scope {
@@ -463,9 +463,9 @@ extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *Ar
// CHECK-NEXT: cleanup normal {
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_CREATE]], %[[ARG_INT_CREATE]], %[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>)
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("ArgHSE") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("ArgInt") <{modifiers = #acc<data_clause_modifier zero>}>
- // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_create) name("ArgHSEPtr[1:1]") <{modifiers = #acc<data_clause_modifier zero>}>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_CREATE]] : !cir.ptr<!rec_HasSideEffects>) dataClause(acc_create) name("ArgHSE") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_CREATE]] : !cir.ptr<!s32i>) dataClause(acc_create) name("ArgInt") <modifiers = zero>
+ // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_CREATE]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) bounds(%[[BOUND1]]) dataClause(acc_create) name("ArgHSEPtr[1:1]") <modifiers = zero>
// CHECK-NEXT: cir.yield
// CHECK-NEXT: }
}
diff --git a/clang/test/CIR/CodeGenOpenACC/enter-data.c b/clang/test/CIR/CodeGenOpenACC/enter-data.c
index 208a6ef155e31..f1e04f29abd8e 100644
--- a/clang/test/CIR/CodeGenOpenACC/enter-data.c
+++ b/clang/test/CIR/CodeGenOpenACC/enter-data.c
@@ -11,27 +11,27 @@ void acc_data(int parmVar, int *ptrParmVar) {
// CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
#pragma acc enter data copyin(readonly, alwaysin: parmVar)
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier readonly,alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <modifiers = "readonly,alwaysin"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
#pragma acc enter data copyin(readonly, alwaysin: parmVar) async
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier readonly,alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async structured(false) name("parmVar") <modifiers = "readonly,alwaysin"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
#pragma acc enter data async copyin(readonly, alwaysin: parmVar)
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier readonly,alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async structured(false) name("parmVar") <modifiers = "readonly,alwaysin"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
#pragma acc enter data copyin(readonly, alwaysin: parmVar) async(parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = cir.builtin_int_cast %[[PARM_LOAD]]
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier readonly,alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) structured(false) name("parmVar") <modifiers = "readonly,alwaysin"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
#pragma acc enter data async(parmVar) copyin(readonly, alwaysin: parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = cir.builtin_int_cast %[[PARM_LOAD]]
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier readonly,alwaysin>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) structured(false) name("parmVar") <modifiers = "readonly,alwaysin"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
#pragma acc enter data create(parmVar)
@@ -39,17 +39,17 @@ void acc_data(int parmVar, int *ptrParmVar) {
// CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
#pragma acc enter data create(zero: parmVar)
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
#pragma acc enter data create(zero: parmVar) async
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async structured(false) name("parmVar") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data async dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
#pragma acc enter data create(zero: parmVar) async(parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = cir.builtin_int_cast %[[PARM_LOAD]]
- // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) structured(false) name("parmVar") <modifiers = zero> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
#pragma acc enter data attach(ptrParmVar)
diff --git a/clang/test/CIR/CodeGenOpenACC/exit-data.c b/clang/test/CIR/CodeGenOpenACC/exit-data.c
index d1bbc3cbcb3f2..902ef0b1e139a 100644
--- a/clang/test/CIR/CodeGenOpenACC/exit-data.c
+++ b/clang/test/CIR/CodeGenOpenACC/exit-data.c
@@ -12,33 +12,33 @@ void acc_data(int parmVar, int *ptrParmVar) {
// CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar")
#pragma acc exit data copyout(zero, alwaysout: parmVar)
- // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_copyout) structured(false) name("parmVar") <modifiers = "zero,alwaysout"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
- // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <modifiers = "zero,alwaysout">
#pragma acc exit data copyout(zero, alwaysout: parmVar) async
- // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async dataClause(acc_copyout) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async dataClause(acc_copyout) structured(false) name("parmVar") <modifiers = "zero,alwaysout"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
- // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <modifiers = "zero,alwaysout">
#pragma acc exit data async copyout(zero, alwaysout: parmVar)
- // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async dataClause(acc_copyout) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async dataClause(acc_copyout) structured(false) name("parmVar") <modifiers = "zero,alwaysout"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
- // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <modifiers = "zero,alwaysout">
#pragma acc exit data finalize copyout(zero, alwaysout: parmVar) async(parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = cir.builtin_int_cast %[[PARM_LOAD]]
- // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) dataClause(acc_copyout) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) dataClause(acc_copyout) structured(false) name("parmVar") <modifiers = "zero,alwaysout"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>) finalize
- // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <modifiers = "zero,alwaysout">
#pragma acc exit data async(parmVar) copyout(zero, alwaysout: parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = cir.builtin_int_cast %[[PARM_LOAD]]
- // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) dataClause(acc_copyout) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}> -> !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) dataClause(acc_copyout) structured(false) name("parmVar") <modifiers = "zero,alwaysout"> -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
- // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <{modifiers = #acc<data_clause_modifier zero,alwaysout>}>
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) structured(false) name("parmVar") <modifiers = "zero,alwaysout">
#pragma acc exit data delete(parmVar) finalize
// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) dataClause(acc_delete) structured(false) name("parmVar") -> !cir.ptr<!s32i>
More information about the cfe-commits
mailing list