https://github.com/erichkeane updated
https://github.com/llvm/llvm-project/pull/146146
>From 27f2b1d2af4fb5f5befd5709c199ae4616d676e2 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Fri, 27 Jun 2025 12:06:19 -0700
Subject: [PATCH 1/2] [OpenACC][CIR] Implement enter-data + clause lowering
'enter data' is a new construct type that requires one of the data
clauses, so we had to wait for all clauses to be ready before we could
commit this. Most of the clauses are simple, but there is a little bit
of work to get 'async' and 'wait' to have similar interfaces in the ACC
dialect, where helpers were added.
---
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 68 +++++++---
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 6 +-
clang/test/CIR/CodeGenOpenACC/enter-data.c | 125 ++++++++++++++++++
.../mlir/Dialect/OpenACC/OpenACCOps.td | 19 +++
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 47 +++++++
5 files changed, 247 insertions(+), 18 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenACC/enter-data.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index b17994ee8771e..3546b6562c021 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -376,10 +376,19 @@ class OpenACCClauseCIREmitter final
// on all operation types.
mlir::ArrayAttr getAsyncOnlyAttr() {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
- mlir::acc::KernelsOp, mlir::acc::DataOp>)
+ mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOnlyAttr();
- else if constexpr (isCombinedType<OpTy>)
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ if (!operation.getAsyncAttr())
+ return mlir::ArrayAttr{};
+
+ llvm::SmallVector<mlir::Attribute> devTysTemp;
+ devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(), mlir::acc::DeviceType::None));
+ return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
+ } else if constexpr (isCombinedType<OpTy>) {
return operation.computeOp.getAsyncOnlyAttr();
+ }
// Note: 'wait' has async as well, but it cannot have data clauses, so we
// don't have to handle them here.
@@ -391,10 +400,19 @@ class OpenACCClauseCIREmitter final
// on all operation types.
mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
- mlir::acc::KernelsOp, mlir::acc::DataOp>)
+ mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOperandsDeviceTypeAttr();
- else if constexpr (isCombinedType<OpTy>)
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ if (!operation.getAsyncOperand())
+ return mlir::ArrayAttr{};
+
+ llvm::SmallVector<mlir::Attribute> devTysTemp;
+ devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(), mlir::acc::DeviceType::None));
+ return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
+ } else if constexpr (isCombinedType<OpTy>) {
return operation.computeOp.getAsyncOperandsDeviceTypeAttr();
+ }
// Note: 'wait' has async as well, but it cannot have data clauses, so we
// don't have to handle them here.
@@ -409,6 +427,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>)
return operation.getAsyncOperands();
+ else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>)
+ return operation.getAsyncOperandMutable();
else if constexpr (isCombinedType<OpTy>)
return operation.computeOp.getAsyncOperands();
@@ -542,10 +562,11 @@ class OpenACCClauseCIREmitter final
void VisitAsyncClause(const OpenACCAsyncClause &clause) {
hasAsyncClause = true;
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
- mlir::acc::KernelsOp, mlir::acc::DataOp>) {
- if (!clause.hasIntExpr())
+ mlir::acc::KernelsOp, mlir::acc::DataOp,
+ mlir::acc::EnterDataOp>) {
+ if (!clause.hasIntExpr()) {
operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
- else {
+ } else {
mlir::Value intExpr;
{
@@ -572,8 +593,8 @@ class OpenACCClauseCIREmitter final
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Combined constructs remain. Data, enter data, exit data,
- // update constructs remain.
+ // unreachable. Combined constructs remain. Exit data,update constructs
+ // remain.
return clauseNotImplemented(clause);
}
}
@@ -604,7 +625,7 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp, mlir::acc::InitOp,
mlir::acc::ShutdownOp, mlir::acc::SetOp,
mlir::acc::DataOp, mlir::acc::WaitOp,
- mlir::acc::HostDataOp>) {
+ mlir::acc::HostDataOp, mlir::acc::EnterDataOp>)
{
operation.getIfCondMutable().append(
createCondition(clause.getConditionExpr()));
} else if constexpr (isCombinedType<OpTy>) {
@@ -659,7 +680,8 @@ class OpenACCClauseCIREmitter final
void VisitWaitClause(const OpenACCWaitClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
mlir::acc::SerialOp,
- mlir::acc::KernelsOp, mlir::acc::DataOp>) {
+ mlir::acc::KernelsOp, mlir::acc::DataOp,
+ mlir::acc::EnterDataOp>) {
if (!clause.hasExprs()) {
operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
} else {
@@ -866,11 +888,16 @@ class OpenACCClauseCIREmitter final
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::CopyinOp>(
+ var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
+ /*structured=*/false, /*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. enter-data, declare constructs remain.
+ // unreachable. declare construct remains.
return clauseNotImplemented(clause);
}
}
@@ -900,11 +927,16 @@ class OpenACCClauseCIREmitter final
var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::CreateOp>(
+ var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
+ /*structured=*/false, /*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. enter-data, declare constructs remain.
+ // unreachable. declare construct remains.
return clauseNotImplemented(clause);
}
}
@@ -974,12 +1006,15 @@ class OpenACCClauseCIREmitter final
addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
/*implicit=*/false);
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::AttachOp>(
+ var, mlir::acc::DataClause::acc_attach, {},
+ /*structured=*/false, /*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
- // TODO: When we've implemented this for everything, switch this to an
- // unreachable. enter data remains.
- return clauseNotImplemented(clause);
+ llvm_unreachable("Unknown construct kind in VisitAttachClause");
}
}
};
@@ -1018,6 +1053,7 @@ EXPL_SPEC(mlir::acc::ShutdownOp)
EXPL_SPEC(mlir::acc::SetOp)
EXPL_SPEC(mlir::acc::WaitOp)
EXPL_SPEC(mlir::acc::HostDataOp)
+EXPL_SPEC(mlir::acc::EnterDataOp)
#undef EXPL_SPEC
template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 1feefa55eb270..10a5601476f4e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -250,8 +250,10 @@ mlir::LogicalResult
CIRGenFunction::emitOpenACCHostDataConstruct(
mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
const OpenACCEnterDataConstruct &s) {
- cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
- return mlir::failure();
+ mlir::Location start = getLoc(s.getSourceRange().getBegin());
+ emitOpenACCOp<EnterDataOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+ s.clauses());
+ return mlir::success();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
const OpenACCExitDataConstruct &s) {
diff --git a/clang/test/CIR/CodeGenOpenACC/enter-data.c
b/clang/test/CIR/CodeGenOpenACC/enter-data.c
new file mode 100644
index 0000000000000..1785fba1a1059
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/enter-data.c
@@ -0,0 +1,125 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir
-fclangir %s -o - | FileCheck %s
+void acc_data(int parmVar, int *ptrParmVar) {
+ // CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]:
!cir.ptr<!s32i>{{.*}}) {
+ // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar",
init]
+ // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>,
!cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>
+ // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>,
!cir.ptr<!cir.ptr<!s32i>>
+
+#pragma acc enter data copyin(parmVar)
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // 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>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
readonly,alwaysin>, name = "parmVar", structured = false}
+ // 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 -> !cir.ptr<!s32i> {modifiers =
#acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured =
false}
+ // 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 -> !cir.ptr<!s32i> {modifiers =
#acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured =
false}
+ // 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:.*]] = builtin.unrealized_conversion_cast
%[[PARM_LOAD]]
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers =
#acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured =
false}
+ // 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:.*]] = builtin.unrealized_conversion_cast
%[[PARM_LOAD]]
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers =
#acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured =
false}
+ // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32)
dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
+
+#pragma acc enter data create(parmVar)
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // 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>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier
zero>, name = "parmVar", structured = false}
+ // 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 -> !cir.ptr<!s32i> {modifiers =
#acc<data_clause_modifier zero>, name = "parmVar", structured = false}
+ // 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:.*]] = builtin.unrealized_conversion_cast
%[[PARM_LOAD]]
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers =
#acc<data_clause_modifier zero>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32)
dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
+
+#pragma acc enter data attach(ptrParmVar)
+ // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] :
!cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar",
structured = false}
+ // CHECK-NEXT: acc.enter_data dataOperands(%[[ATTACH1]] :
!cir.ptr<!cir.ptr<!s32i>>)
+
+#pragma acc enter data attach(ptrParmVar) async
+ // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] :
!cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name =
"ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data async dataOperands(%[[ATTACH1]] :
!cir.ptr<!cir.ptr<!s32i>>)
+
+#pragma acc enter data attach(ptrParmVar) async(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast
%[[PARM_LOAD]]
+ // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] :
!cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) ->
!cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32)
dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)
+
+#pragma acc enter data if (parmVar == 1) copyin(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) dataOperands(%[[COPYIN1]] :
!cir.ptr<!s32i>)
+
+#pragma acc enter data async if (parmVar == 1) copyin(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async
dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
+
+#pragma acc enter data if (parmVar == 1) async(parmVar) copyin(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast
%[[PARM_LOAD]]
+ // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] :
!cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {name =
"parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async(%[[PARM_CAST]] : si32)
dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)
+
+#pragma acc enter data wait create(parmVar)
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data wait dataOperands(%[[CREATE1]] :
!cir.ptr<!s32i>)
+
+#pragma acc enter data wait(1) create(parmVar)
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE_CONST]]
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data wait(%[[ONE_CAST]] : si32)
dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
+
+#pragma acc enter data wait(parmVar, 1, 2) create(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast
%[[PARM_LOAD]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE_CONST]]
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO_CONST]]
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data wait(%[[PARM_CAST]], %[[ONE_CAST]],
%[[TWO_CAST]] : si32, si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
+
+#pragma acc enter data wait(devnum: parmVar: 1, 2) create(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast
%[[PARM_LOAD]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast
%[[ONE_CONST]]
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast
%[[TWO_CONST]]
+ // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] :
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.enter_data wait_devnum(%[[PARM_CAST]] : si32)
wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) dataOperands(%[[CREATE1]] :
!cir.ptr<!s32i>)
+
+}
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 8cbdf710cfa6e..3403e158c9f58 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2010,6 +2010,25 @@ def OpenACC_EnterDataOp : OpenACC_Op<"enter_data",
/// The i-th data operand passed.
Value getDataOperand(unsigned i);
+
+ /// Add an entry to the 'async-only' attribute (clause spelled without
+ /// arguments). DeviceType array is supplied even though it should always
be
+ /// empty, so this can mirror other versions of this function.
+ void addAsyncOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
+ /// Add a value to the 'async'. DeviceType array is supplied even though it
+ /// should always be empty, so this can mirror other versions of this
+ /// function.
+ void addAsyncOperand(MLIRContext *, mlir::Value,
+ llvm::ArrayRef<DeviceType>);
+ /// Add an entry to the 'wait-only' attribute (clause spelled without
+ /// arguments). DeviceType array is supplied even though it should always
be
+ /// empty, so this can mirror other versions of this function.
+ void addWaitOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
+ /// Add an array-like entry to the 'wait'. DeviceType array is supplied
+ /// even though it should always be empty, so this can mirror other
versions
+ /// of this function.
+ void addWaitOperands(MLIRContext *, bool hasDevnum, mlir::ValueRange,
+ llvm::ArrayRef<DeviceType>);
}];
let assemblyFormat = [{
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 37acb6acbfa91..f0516ef0f0f62 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -3218,6 +3218,53 @@ void
EnterDataOp::getCanonicalizationPatterns(RewritePatternSet &results,
results.add<RemoveConstantIfCondition<EnterDataOp>>(context);
}
+void EnterDataOp::addAsyncOnly(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ assert(effectiveDeviceTypes.empty());
+ assert(!getAsyncAttr());
+ assert(!getAsyncOperand());
+
+ setAsyncAttr(mlir::UnitAttr::get(context));
+}
+
+void EnterDataOp::addAsyncOperand(
+ MLIRContext *context, mlir::Value newValue,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ assert(effectiveDeviceTypes.empty());
+ assert(!getAsyncAttr());
+ assert(!getAsyncOperand());
+
+ getAsyncOperandMutable().append(newValue);
+}
+
+void EnterDataOp::addWaitOnly(MLIRContext *context,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes)
{
+ assert(effectiveDeviceTypes.empty());
+ assert(!getWaitAttr());
+ assert(getWaitOperands().empty());
+ assert(!getWaitDevnum());
+
+ setWaitAttr(mlir::UnitAttr::get(context));
+}
+
+void EnterDataOp::addWaitOperands(
+ MLIRContext *context, bool hasDevnum, mlir::ValueRange newValues,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ assert(effectiveDeviceTypes.empty());
+ assert(!getWaitAttr());
+ assert(getWaitOperands().empty());
+ assert(!getWaitDevnum());
+
+ // if hasDevnum, the first value is the devnum. The 'rest' go into the
+ // operands list.
+ if (hasDevnum) {
+ getWaitDevnumMutable().append(newValues.front());
+ newValues = newValues.drop_front();
+ }
+
+ getWaitOperandsMutable().append(newValues);
+}
+
//===----------------------------------------------------------------------===//
// AtomicReadOp
//===----------------------------------------------------------------------===//
>From d409bb9072f12d3896516d4fa4af78adf3faefa2 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Fri, 27 Jun 2025 13:39:29 -0700
Subject: [PATCH 2/2] Add missing space in comment
---
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 3546b6562c021..d982cc92d9b4b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -593,7 +593,7 @@ class OpenACCClauseCIREmitter final
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Combined constructs remain. Exit data,update constructs
+ // unreachable. Combined constructs remain. Exit data, update constructs
// remain.
return clauseNotImplemented(clause);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
