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

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clangir

Author: Erich Keane (erichkeane)

<details>
<summary>Changes</summary>

…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

---
Full diff: https://github.com/llvm/llvm-project/pull/228499.diff


2 Files Affected:

- (modified) clang/test/CIR/CodeGenOpenACC/enter-data.c (+29) 
- (modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+15-6) 


``````````diff
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.

``````````

</details>


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