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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
