https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/145770
Some of the 'data' clauses can have a 'modifier-list' which specifies one of a
few keywords from a list. This patch adds support for lowering them following
#144806.
We have to keep a separate enum from MLIR, since we have to keep 'always'
around for semantic reasons, whereas the dialect doesn't differentiate these.
This patch ensures we get these right for the only applicable clause so far,
which is 'copy'.
>From 188633e5eed60d0db90cd4256ed0f5b13a369082 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Wed, 25 Jun 2025 11:55:50 -0700
Subject: [PATCH] [OpenACC][CIR] Implement 'modifier-list' lowering
Some of the 'data' clauses can have a 'modifier-list' which specifies
one of a few keywords from a list. This patch adds support for lowering
them following #144806.
We have to keep a separate enum from MLIR, since we have to keep
'always' around for semantic reasons, whereas the dialect doesn't
differentiate these.
This patch ensures we get these right for the only applicable clause so
far, which is 'copy'.
---
clang/include/clang/Basic/OpenACCKinds.h | 19 +++--
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 47 ++++++++----
clang/test/CIR/CodeGenOpenACC/combined-copy.c | 73 +++++++++++++++++--
clang/test/CIR/CodeGenOpenACC/compute-copy.c | 59 +++++++++++++--
4 files changed, 160 insertions(+), 38 deletions(-)
diff --git a/clang/include/clang/Basic/OpenACCKinds.h
b/clang/include/clang/Basic/OpenACCKinds.h
index a7a29e2add20a..938385679e3ab 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -634,16 +634,19 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream
&Out,
}
// Represents the 'modifier' of a 'modifier-list', as applied to copy, copyin,
-// copyout, and create. Implemented as a 'bitmask'
+// copyout, and create. Implemented as a 'bitmask'.
+// Note: This attempts to synchronize with mlir::acc::DataClauseModifier,
+// however has to store `Always` separately(whereas MLIR has it as AlwaysIn &
+// AlwaysOut). However, we keep them in sync so that we can cast between them.
enum class OpenACCModifierKind : uint8_t {
Invalid = 0,
- Always = 1 << 0,
- AlwaysIn = 1 << 1,
- AlwaysOut = 1 << 2,
- Readonly = 1 << 3,
- Zero = 1 << 4,
- Capture = 1 << 5,
- LLVM_MARK_AS_BITMASK_ENUM(Capture)
+ Zero = 1 << 0,
+ Readonly = 1 << 1,
+ AlwaysIn = 1 << 2,
+ AlwaysOut = 1 << 3,
+ Capture = 1 << 4,
+ Always = 1 << 5,
+ LLVM_MARK_AS_BITMASK_ENUM(Always)
};
inline bool isOpenACCModifierBitSet(OpenACCModifierKind List,
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 899e91574e917..c93ca803b14ae 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -286,16 +286,28 @@ class OpenACCClauseCIREmitter final
std::move(bounds)};
}
+ mlir::acc::DataClauseModifier
+ convertModifiers(OpenACCModifierKind modifiers) {
+ using namespace mlir::acc;
+ DataClauseModifier mlirModifiers{};
+
+ // The MLIR representation of this represents `always` as `alwaysin` +
+ // `alwaysout`. So do a small fixup here.
+ if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) {
+ mlirModifiers = mlirModifiers | DataClauseModifier::always;
+ modifiers &= ~OpenACCModifierKind::Always;
+ }
+
+ mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
+ return mlirModifiers;
+ }
+
template <typename BeforeOpTy, typename AfterOpTy>
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
- bool structured, bool implicit) {
+ OpenACCModifierKind modifiers, bool structured,
+ bool implicit) {
DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
- // TODO: OpenACC: we should comprehend the 'modifier-list' here for the
data
- // operand. At the moment, we don't have a uniform way to assign these
- // properly, and the dialect cannot represent anything other than
'readonly'
- // and 'zero' on copyin/copyout/create, so for now, we skip it.
-
auto beforeOp =
builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue,
structured,
implicit, opInfo.name, opInfo.bounds);
@@ -323,6 +335,8 @@ class OpenACCClauseCIREmitter final
// Set the 'rest' of the info for both operations.
beforeOp.setDataClause(dataClause);
afterOp.setDataClause(dataClause);
+ beforeOp.setModifiers(convertModifiers(modifiers));
+ afterOp.setModifiers(convertModifiers(modifiers));
// Make sure we record these, so 'async' values can be updated later.
dataOperands.push_back(beforeOp.getOperation());
@@ -331,7 +345,8 @@ class OpenACCClauseCIREmitter final
template <typename BeforeOpTy>
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
- bool structured, bool implicit) {
+ OpenACCModifierKind modifiers, bool structured,
+ bool implicit) {
DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
auto beforeOp =
builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue,
structured,
@@ -340,6 +355,8 @@ class OpenACCClauseCIREmitter final
// Set the 'rest' of the info for the operation.
beforeOp.setDataClause(dataClause);
+ beforeOp.setModifiers(convertModifiers(modifiers));
+
// Make sure we record these, so 'async' values can be updated later.
dataOperands.push_back(beforeOp.getOperation());
}
@@ -818,7 +835,8 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp>) {
for (auto var : clause.getVarList())
addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
- var, mlir::acc::DataClause::acc_copy, /*structured=*/true,
+ var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
+ /*structured=*/true,
/*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
@@ -833,8 +851,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
for (auto var : clause.getVarList())
addDataOperand<mlir::acc::UseDeviceOp>(
- var, mlir::acc::DataClause::acc_use_device,
- /*structured=*/true, /*implicit=*/false);
+ var, mlir::acc::DataClause::acc_use_device, {},
/*structured=*/true,
+ /*implicit=*/false);
} else {
llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
}
@@ -845,7 +863,8 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp>) {
for (auto var : clause.getVarList())
addDataOperand<mlir::acc::DevicePtrOp>(
- var, mlir::acc::DataClause::acc_deviceptr, /*structured=*/true,
+ var, mlir::acc::DataClause::acc_deviceptr, {},
+ /*structured=*/true,
/*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
@@ -861,7 +880,7 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp>) {
for (auto var : clause.getVarList())
addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
- var, mlir::acc::DataClause::acc_no_create, /*structured=*/true,
+ var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
/*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
@@ -877,7 +896,7 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp>) {
for (auto var : clause.getVarList())
addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
- var, mlir::acc::DataClause::acc_present, /*structured=*/true,
+ var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
/*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
@@ -893,7 +912,7 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp>) {
for (auto var : clause.getVarList())
addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
- var, mlir::acc::DataClause::acc_attach, /*structured=*/true,
+ var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
/*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c
b/clang/test/CIR/CodeGenOpenACC/combined-copy.c
index 72471d4ec7874..4d609c747f16d 100644
--- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c
+++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c
@@ -77,29 +77,29 @@ void acc_compute(int parmVar) {
// these do nothing to the IR.
#pragma acc parallel loop copy(alwaysin: localVar1) copy(alwaysout: localVar2)
copy(always: localVar3)
for(int i = 0; i < 5; ++i);
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
name = "localVar1"} loc
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] :
!cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>,
name = "localVar2"} loc
- // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause
acc_copy>, name = "localVar3"} loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] :
!cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"}
loc
// CHECK-NEXT: 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<data_clause
acc_copy>, name = "localVar3"} loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to
varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>,
name = "localVar2"} loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>,
name = "localVar1"} loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to
varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"}
loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to
varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} 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>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
name = "localVar1"} loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always>, name = "localVar1"} 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<data_clause acc_copy>,
name = "localVar1"} loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc
short *localPointer;
float localArray[5];
@@ -1102,3 +1102,60 @@ void copy_member_of_array_element_member() {
// CHECK-NEXT: } loc
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[GETB]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>,
name = "outer.inner[2].b"}
}
+
+void modifier_list() {
+ // CHECK: cir.func @modifier_list() {
+ int localVar;
+ // CHECK-NEXT: %[[LOCALVAR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>,
["localVar"]
+
+#pragma acc parallel loop copy(always:localVar)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always>, name = "localVar"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!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.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"}
+#pragma acc serial loop copy(alwaysin:localVar)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"}
+ // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] :
!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.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"}
+#pragma acc kernels loop copy(alwaysout:localVar)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"}
+ // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] :
!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(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"}
+#pragma acc parallel loop copy(capture:localVar)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier capture>, name = "localVar"}
+ // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] :
!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.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"}
+#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>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"}
+ // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] :
!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.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name =
"localVar"}
+}
diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copy.c
b/clang/test/CIR/CodeGenOpenACC/compute-copy.c
index 888bad29caa7c..2fba123b256c9 100644
--- a/clang/test/CIR/CodeGenOpenACC/compute-copy.c
+++ b/clang/test/CIR/CodeGenOpenACC/compute-copy.c
@@ -69,23 +69,23 @@ void acc_compute(int parmVar) {
// these do nothing to the IR.
#pragma acc parallel copy(alwaysin: localVar1) copy(alwaysout: localVar2)
copy(always: localVar3)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
name = "localVar1"} loc
- // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] :
!cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>,
name = "localVar2"} loc
- // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause
acc_copy>, name = "localVar3"} loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc
+ // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] :
!cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc
+ // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] :
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"}
loc
// CHECK-NEXT: 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<data_clause
acc_copy>, name = "localVar3"} loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to
varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>,
name = "localVar2"} loc
- // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>,
name = "localVar1"} loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to
varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"}
loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to
varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc
#pragma acc serial copy(always, alwaysin, alwaysout: localVar1)
;
- // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
name = "localVar1"} loc
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always>, name = "localVar1"} 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<data_clause acc_copy>,
name = "localVar1"} loc
+ // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to
varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc
short *localPointer;
float localArray[5];
@@ -897,3 +897,46 @@ void acc_compute_members() {
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] :
!cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]])
to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>)
{dataClause = #acc<data_clause acc_copy>, name =
"localStruct.ptrPtrMember[1:3][1:1]"}
}
+
+void modifier_list() {
+ // CHECK: cir.func @modifier_list() {
+ int localVar;
+ // CHECK-NEXT: %[[LOCALVAR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>,
["localVar"]
+
+#pragma acc parallel copy(always:localVar)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always>, name = "localVar"}
+ // 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<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"}
+
+#pragma acc serial copy(alwaysin:localVar)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"}
+ // 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<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"}
+#pragma acc kernels copy(alwaysout:localVar)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"}
+ // 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<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"}
+#pragma acc parallel copy(capture:localVar)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier capture>, name = "localVar"}
+ // 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<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"}
+#pragma acc serial copy(capture, always, alwaysin, alwaysout:localVar)
+ ;
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>,
modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"}
+ // 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<data_clause
acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name =
"localVar"}
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
