[clang] [mlir] [CIR][OpenACC] Correct enter-data wait merging behavior for all but d… (PR #228499)

[Erich Keane via cfe-commits]

https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/228499

…evnum.

We should be able to merge 'wait'.  On the bug report, we determined that 'empty wait' should effectively override everything, so this patch implements this.  Also, waits without a devnum are the same as just the same lists being merged.

The dialect still needs to figure out how to represent multiple dev-nums, so this is left as an assert, and clang just won't be able to get this to work yet.

Fixes: #227981

>From 624c2f309137fa16a079d9e87c8f8d4e5ab8e0f9 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 2 Oct 2026 08:58:04 -0700
Subject: [PATCH] [CIR][OpenACC] Correct enter-data wait merging behavior for
 all but devnum.

We should be able to merge 'wait'.  On the bug report, we determined
that 'empty wait' should effectively override everything, so this patch
implements this.  Also, waits without a devnum are the same as just the
same lists being merged.

The dialect still needs to figure out how to represent multiple
dev-nums, so this is left as an assert, and clang just won't be able to
get this to work yet.

Fixes: #227981
---
 clang/test/CIR/CodeGenOpenACC/enter-data.c | 29 ++++++++++++++++++++++
 mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp    | 21 +++++++++++-----
 2 files changed, 44 insertions(+), 6 deletions(-)

diff --git a/clang/test/CIR/CodeGenOpenACC/enter-data.c b/clang/test/CIR/CodeGenOpenACC/enter-data.c
index 136bd9690c4438..c0feda22a63c5c 100644
--- a/clang/test/CIR/CodeGenOpenACC/enter-data.c
+++ b/clang/test/CIR/CodeGenOpenACC/enter-data.c
@@ -123,3 +123,32 @@ void acc_data(int parmVar, int *ptrParmVar) {
   // CHECK-NEXT: acc.enter_data wait_devnum(%[[PARM_CAST]] : si32) wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)
 
 }
+
+void gh228290(int i) {
+  // CHECK: cir.func{{.*}}@gh228290(
+
+#pragma acc enter data copyin(i) wait wait
+// CHECK: acc.enter_data wait dataOperands({{.*}})
+
+#pragma acc enter data copyin(i) wait wait(1)
+// CHECK: acc.enter_data wait dataOperands({{.*}})
+
+#pragma acc enter data copyin(i) wait(1) wait
+// CHECK: acc.enter_data wait dataOperands({{.*}})
+
+#pragma acc enter data copyin(i) wait(1) wait(2)
+// CHECK: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32
+// CHECK: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
+// CHECK: %[[TWO_CAST:.*]] = cir.builtin_int_cast %[[TWO_CAST]] : !s32i -> si32
+// CHECK: acc.enter_data wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) dataOperands({{.*}})
+
+#pragma acc enter data copyin(i) wait(devnum:1: 1) wait(devnum:1: 2)
+// CHECK: FAIL
+//
+//#pragma acc enter data copyin(i) wait(devnum:1: 1) wait(devnum:2: 2)
+//  ;
+//  // CHECK: FAIL
+
+}
+
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 0d9ef4abc4bd50..aeefbbfd4132ff 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -4570,20 +4570,29 @@ void EnterDataOp::addAsyncOperand(
 void EnterDataOp::addWaitOnly(MLIRContext *context,
                               llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
   assert(effectiveDeviceTypes.empty());
-  assert(!getWaitAttr());
-  assert(getWaitOperands().empty());
-  assert(!getWaitDevnum());
+
+  if (getWaitAttr())
+    return;
 
   setWaitAttr(mlir::UnitAttr::get(context));
+
+  getWaitDevnumMutable().clear();
+  getWaitOperandsMutable().clear();
 }
 
 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 (getWaitAttr())
+    return;
+
+  // FIXME: At one point we need to figure out how to support multiple devnums
+  // here.  For now, assert.  Eventually we probably want to make dev-num and
+  // operands work in 'lock-step', so that getWaitDevnum().size() ==
+  // getWaitOperandsMutable().size().
+  assert(getWaitDevnum().empty() && "Merging devnum not yet implemented");
 
   // if hasDevnum, the first value is the devnum. The 'rest' go into the
   // operands list.