https://github.com/erichkeane updated
https://github.com/llvm/llvm-project/pull/145976
>From b1c4ebc8e9efba38f2cda2696f60cc4a86c3fd89 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Thu, 26 Jun 2025 09:02:15 -0700
Subject: [PATCH 1/2] [OpenACC][CIR] Implement copyin/copyout/create lowering
for compute/combined
This patch does the lowering of copyin (represented as a
acc.copyin/acc.delete), copyout (acc.create/acc.copyin), and create
(acc.create/acc.delete).
Additionally, it found a few problems with #144806, so it fixes those as
well.
---
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 51 ++++++
.../combined-copyin-copyout-create.c | 160 ++++++++++++++++++
.../compute-copyin-copyout-create.c | 128 ++++++++++++++
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 6 +-
mlir/test/Dialect/OpenACC/ops.mlir | 12 +-
5 files changed, 354 insertions(+), 3 deletions(-)
create mode 100644
clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c
create mode 100644
clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 1454cee336a09..fe4145959b206 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -858,6 +858,57 @@ class OpenACCClauseCIREmitter final
}
}
+ void VisitCopyInClause(const OpenACCCopyInClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
+ mlir::acc::KernelsOp>) {
+ for (auto var : clause.getVarList())
+ addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
+ var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
+ /*structured=*/true,
+ /*implicit=*/false);
+ } else if constexpr (isCombinedType<OpTy>) {
+ applyToComputeOp(clause);
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. data, declare, combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
+
+ void VisitCopyOutClause(const OpenACCCopyOutClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
+ mlir::acc::KernelsOp>) {
+ for (auto var : clause.getVarList())
+ addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
+ var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
+ /*structured=*/true,
+ /*implicit=*/false);
+ } else if constexpr (isCombinedType<OpTy>) {
+ applyToComputeOp(clause);
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. data, declare, combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
+
+ void VisitCreateClause(const OpenACCCreateClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
+ mlir::acc::KernelsOp>) {
+ for (auto var : clause.getVarList())
+ addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
+ var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
+ /*structured=*/true,
+ /*implicit=*/false);
+ } else if constexpr (isCombinedType<OpTy>) {
+ applyToComputeOp(clause);
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. data, declare, combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
+
void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
for (auto var : clause.getVarList())
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c
b/clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c
new file mode 100644
index 0000000000000..d6179c012ee91
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/combined-copyin-copyout-create.c
@@ -0,0 +1,160 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir
-fclangir %s -o - | FileCheck %s
+
+void acc_combined(int parmVar) {
+ // CHECK: cir.func{{.*}} @acc_combined(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar",
init]
+
+ int localVar1;
+ // CHECK-NEXT: %[[LV1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["localVar1"]
+ float localVar2;
+ // CHECK-NEXT: %[[LV2:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>,
["localVar2"]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]]
+#pragma acc parallel loop copyin(parmVar) copyout(localVar1) create(localVar2)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, name = "localVar1"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "localVar2"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]],
%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>,
!cir.ptr<!cir.float>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!cir.float>)
{dataClause = #acc<data_clause acc_create>, name = "localVar2"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[LV1]] : !cir.ptr<!s32i>) {name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, name = "parmVar"}
+
+#pragma acc serial loop copyin(parmVar, localVar1)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
+ // 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: } loc
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, name = "parmVar"}
+
+#pragma acc kernels loop copyout(parmVar, localVar1)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, name = "localVar1"}
+ // 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: } 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"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar"}
+
+#pragma acc parallel loop create (parmVar, localVar2)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "localVar2"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[CREATE1]],
%[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!cir.float>)
{dataClause = #acc<data_clause acc_create>, name = "localVar2"}
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_create>, name = "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>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always>, name = "localVar1"}
+ // 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: } loc
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>,
name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier capture>,
name = "parmVar"}
+
+#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>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"}
+ // 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: } loc
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to
varPtr(%[[LV1]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier
always>, name = "localVar1"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier
capture>, name = "parmVar"}
+
+#pragma acc parallel loop create(capture: parmVar)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
capture>, name = "parmVar"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[CREATE1]] :
!cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.loop combined(parallel) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier capture>,
name = "parmVar"}
+
+#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>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,capture>, name = "localVar1"}
+ // 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: } loc
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,capture>, name = "parmVar"}
+
+#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>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "localVar1"}
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {modifiers =
#acc<data_clause_modifier always,readonly,capture>, name = "localVar2"}
+ // 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: } loc
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>)
{dataClause = #acc<data_clause acc_copyin>, modifiers =
#acc<data_clause_modifier always,readonly,capture>, name = "localVar2"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "parmVar"}
+
+#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>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name
= "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name
= "localVar1"}
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name
= "localVar2"}
+ // 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: } 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>) {modifiers = #acc<data_clause_modifier
always,zero,capture>, name = "localVar2"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to
varPtr(%[[LV1]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier
always,zero,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier
always,zero,capture>, name = "parmVar"}
+
+#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>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
zero,capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
zero,capture>, name = "localVar1"}
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {modifiers =
#acc<data_clause_modifier zero,capture>, name = "localVar2"}
+ // 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: } loc
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE3]] : !cir.ptr<!cir.float>)
{dataClause = #acc<data_clause acc_create>, modifiers =
#acc<data_clause_modifier zero,capture>, name = "localVar2"}
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier
zero,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier
zero,capture>, name = "parmVar"}
+}
diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c
b/clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c
new file mode 100644
index 0000000000000..2180a3370939e
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/compute-copyin-copyout-create.c
@@ -0,0 +1,128 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir
-fclangir %s -o - | FileCheck %s
+
+void acc_compute(int parmVar) {
+ // CHECK: cir.func{{.*}} @acc_compute(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar",
init]
+
+ int localVar1;
+ // CHECK-NEXT: %[[LV1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["localVar1"]
+ float localVar2;
+ // CHECK-NEXT: %[[LV2:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>,
["localVar2"]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]]
+
+#pragma acc parallel copyin(parmVar) copyout(localVar1) create(localVar2)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, name = "localVar1"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "localVar2"}
+ // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]], %[[CREATE1]],
%[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!cir.float>)
{dataClause = #acc<data_clause acc_create>, name = "localVar2"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[LV1]] : !cir.ptr<!s32i>) {name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, name = "parmVar"}
+
+#pragma acc serial copyin(parmVar, localVar1)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
+ // 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<data_clause acc_copyin>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, name = "parmVar"}
+
+#pragma acc kernels copyout(parmVar, localVar1)
+ ;
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, name = "localVar1"}
+ // 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"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar"}
+
+#pragma acc parallel create (parmVar, localVar2)
+ ;
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "localVar2"}
+ // CHECK-NEXT: acc.parallel dataOperands(%[[CREATE1]], %[[CREATE2]] :
!cir.ptr<!s32i>, !cir.ptr<!cir.float>) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!cir.float>)
{dataClause = #acc<data_clause acc_create>, name = "localVar2"}
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_create>, name = "parmVar"}
+
+#pragma acc serial copyin(capture: parmVar) copyin(always: localVar1)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always>, name = "localVar1"}
+ // 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<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always>,
name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier capture>,
name = "parmVar"}
+
+#pragma acc kernels copyout(capture: parmVar) copyout(always: localVar1)
+ ;
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"}
+ // 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>) {modifiers = #acc<data_clause_modifier
always>, name = "localVar1"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier
capture>, name = "parmVar"}
+
+#pragma acc parallel create(capture: parmVar)
+ ;
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
capture>, name = "parmVar"}
+ // 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<data_clause acc_create>, modifiers = #acc<data_clause_modifier capture>,
name = "parmVar"}
+
+#pragma acc serial copyin(capture, always: parmVar, localVar1)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,capture>, name = "localVar1"}
+ // 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<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,capture>, name = "parmVar"}
+
+#pragma acc kernels copyin(readonly, always, alwaysin, capture: parmVar,
localVar1, localVar2)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "localVar1"}
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {modifiers =
#acc<data_clause_modifier always,readonly,capture>, name = "localVar2"}
+ // 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<data_clause acc_copyin>, modifiers =
#acc<data_clause_modifier always,readonly,capture>, name = "localVar2"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier
always,readonly,capture>, name = "parmVar"}
+
+#pragma acc parallel copyout(zero, always, alwaysout, capture: parmVar,
localVar1, localVar2)
+ ;
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name
= "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name
= "localVar1"}
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause
acc_copyout>, modifiers = #acc<data_clause_modifier always,zero,capture>, name
= "localVar2"}
+ // 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>) {modifiers = #acc<data_clause_modifier
always,zero,capture>, name = "localVar2"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to
varPtr(%[[LV1]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier
always,zero,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to
varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier
always,zero,capture>, name = "parmVar"}
+
+#pragma acc serial create(zero, capture: parmVar, localVar1, localVar2)
+ ;
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
zero,capture>, name = "parmVar"}
+ // CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
zero,capture>, name = "localVar1"}
+ // CHECK-NEXT: %[[CREATE3:.*]] = acc.create varPtr(%[[LV2]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {modifiers =
#acc<data_clause_modifier zero,capture>, name = "localVar2"}
+ // 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<data_clause acc_create>, modifiers =
#acc<data_clause_modifier zero,capture>, name = "localVar2"}
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier
zero,capture>, name = "localVar1"}
+ // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause
= #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier
zero,capture>, name = "parmVar"}
+}
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index e08fc263e29cc..37acb6acbfa91 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -575,8 +575,10 @@ LogicalResult acc::CreateOp::verify() {
return failure();
if (failed(checkVarAndAccVar(*this)))
return failure();
+ // this op is the entry part of copyout, so it also needs to allow all
+ // modifiers allowed on copyout.
if (failed(checkValidModifier(*this, acc::DataClauseModifier::zero |
- acc::DataClauseModifier::alwaysout |
+ acc::DataClauseModifier::always |
acc::DataClauseModifier::capture)))
return failure();
return success();
@@ -708,7 +710,7 @@ LogicalResult acc::DeleteOp::verify() {
// allowed on either case.
if (failed(checkValidModifier(*this, acc::DataClauseModifier::zero |
acc::DataClauseModifier::readonly |
- acc::DataClauseModifier::alwaysin |
+ acc::DataClauseModifier::always |
acc::DataClauseModifier::capture)))
return failure();
return success();
diff --git a/mlir/test/Dialect/OpenACC/ops.mlir
b/mlir/test/Dialect/OpenACC/ops.mlir
index c1d8276d904bb..7bb6cf43e49a7 100644
--- a/mlir/test/Dialect/OpenACC/ops.mlir
+++ b/mlir/test/Dialect/OpenACC/ops.mlir
@@ -924,15 +924,25 @@ func.func @testdataop(%a: memref<f32>, %b: memref<f32>,
%c: memref<f32>) -> () {
func.func @testdataopmodifiers(%a: memref<f32>, %b: memref<f32>, %c:
memref<f32>) -> () {
%0 = acc.create varPtr(%a : memref<f32>) -> memref<f32> {modifiers =
#acc<data_clause_modifier capture,zero>}
%1 = acc.copyin varPtr(%b : memref<f32>) -> memref<f32> {modifiers =
#acc<data_clause_modifier readonly,capture,always>}
- acc.data dataOperands(%0, %1 : memref<f32>, memref<f32>) {
+ %2 = acc.copyin varPtr(%c : memref<f32>) -> memref<f32> {modifiers =
#acc<data_clause_modifier always>}
+ %3 = acc.create varPtr(%c : memref<f32>) -> memref<f32> {modifiers =
#acc<data_clause_modifier always>}
+ acc.data dataOperands(%0, %1, %2, %3 : memref<f32>, memref<f32>,
memref<f32>, memref<f32>) {
}
acc.copyout accPtr(%0 : memref<f32>) to varPtr(%a : memref<f32>) {modifiers
= #acc<data_clause_modifier zero,capture,always>}
+ acc.delete accPtr(%2 : memref<f32>) {modifiers = #acc<data_clause_modifier
always>}
+ acc.copyout accPtr(%3 : memref<f32>) to varPtr(%c : memref<f32>) {modifiers
= #acc<data_clause_modifier always>}
+
func.return
}
+
// CHECK: func @testdataopmodifiers(%[[ARGA:.*]]: memref<f32>,
%[[ARGB:.*]]: memref<f32>, %[[ARGC:.*]]: memref<f32>) {
// CHECK: %[[CREATEA:.*]] = acc.create varPtr(%[[ARGA]] : memref<f32>) ->
memref<f32> {modifiers = #acc<data_clause_modifier zero,capture>}
// CHECK: %[[COPYINB:.*]] = acc.copyin varPtr(%[[ARGB]] : memref<f32>) ->
memref<f32> {modifiers = #acc<data_clause_modifier always,readonly,capture>}
+// CHECK: %[[COPYINC:.*]] = acc.copyin varPtr(%[[ARGC]] : memref<f32>) ->
memref<f32> {modifiers = #acc<data_clause_modifier always>}
+// CHECK: %[[CREATEC:.*]] = acc.create varPtr(%[[ARGC]] : memref<f32>) ->
memref<f32> {modifiers = #acc<data_clause_modifier always>}
// CHECK: acc.copyout accPtr(%[[CREATEA]] : memref<f32>) to
varPtr(%[[ARGA]] : memref<f32>) {modifiers = #acc<data_clause_modifier
always,zero,capture>}
+// CHECK: acc.delete accPtr(%[[COPYINC]] : memref<f32>) {modifiers =
#acc<data_clause_modifier always>}
+// CHECK: acc.copyout accPtr(%[[CREATEC]] : memref<f32>) to
varPtr(%[[ARGC]] : memref<f32>) {modifiers = #acc<data_clause_modifier always>}
// -----
>From 1e5ba1cf2ffde8671a58473417615583f8eaf834 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Fri, 27 Jun 2025 06:23:00 -0700
Subject: [PATCH 2/2] Fix mlir test, change type of 'var' in the loop
---
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 18 +++++++++---------
mlir/test/Dialect/OpenACC/invalid.mlir | 4 ++--
2 files changed, 11 insertions(+), 11 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index fe4145959b206..cce177056c6e0 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -844,7 +844,7 @@ class OpenACCClauseCIREmitter final
void VisitCopyClause(const OpenACCCopyClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
/*structured=*/true,
@@ -861,7 +861,7 @@ class OpenACCClauseCIREmitter final
void VisitCopyInClause(const OpenACCCopyInClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
/*structured=*/true,
@@ -878,7 +878,7 @@ class OpenACCClauseCIREmitter final
void VisitCopyOutClause(const OpenACCCopyOutClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
/*structured=*/true,
@@ -895,7 +895,7 @@ class OpenACCClauseCIREmitter final
void VisitCreateClause(const OpenACCCreateClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
/*structured=*/true,
@@ -911,7 +911,7 @@ class OpenACCClauseCIREmitter final
void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::UseDeviceOp>(
var, mlir::acc::DataClause::acc_use_device, {},
/*structured=*/true,
/*implicit=*/false);
@@ -923,7 +923,7 @@ class OpenACCClauseCIREmitter final
void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::DevicePtrOp>(
var, mlir::acc::DataClause::acc_deviceptr, {},
/*structured=*/true,
@@ -940,7 +940,7 @@ class OpenACCClauseCIREmitter final
void VisitNoCreateClause(const OpenACCNoCreateClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
/*implicit=*/false);
@@ -956,7 +956,7 @@ class OpenACCClauseCIREmitter final
void VisitPresentClause(const OpenACCPresentClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
/*implicit=*/false);
@@ -972,7 +972,7 @@ class OpenACCClauseCIREmitter final
void VisitAttachClause(const OpenACCAttachClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp>) {
- for (auto var : clause.getVarList())
+ for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
/*implicit=*/false);
diff --git a/mlir/test/Dialect/OpenACC/invalid.mlir
b/mlir/test/Dialect/OpenACC/invalid.mlir
index d85ad2ff80d80..68afd9fccba79 100644
--- a/mlir/test/Dialect/OpenACC/invalid.mlir
+++ b/mlir/test/Dialect/OpenACC/invalid.mlir
@@ -829,5 +829,5 @@ func.func @acc_loop_container() {
// -----
%value = memref.alloc() : memref<f32>
-// expected-error @below {{invalid data clause modifiers: alwaysin}}
-%0 = acc.create varPtr(%value : memref<f32>) -> memref<f32> {modifiers =
#acc<data_clause_modifier zero,capture,always>}
+// expected-error @below {{invalid data clause modifiers: readonly}}
+%0 = acc.create varPtr(%value : memref<f32>) -> memref<f32> {modifiers =
#acc<data_clause_modifier readonly,zero,capture,always>}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
