[Mlir-commits] [mlir] [mlir][OpenACC] Preserve worker rows when combining reductions (PR #210804)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Mon Jul 20 21:00:41 PDT 2026
https://github.com/khaki3 updated https://github.com/llvm/llvm-project/pull/210804
>From 4c316621a473f21607795bafbd1ffa51a9dabe93 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 13:12:09 -0700
Subject: [PATCH 01/22] [OpenACC] Preserve worker rows when combining
reductions
Keep ThreadY active for worker-private sources so every worker row contributes, and reject mixed scopes that require incompatible predication.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 53 +++++++++--
...-worker-reduction-combine-mixed-scope.mlir | 43 +++++++++
...cc-cg-to-gpu-worker-reduction-combine.mlir | 95 +++++++++++++++++++
3 files changed, 181 insertions(+), 10 deletions(-)
create mode 100644 mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
create mode 100644 mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 24aeed6e23bda..6be79b0dd40e4 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1246,7 +1246,30 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
mlir::acc::GPUParallelDimAttr lowestParDim =
mlir::acc::GPUParallelDimAttr::threadXDim(ctx);
if (block) {
- block->walk([&](Operation *op) {
+ std::optional<bool> combineThreadYActive;
+ auto applyCombineParDims =
+ [&](Operation *combineOp, Value src,
+ ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
+ bool isWorkerPrivate =
+ getPrivateScopeForMemref(src) == PrivateMemScope::Worker;
+ for (mlir::acc::GPUParallelDimAttr parDim : combineParDims) {
+ if (parDim.isThreadY()) {
+ if (combineThreadYActive &&
+ *combineThreadYActive != isWorkerPrivate) {
+ combineOp->emitError()
+ << "mixed worker-private and non-worker-private reduction "
+ "combines require incompatible ThreadY predication";
+ hasFailed = true;
+ return failure();
+ }
+ combineThreadYActive = isWorkerPrivate;
+ } else {
+ mlir::acc::removeParDim(ancestorParDims, parDim);
+ }
+ }
+ return success();
+ };
+ block->walk([&](Operation *op) -> WalkResult {
// Check stores to acc.private_local - add the privatize's par_dims
// as active dims so predication is correct for per-worker/gang memory.
if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(op)) {
@@ -1277,17 +1300,17 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
// kernel and loop in combined constructs.
if (acc::ReductionCombineOp reductionCombineOp =
dyn_cast<acc::ReductionCombineOp>(op)) {
- for (mlir::acc::GPUParallelDimAttr parDim :
- getReductionCombineParDims(reductionCombineOp)) {
- mlir::acc::removeParDim(ancestorParDims, parDim);
- }
+ if (failed(applyCombineParDims(
+ reductionCombineOp, reductionCombineOp.getSrcMemref(),
+ getReductionCombineParDims(reductionCombineOp))))
+ return WalkResult::interrupt();
}
if (acc::ReductionCombineRegionOp combineRegionOp =
dyn_cast<acc::ReductionCombineRegionOp>(op)) {
- for (mlir::acc::GPUParallelDimAttr parDim :
- getReductionCombineParDims(combineRegionOp)) {
- mlir::acc::removeParDim(ancestorParDims, parDim);
- }
+ if (failed(applyCombineParDims(
+ combineRegionOp, combineRegionOp.getSrcVar(),
+ getReductionCombineParDims(combineRegionOp))))
+ return WalkResult::interrupt();
}
// An array accumulate reduces across its par_dims via gpu.all_reduce, so
// all those threads must execute it - treat them as active (unlike the
@@ -1301,6 +1324,14 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
return WalkResult::advance();
});
+ if (combineThreadYActive) {
+ mlir::acc::GPUParallelDimAttr threadY =
+ mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
+ if (*combineThreadYActive)
+ mlir::acc::insertParDim(ancestorParDims, threadY);
+ else
+ mlir::acc::removeParDim(ancestorParDims, threadY);
+ }
}
// Obtain launch dimensions
@@ -1747,7 +1778,7 @@ ACCCGToGPULowering::getPrivateMemScope(acc::PrivatizeOp privatizeOp) {
}
if (hasThreadX)
return PrivateMemScope::Thread;
- if (hasBlock && hasThreadY)
+ if (hasThreadY)
return PrivateMemScope::Worker;
if (hasBlock)
return PrivateMemScope::Gang;
@@ -1871,6 +1902,8 @@ void ACCCGToGPULowering::processPredicateRegion(
SmallVector<mlir::acc::GPUParallelDimAttr>>
parDimsPair = computeActiveAndInactiveParDims(
interOp, &interOp.getRegion().front());
+ if (hasFailed)
+ return;
// If ThreadY reduction exists, subgroup alignment is applied
// (blockDim.x = subgroupSize), so ThreadX lanes exist even without explicit
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
new file mode 100644
index 0000000000000..a0d381740c034
--- /dev/null
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -0,0 +1,43 @@
+// RUN: mlir-opt %s --pass-pipeline="builtin.module(func.func(acc-cg-to-gpu))" \
+// RUN: -verify-diagnostics
+
+func.func @mixed_scope_worker_reduction_combine(
+ %other: memref<i32>, %result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ // expected-error at +1 {{failed to legalize operation 'acc.compute_region' that was explicitly marked illegal}}
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %other_arg = %other,
+ %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c0_i32 = arith.constant 0 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
+ memref.store %c0_i32, %local[] : memref<i32>
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[thread_y]>}
+ acc.predicate_region {
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ // expected-error at +1 {{mixed worker-private and non-worker-private reduction combines require incompatible ThreadY predication}}
+ acc.reduction_combine %other_arg into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
new file mode 100644
index 0000000000000..b2fd246d33d1b
--- /dev/null
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -0,0 +1,95 @@
+// RUN: mlir-opt %s --pass-pipeline="builtin.module(func.func(acc-cg-to-gpu))" | FileCheck %s
+
+// A worker-private reduction has one shared slot per ThreadY row. The combine
+// must keep ThreadY active and predicate only ThreadX.
+
+// CHECK-LABEL: func.func @worker_reduction_combine
+// CHECK: gpu.launch {{.*}} threads([[TID_X:%[^,]+]], [[TID_Y:%[^,]+]],
+// CHECK-NOT: arith.cmpi eq, [[TID_Y]]
+// CHECK: %[[IS_X_ZERO:.*]] = arith.cmpi eq, [[TID_X]],
+// CHECK-NOT: arith.andi
+// CHECK: scf.if %[[IS_X_ZERO]]
+// CHECK: acc.atomic.update
+
+func.func @worker_reduction_combine(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c0_i32 = arith.constant 0 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
+ memref.store %c0_i32, %local[] : memref<i32>
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[thread_y]>}
+ acc.predicate_region {
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @worker_reduction_combine_region
+// CHECK: gpu.launch {{.*}} threads([[REGION_TID_X:%[^,]+]], [[REGION_TID_Y:%[^,]+]],
+// CHECK-NOT: arith.cmpi eq, [[REGION_TID_Y]]
+// CHECK: %[[REGION_IS_X_ZERO:.*]] = arith.cmpi eq, [[REGION_TID_X]],
+// CHECK-NOT: arith.andi
+// CHECK: scf.if %[[REGION_IS_X_ZERO]]
+// CHECK: arith.addi
+
+func.func @worker_reduction_combine_region(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c0_i32 = arith.constant 0 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
+ memref.store %c0_i32, %local[] : memref<i32>
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[thread_y]>}
+ acc.predicate_region {
+ acc.reduction_combine_region %local into %result_arg : memref<i32> {
+ %lhs = memref.load %result_arg[] : memref<i32>
+ %rhs = memref.load %local[] : memref<i32>
+ %sum = arith.addi %lhs, %rhs : i32
+ memref.store %sum, %result_arg[] : memref<i32>
+ acc.yield
+ } {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
>From d763e8065946687750b2c78fcf5393167b7de962 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 15:08:39 -0700
Subject: [PATCH 02/22] [OpenACC] Report mixed reduction scopes as NYI
Use the standard OpenACC unsupported-feature diagnostic for incompatible ThreadY predication.
---
mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 7 ++++---
1 file changed, 4 insertions(+), 3 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 6be79b0dd40e4..378fa17b24850 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1256,9 +1256,10 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
if (parDim.isThreadY()) {
if (combineThreadYActive &&
*combineThreadYActive != isWorkerPrivate) {
- combineOp->emitError()
- << "mixed worker-private and non-worker-private reduction "
- "combines require incompatible ThreadY predication";
+ (void)accSupport.emitNYI(
+ combineOp->getLoc(),
+ "mixed worker-private and non-worker-private reduction "
+ "combines require incompatible ThreadY predication");
hasFailed = true;
return failure();
}
>From 1fe76c2f1ea14bd036598286c2d9c0ea8d2291ef Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 16:50:59 -0700
Subject: [PATCH 03/22] [OpenACC] Isolate ThreadY predicate requirements
Propagate nested worker requirements without broadening unrelated side effects, and reject incompatible operations sharing one predicate region.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 126 +++++++++++++++---
...-worker-reduction-combine-mixed-scope.mlir | 40 +++++-
...cc-cg-to-gpu-worker-reduction-combine.mlir | 46 +++++++
3 files changed, 189 insertions(+), 23 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 378fa17b24850..34131dc3753a3 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1246,27 +1246,111 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
mlir::acc::GPUParallelDimAttr lowestParDim =
mlir::acc::GPUParallelDimAttr::threadXDim(ctx);
if (block) {
- std::optional<bool> combineThreadYActive;
+ struct ThreadYRequirement {
+ bool active = false;
+ bool inactive = false;
+ };
+ auto analyzeThreadYRequirements =
+ [&](auto &&self,
+ Block &predicateBlock) -> FailureOr<ThreadYRequirement> {
+ ThreadYRequirement requirement;
+ auto requireThreadY = [&](Operation *op, bool active) {
+ if ((active && requirement.inactive) ||
+ (!active && requirement.active)) {
+ (void)accSupport.emitNYI(
+ op->getLoc(),
+ "operations in the same predicate region require incompatible "
+ "ThreadY predication");
+ hasFailed = true;
+ return failure();
+ }
+ requirement.active |= active;
+ requirement.inactive |= !active;
+ return success();
+ };
+ auto applyCombineRequirement =
+ [&](Operation *combineOp, Value src,
+ ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
+ if (llvm::none_of(combineParDims,
+ [](auto parDim) { return parDim.isThreadY(); }))
+ return success();
+ bool isWorkerPrivate =
+ getPrivateScopeForMemref(src) == PrivateMemScope::Worker;
+ return requireThreadY(combineOp, isWorkerPrivate);
+ };
+
+ for (Operation &nestedOp : predicateBlock) {
+ if (acc::PredicateRegionOp nestedPredicate =
+ dyn_cast<acc::PredicateRegionOp>(nestedOp)) {
+ FailureOr<ThreadYRequirement> nestedRequirement =
+ self(self, nestedPredicate.getRegion().front());
+ if (failed(nestedRequirement))
+ return failure();
+ // An active descendant must not be excluded by this region's
+ // predicate. Inactive descendants apply their own predicate.
+ if (nestedRequirement->active &&
+ failed(requireThreadY(&nestedOp, /*active=*/true)))
+ return failure();
+ continue;
+ }
+ if (acc::ReductionCombineOp combineOp =
+ dyn_cast<acc::ReductionCombineOp>(nestedOp)) {
+ if (failed(applyCombineRequirement(
+ combineOp, combineOp.getSrcMemref(),
+ getReductionCombineParDims(combineOp))))
+ return failure();
+ continue;
+ }
+ if (acc::ReductionCombineRegionOp combineRegionOp =
+ dyn_cast<acc::ReductionCombineRegionOp>(nestedOp)) {
+ if (failed(applyCombineRequirement(
+ combineRegionOp, combineRegionOp.getSrcVar(),
+ getReductionCombineParDims(combineRegionOp))))
+ return failure();
+ continue;
+ }
+ if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(nestedOp)) {
+ bool threadYActive = false;
+ if (acc::PrivateLocalOp privateLocal =
+ storeOp.getMemref().getDefiningOp<acc::PrivateLocalOp>()) {
+ if (acc::GPUParallelDimsAttr parDims =
+ getPrivatizeOp(privateLocal, computeRegion)
+ .getParDimsAttr()) {
+ threadYActive = llvm::any_of(parDims.getArray(), [](auto parDim) {
+ return parDim.isThreadY();
+ });
+ }
+ }
+ if (failed(requireThreadY(storeOp, threadYActive)))
+ return failure();
+ continue;
+ }
+ if (acc::ReductionAccumulateArrayOp accumulateArrayOp =
+ dyn_cast<acc::ReductionAccumulateArrayOp>(nestedOp)) {
+ bool threadYActive =
+ llvm::any_of(accumulateArrayOp.getParDims().getArray(),
+ [](auto parDim) { return parDim.isThreadY(); });
+ if (failed(requireThreadY(accumulateArrayOp, threadYActive)))
+ return failure();
+ continue;
+ }
+ if (!isMemoryEffectFree(&nestedOp) &&
+ failed(requireThreadY(&nestedOp, /*active=*/false)))
+ return failure();
+ }
+ return requirement;
+ };
+
+ FailureOr<ThreadYRequirement> threadYRequirement =
+ analyzeThreadYRequirements(analyzeThreadYRequirements, *block);
+ if (failed(threadYRequirement))
+ return {};
+
auto applyCombineParDims =
- [&](Operation *combineOp, Value src,
- ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
- bool isWorkerPrivate =
- getPrivateScopeForMemref(src) == PrivateMemScope::Worker;
+ [&](ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
for (mlir::acc::GPUParallelDimAttr parDim : combineParDims) {
- if (parDim.isThreadY()) {
- if (combineThreadYActive &&
- *combineThreadYActive != isWorkerPrivate) {
- (void)accSupport.emitNYI(
- combineOp->getLoc(),
- "mixed worker-private and non-worker-private reduction "
- "combines require incompatible ThreadY predication");
- hasFailed = true;
- return failure();
- }
- combineThreadYActive = isWorkerPrivate;
- } else {
+ if (!parDim.isThreadY())
mlir::acc::removeParDim(ancestorParDims, parDim);
- }
}
return success();
};
@@ -1302,14 +1386,12 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
if (acc::ReductionCombineOp reductionCombineOp =
dyn_cast<acc::ReductionCombineOp>(op)) {
if (failed(applyCombineParDims(
- reductionCombineOp, reductionCombineOp.getSrcMemref(),
getReductionCombineParDims(reductionCombineOp))))
return WalkResult::interrupt();
}
if (acc::ReductionCombineRegionOp combineRegionOp =
dyn_cast<acc::ReductionCombineRegionOp>(op)) {
if (failed(applyCombineParDims(
- combineRegionOp, combineRegionOp.getSrcVar(),
getReductionCombineParDims(combineRegionOp))))
return WalkResult::interrupt();
}
@@ -1325,10 +1407,10 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
return WalkResult::advance();
});
- if (combineThreadYActive) {
+ if (threadYRequirement->active || threadYRequirement->inactive) {
mlir::acc::GPUParallelDimAttr threadY =
mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
- if (*combineThreadYActive)
+ if (threadYRequirement->active)
mlir::acc::insertParDim(ancestorParDims, threadY);
else
mlir::acc::removeParDim(ancestorParDims, threadY);
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index a0d381740c034..e80b5afa26d20 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -30,7 +30,7 @@ func.func @mixed_scope_worker_reduction_combine(
acc.predicate_region {
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
- // expected-error at +1 {{mixed worker-private and non-worker-private reduction combines require incompatible ThreadY predication}}
+ // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
acc.reduction_combine %other_arg into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
@@ -41,3 +41,41 @@ func.func @mixed_scope_worker_reduction_combine(
}
return
}
+
+func.func @worker_combine_with_single_store(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ // expected-error at +1 {{failed to legalize operation 'acc.compute_region' that was explicitly marked illegal}}
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c7_i32 = arith.constant 7 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
+ memref.store %c7_i32, %local[] : memref<i32>
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[thread_y]>}
+ acc.predicate_region {
+ memref.store %c7_i32, %result_arg[] : memref<i32>
+ // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index b2fd246d33d1b..afd6108d6e670 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -93,3 +93,49 @@ func.func @worker_reduction_combine_region(%result: memref<i32>) {
}
return
}
+
+// Nested predicate regions choose their ThreadY predicates independently. The
+// outer region must keep ThreadY active so it does not exclude worker rows
+// before the nested worker-private combine is reached.
+func.func @nested_worker_reduction_combines(
+ %other: memref<i32>, %result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %other_arg = %other,
+ %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c0_i32 = arith.constant 0 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
+ memref.store %c0_i32, %local[] : memref<i32>
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[thread_y]>}
+ acc.predicate_region {
+ acc.predicate_region {
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ acc.predicate_region {
+ acc.reduction_combine %other_arg into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
>From 477f10e0472195d6bb55598bf364a1655682ab6f Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 17:04:51 -0700
Subject: [PATCH 04/22] [OpenACC] Propagate nested ThreadY requirements
Analyze control-flow regions recursively and recognize aliased worker-private stores so their worker rows remain active.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 32 ++++---
...cc-cg-to-gpu-worker-reduction-combine.mlir | 86 +++++++++++++++++++
2 files changed, 107 insertions(+), 11 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 34131dc3753a3..98c697f52668e 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1310,17 +1310,8 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
continue;
}
if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(nestedOp)) {
- bool threadYActive = false;
- if (acc::PrivateLocalOp privateLocal =
- storeOp.getMemref().getDefiningOp<acc::PrivateLocalOp>()) {
- if (acc::GPUParallelDimsAttr parDims =
- getPrivatizeOp(privateLocal, computeRegion)
- .getParDimsAttr()) {
- threadYActive = llvm::any_of(parDims.getArray(), [](auto parDim) {
- return parDim.isThreadY();
- });
- }
- }
+ bool threadYActive = getPrivateScopeForMemref(storeOp.getMemref()) ==
+ PrivateMemScope::Worker;
if (failed(requireThreadY(storeOp, threadYActive)))
return failure();
continue;
@@ -1334,6 +1325,25 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
return failure();
continue;
}
+ if (nestedOp.getNumRegions() != 0) {
+ for (Region ®ion : nestedOp.getRegions()) {
+ for (Block &nestedBlock : region) {
+ FailureOr<ThreadYRequirement> nestedRequirement =
+ self(self, nestedBlock);
+ if (failed(nestedRequirement))
+ return failure();
+ // Other region-bearing operations do not introduce independent
+ // ACC predication, so both requirements apply to this region.
+ if (nestedRequirement->active &&
+ failed(requireThreadY(&nestedOp, /*active=*/true)))
+ return failure();
+ if (nestedRequirement->inactive &&
+ failed(requireThreadY(&nestedOp, /*active=*/false)))
+ return failure();
+ }
+ }
+ continue;
+ }
if (!isMemoryEffectFree(&nestedOp) &&
failed(requireThreadY(&nestedOp, /*active=*/false)))
return failure();
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index afd6108d6e670..8a2dd6bd470f8 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -97,6 +97,12 @@ func.func @worker_reduction_combine_region(%result: memref<i32>) {
// Nested predicate regions choose their ThreadY predicates independently. The
// outer region must keep ThreadY active so it does not exclude worker rows
// before the nested worker-private combine is reached.
+// CHECK-LABEL: func.func @nested_worker_reduction_combines
+// CHECK: gpu.launch {{.*}} threads([[NESTED_TX:%[^,]+]], [[NESTED_TY:%[^,]+]],
+// CHECK-NOT: arith.cmpi eq, [[NESTED_TY]]
+// CHECK: %[[NESTED_TX_ZERO:.*]] = arith.cmpi eq, [[NESTED_TX]],
+// CHECK-NOT: arith.andi
+// CHECK: scf.if %[[NESTED_TX_ZERO]]
func.func @nested_worker_reduction_combines(
%other: memref<i32>, %result: memref<i32>) {
%c1 = arith.constant 1 : index
@@ -139,3 +145,83 @@ func.func @nested_worker_reduction_combines(
}
return
}
+
+// CHECK-LABEL: func.func @worker_combine_in_scf_if
+// CHECK: gpu.launch {{.*}} threads([[IF_TX:%[^,]+]], [[IF_TY:%[^,]+]],
+// CHECK-NOT: arith.cmpi eq, [[IF_TY]]
+// CHECK: %[[IF_TX_ZERO:.*]] = arith.cmpi eq, [[IF_TX]],
+// CHECK-NOT: arith.andi
+// CHECK: scf.if %[[IF_TX_ZERO]]
+func.func @worker_combine_in_scf_if(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %true = arith.constant true
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ acc.predicate_region {
+ scf.if %true {
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
+
+// CHECK-LABEL: func.func @aliased_worker_store
+// CHECK: gpu.launch {{.*}} threads([[ALIAS_TX:%[^,]+]], [[ALIAS_TY:%[^,]+]],
+// CHECK-NOT: arith.cmpi eq, [[ALIAS_TY]]
+// CHECK: %[[ALIAS_TX_ZERO:.*]] = arith.cmpi eq, [[ALIAS_TX]],
+// CHECK-NOT: arith.andi
+// CHECK: scf.if %[[ALIAS_TX_ZERO]]
+// CHECK: memref.store
+// CHECK: acc.atomic.update
+func.func @aliased_worker_store(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c7_i32 = arith.constant 7 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ %cast = memref.cast %local : memref<i32> to memref<i32>
+ acc.predicate_region {
+ memref.store %c7_i32, %cast[] : memref<i32>
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
>From 2e234b5cb4dac8aa7cf1ab07b63984b7cb9f329f Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 17:16:55 -0700
Subject: [PATCH 05/22] [OpenACC] Account for predicate operation effects
Preserve all privatized ThreadY dimensions through aliases and reject intrinsically effecting region operations that cannot share worker predication.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 33 +++++++++++-----
...-worker-reduction-combine-mixed-scope.mlir | 38 +++++++++++++++++++
...cc-cg-to-gpu-worker-reduction-combine.mlir | 36 ++++++++++++++++++
3 files changed, 97 insertions(+), 10 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 98c697f52668e..20e8977c907f3 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1268,15 +1268,22 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
requirement.inactive |= !active;
return success();
};
+ auto hasThreadYPrivateStorage = [&](Value memref) {
+ acc::PrivatizeOp privatize = getPrivatizeForMemref(memref);
+ if (!privatize)
+ return false;
+ acc::GPUParallelDimsAttr parDims = privatize.getParDimsAttr();
+ return parDims && llvm::any_of(parDims.getArray(), [](auto parDim) {
+ return parDim.isThreadY();
+ });
+ };
auto applyCombineRequirement =
[&](Operation *combineOp, Value src,
ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
if (llvm::none_of(combineParDims,
[](auto parDim) { return parDim.isThreadY(); }))
return success();
- bool isWorkerPrivate =
- getPrivateScopeForMemref(src) == PrivateMemScope::Worker;
- return requireThreadY(combineOp, isWorkerPrivate);
+ return requireThreadY(combineOp, hasThreadYPrivateStorage(src));
};
for (Operation &nestedOp : predicateBlock) {
@@ -1310,9 +1317,8 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
continue;
}
if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(nestedOp)) {
- bool threadYActive = getPrivateScopeForMemref(storeOp.getMemref()) ==
- PrivateMemScope::Worker;
- if (failed(requireThreadY(storeOp, threadYActive)))
+ if (failed(requireThreadY(
+ storeOp, hasThreadYPrivateStorage(storeOp.getMemref()))))
return failure();
continue;
}
@@ -1326,6 +1332,15 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
continue;
}
if (nestedOp.getNumRegions() != 0) {
+ bool hasOwnEffects = true;
+ if (auto effectOp = dyn_cast<MemoryEffectOpInterface>(&nestedOp)) {
+ hasOwnEffects = !effectOp.hasNoEffect();
+ } else if (nestedOp.hasTrait<OpTrait::HasRecursiveMemoryEffects>()) {
+ hasOwnEffects = false;
+ }
+ if (hasOwnEffects &&
+ failed(requireThreadY(&nestedOp, /*active=*/false)))
+ return failure();
for (Region ®ion : nestedOp.getRegions()) {
for (Block &nestedBlock : region) {
FailureOr<ThreadYRequirement> nestedRequirement =
@@ -1368,10 +1383,8 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
// Check stores to acc.private_local - add the privatize's par_dims
// as active dims so predication is correct for per-worker/gang memory.
if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(op)) {
- if (auto privateLocalOp =
- storeOp.getMemref().getDefiningOp<acc::PrivateLocalOp>()) {
- acc::PrivatizeOp privatizeOp =
- getPrivatizeOp(privateLocalOp, computeRegion);
+ if (acc::PrivatizeOp privatizeOp =
+ getPrivatizeForMemref(storeOp.getMemref())) {
if (mlir::acc::GPUParallelDimsAttr parDimsAttr =
privatizeOp.getParDimsAttr()) {
for (auto parDim : parDimsAttr.getArray())
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index e80b5afa26d20..19e1b3106d966 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -79,3 +79,41 @@ func.func @worker_combine_with_single_store(%result: memref<i32>) {
}
return
}
+
+func.func @worker_combine_with_atomic_update(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ // expected-error at +1 {{failed to legalize operation 'acc.compute_region' that was explicitly marked illegal}}
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c1_i32 = arith.constant 1 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ acc.predicate_region {
+ acc.atomic.update %result_arg : memref<i32> {
+ ^bb0(%current: i32):
+ %next = arith.addi %current, %c1_i32 : i32
+ acc.yield %next : i32
+ }
+ // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index 8a2dd6bd470f8..ba8f16d3966df 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -225,3 +225,39 @@ func.func @aliased_worker_store(%result: memref<i32>) {
}
return
}
+
+// CHECK-LABEL: func.func @aliased_thread_store
+// CHECK: gpu.launch {{.*}} threads([[THREAD_TX:%[^,]+]], [[THREAD_TY:%[^,]+]],
+// CHECK-NOT: arith.cmpi eq, [[THREAD_TX]]
+// CHECK-NOT: arith.cmpi eq, [[THREAD_TY]]
+// CHECK: memref.store
+func.func @aliased_thread_store() {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y, thread_x]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private)
+ : (!acc.private_type<memref<i32>>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c7_i32 = arith.constant 7 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ %cast = memref.cast %local : memref<i32> to memref<i32>
+ acc.predicate_region {
+ memref.store %c7_i32, %cast[] : memref<i32>
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
>From ae0d329822ede9a6d80731ebd45abd7e520ef5ae Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 17:26:21 -0700
Subject: [PATCH 06/22] [OpenACC] Require proven private aliases
Follow only alias-preserving view operations and treat every non-private combine as ThreadY-inactive to avoid broadening ambiguous global effects.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 45 +++++++++----------
...-worker-reduction-combine-mixed-scope.mlir | 6 ++-
2 files changed, 25 insertions(+), 26 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 20e8977c907f3..9e9114e9aa174 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1277,14 +1277,9 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
return parDim.isThreadY();
});
};
- auto applyCombineRequirement =
- [&](Operation *combineOp, Value src,
- ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
- if (llvm::none_of(combineParDims,
- [](auto parDim) { return parDim.isThreadY(); }))
- return success();
- return requireThreadY(combineOp, hasThreadYPrivateStorage(src));
- };
+ auto applyCombineRequirement = [&](Operation *combineOp, Value src) {
+ return requireThreadY(combineOp, hasThreadYPrivateStorage(src));
+ };
for (Operation &nestedOp : predicateBlock) {
if (acc::PredicateRegionOp nestedPredicate =
@@ -1302,17 +1297,15 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
if (acc::ReductionCombineOp combineOp =
dyn_cast<acc::ReductionCombineOp>(nestedOp)) {
- if (failed(applyCombineRequirement(
- combineOp, combineOp.getSrcMemref(),
- getReductionCombineParDims(combineOp))))
+ if (failed(
+ applyCombineRequirement(combineOp, combineOp.getSrcMemref())))
return failure();
continue;
}
if (acc::ReductionCombineRegionOp combineRegionOp =
dyn_cast<acc::ReductionCombineRegionOp>(nestedOp)) {
- if (failed(applyCombineRequirement(
- combineRegionOp, combineRegionOp.getSrcVar(),
- getReductionCombineParDims(combineRegionOp))))
+ if (failed(applyCombineRequirement(combineRegionOp,
+ combineRegionOp.getSrcVar())))
return failure();
continue;
}
@@ -1893,18 +1886,22 @@ ACCCGToGPULowering::getPrivateMemScope(acc::PrivatizeOp privatizeOp) {
/// Walks back from a memref use to its defining `acc.private_local`, if any.
static acc::PrivateLocalOp getPrivateLocalForMemref(Value memref) {
- llvm::SmallVector<Value, 8> worklist{memref};
- llvm::SmallPtrSet<Value, 8> seen;
- while (!worklist.empty()) {
- Value v = worklist.pop_back_val();
- if (!seen.insert(v).second)
- continue;
- Operation *def = v.getDefiningOp();
- if (!def)
- continue;
+ Value current = memref;
+ while (Operation *def = current.getDefiningOp()) {
if (acc::PrivateLocalOp privateLocal = dyn_cast<acc::PrivateLocalOp>(def))
return privateLocal;
- worklist.append(def->getOperands().begin(), def->getOperands().end());
+ if (ViewLikeOpInterface viewLike = dyn_cast<ViewLikeOpInterface>(def)) {
+ current = viewLike.getViewSource();
+ continue;
+ }
+ if (acc::PartialEntityAccessOpInterface partialAccess =
+ dyn_cast<acc::PartialEntityAccessOpInterface>(def)) {
+ current = partialAccess.getBaseEntity();
+ continue;
+ }
+ // Do not follow arbitrary operands: multi-source operations such as
+ // arith.select do not prove that their result aliases private storage.
+ return nullptr;
}
return nullptr;
}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index 19e1b3106d966..1275e4c77890a 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -32,7 +32,7 @@ func.func @mixed_scope_worker_reduction_combine(
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
// expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
acc.reduction_combine %other_arg into %result_arg <add> : memref<i32>
- {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ {acc.par_dims = #acc<par_dims[block_y, thread_x]>}
}
scf.reduce
} {acc.par_dims = #acc<par_dims[block_y]>}
@@ -59,15 +59,17 @@ func.func @worker_combine_with_single_store(%result: memref<i32>) {
%c0 = arith.constant 0 : index
%c1_inner = arith.constant 1 : index
%c7_i32 = arith.constant 7 : i32
+ %false = arith.constant false
scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
%local = acc.private_local %private_arg
: (!acc.private_type<memref<i32>>) -> memref<i32>
+ %selected = arith.select %false, %local, %result_arg : memref<i32>
scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
memref.store %c7_i32, %local[] : memref<i32>
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
- memref.store %c7_i32, %result_arg[] : memref<i32>
+ memref.store %c7_i32, %selected[] : memref<i32>
// expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
>From 3ed1de6f3b95aee6516db811bbd62dc4c54f9335 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 17:38:25 -0700
Subject: [PATCH 07/22] [OpenACC] Restrict ThreadY predicate broadening
Preserve legacy predicates unless a proven worker combine requires ThreadY, and reject broadening around writes, unknown effects, or incompatible combines.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 223 +++++++++---------
...cc-cg-to-gpu-worker-reduction-combine.mlir | 78 +-----
2 files changed, 107 insertions(+), 194 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 9e9114e9aa174..50635dde561db 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1246,138 +1246,122 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
mlir::acc::GPUParallelDimAttr lowestParDim =
mlir::acc::GPUParallelDimAttr::threadXDim(ctx);
if (block) {
- struct ThreadYRequirement {
- bool active = false;
- bool inactive = false;
+ struct ThreadYBroadeningInfo {
+ bool hasActiveWorkerCombine = false;
+ bool hasExplicitInactiveCombine = false;
+ bool hasBroadeningConflict = false;
+ Operation *diagnosticOp = nullptr;
};
- auto analyzeThreadYRequirements =
- [&](auto &&self,
- Block &predicateBlock) -> FailureOr<ThreadYRequirement> {
- ThreadYRequirement requirement;
- auto requireThreadY = [&](Operation *op, bool active) {
- if ((active && requirement.inactive) ||
- (!active && requirement.active)) {
- (void)accSupport.emitNYI(
- op->getLoc(),
- "operations in the same predicate region require incompatible "
- "ThreadY predication");
- hasFailed = true;
- return failure();
+ auto mergeBroadeningInfo = [](ThreadYBroadeningInfo &dst,
+ const ThreadYBroadeningInfo &src) {
+ dst.hasActiveWorkerCombine |= src.hasActiveWorkerCombine;
+ dst.hasExplicitInactiveCombine |= src.hasExplicitInactiveCombine;
+ dst.hasBroadeningConflict |= src.hasBroadeningConflict;
+ if (!dst.diagnosticOp)
+ dst.diagnosticOp = src.diagnosticOp;
+ };
+ auto isProvenWorkerPrivate = [&](Value memref) {
+ Value current = memref;
+ while (Operation *def = current.getDefiningOp()) {
+ if (acc::PrivateLocalOp privateLocal =
+ dyn_cast<acc::PrivateLocalOp>(def)) {
+ return getPrivateMemScope(getPrivatizeOp(
+ privateLocal, computeRegion)) == PrivateMemScope::Worker;
+ }
+ if (ViewLikeOpInterface viewLike = dyn_cast<ViewLikeOpInterface>(def)) {
+ current = viewLike.getViewSource();
+ continue;
+ }
+ break;
+ }
+ return false;
+ };
+ auto hasUnsafeOwnEffects = [](Operation *op) {
+ if (auto effectOp = dyn_cast<MemoryEffectOpInterface>(op)) {
+ SmallVector<MemoryEffects::EffectInstance> effects;
+ effectOp.getEffects(effects);
+ return llvm::any_of(effects, [](const auto &effect) {
+ return !isa<MemoryEffects::Read>(effect.getEffect());
+ });
+ }
+ return !op->hasTrait<OpTrait::HasRecursiveMemoryEffects>();
+ };
+ auto analyzeThreadYBroadening =
+ [&](auto &&self, Block &predicateBlock) -> ThreadYBroadeningInfo {
+ ThreadYBroadeningInfo info;
+ auto classifyCombine = [&](Operation *combineOp, Value src,
+ ArrayRef<GPUParallelDimAttr> parDims) {
+ bool hasThreadY = llvm::any_of(
+ parDims, [](auto parDim) { return parDim.isThreadY(); });
+ if (hasThreadY && isProvenWorkerPrivate(src)) {
+ info.hasActiveWorkerCombine = true;
+ } else {
+ info.hasExplicitInactiveCombine = true;
+ if (!info.diagnosticOp)
+ info.diagnosticOp = combineOp;
}
- requirement.active |= active;
- requirement.inactive |= !active;
- return success();
- };
- auto hasThreadYPrivateStorage = [&](Value memref) {
- acc::PrivatizeOp privatize = getPrivatizeForMemref(memref);
- if (!privatize)
- return false;
- acc::GPUParallelDimsAttr parDims = privatize.getParDimsAttr();
- return parDims && llvm::any_of(parDims.getArray(), [](auto parDim) {
- return parDim.isThreadY();
- });
- };
- auto applyCombineRequirement = [&](Operation *combineOp, Value src) {
- return requireThreadY(combineOp, hasThreadYPrivateStorage(src));
};
for (Operation &nestedOp : predicateBlock) {
if (acc::PredicateRegionOp nestedPredicate =
dyn_cast<acc::PredicateRegionOp>(nestedOp)) {
- FailureOr<ThreadYRequirement> nestedRequirement =
+ ThreadYBroadeningInfo nestedInfo =
self(self, nestedPredicate.getRegion().front());
- if (failed(nestedRequirement))
- return failure();
- // An active descendant must not be excluded by this region's
- // predicate. Inactive descendants apply their own predicate.
- if (nestedRequirement->active &&
- failed(requireThreadY(&nestedOp, /*active=*/true)))
- return failure();
+ // A worker-active descendant must reach its own predicate. Other
+ // requirements are enforced within the nested predicate region.
+ info.hasActiveWorkerCombine |= nestedInfo.hasActiveWorkerCombine;
continue;
}
if (acc::ReductionCombineOp combineOp =
dyn_cast<acc::ReductionCombineOp>(nestedOp)) {
- if (failed(
- applyCombineRequirement(combineOp, combineOp.getSrcMemref())))
- return failure();
+ classifyCombine(combineOp, combineOp.getSrcMemref(),
+ getReductionCombineParDims(combineOp));
continue;
}
if (acc::ReductionCombineRegionOp combineRegionOp =
dyn_cast<acc::ReductionCombineRegionOp>(nestedOp)) {
- if (failed(applyCombineRequirement(combineRegionOp,
- combineRegionOp.getSrcVar())))
- return failure();
- continue;
- }
- if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(nestedOp)) {
- if (failed(requireThreadY(
- storeOp, hasThreadYPrivateStorage(storeOp.getMemref()))))
- return failure();
- continue;
- }
- if (acc::ReductionAccumulateArrayOp accumulateArrayOp =
- dyn_cast<acc::ReductionAccumulateArrayOp>(nestedOp)) {
- bool threadYActive =
- llvm::any_of(accumulateArrayOp.getParDims().getArray(),
- [](auto parDim) { return parDim.isThreadY(); });
- if (failed(requireThreadY(accumulateArrayOp, threadYActive)))
- return failure();
+ classifyCombine(combineRegionOp, combineRegionOp.getSrcVar(),
+ getReductionCombineParDims(combineRegionOp));
continue;
}
if (nestedOp.getNumRegions() != 0) {
- bool hasOwnEffects = true;
- if (auto effectOp = dyn_cast<MemoryEffectOpInterface>(&nestedOp)) {
- hasOwnEffects = !effectOp.hasNoEffect();
- } else if (nestedOp.hasTrait<OpTrait::HasRecursiveMemoryEffects>()) {
- hasOwnEffects = false;
+ if (hasUnsafeOwnEffects(&nestedOp)) {
+ info.hasBroadeningConflict = true;
+ if (!info.diagnosticOp)
+ info.diagnosticOp = &nestedOp;
}
- if (hasOwnEffects &&
- failed(requireThreadY(&nestedOp, /*active=*/false)))
- return failure();
for (Region ®ion : nestedOp.getRegions()) {
- for (Block &nestedBlock : region) {
- FailureOr<ThreadYRequirement> nestedRequirement =
- self(self, nestedBlock);
- if (failed(nestedRequirement))
- return failure();
- // Other region-bearing operations do not introduce independent
- // ACC predication, so both requirements apply to this region.
- if (nestedRequirement->active &&
- failed(requireThreadY(&nestedOp, /*active=*/true)))
- return failure();
- if (nestedRequirement->inactive &&
- failed(requireThreadY(&nestedOp, /*active=*/false)))
- return failure();
- }
+ for (Block &nestedBlock : region)
+ mergeBroadeningInfo(info, self(self, nestedBlock));
}
continue;
}
- if (!isMemoryEffectFree(&nestedOp) &&
- failed(requireThreadY(&nestedOp, /*active=*/false)))
- return failure();
+ if (hasUnsafeOwnEffects(&nestedOp)) {
+ info.hasBroadeningConflict = true;
+ if (!info.diagnosticOp)
+ info.diagnosticOp = &nestedOp;
+ }
}
- return requirement;
+ return info;
};
- FailureOr<ThreadYRequirement> threadYRequirement =
- analyzeThreadYRequirements(analyzeThreadYRequirements, *block);
- if (failed(threadYRequirement))
- return {};
+ ThreadYBroadeningInfo threadYInfo =
+ analyzeThreadYBroadening(analyzeThreadYBroadening, *block);
auto applyCombineParDims =
[&](ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
- for (mlir::acc::GPUParallelDimAttr parDim : combineParDims) {
- if (!parDim.isThreadY())
- mlir::acc::removeParDim(ancestorParDims, parDim);
- }
+ for (mlir::acc::GPUParallelDimAttr parDim : combineParDims)
+ mlir::acc::removeParDim(ancestorParDims, parDim);
return success();
};
block->walk([&](Operation *op) -> WalkResult {
// Check stores to acc.private_local - add the privatize's par_dims
// as active dims so predication is correct for per-worker/gang memory.
if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(op)) {
- if (acc::PrivatizeOp privatizeOp =
- getPrivatizeForMemref(storeOp.getMemref())) {
+ if (auto privateLocalOp =
+ storeOp.getMemref().getDefiningOp<acc::PrivateLocalOp>()) {
+ acc::PrivatizeOp privatizeOp =
+ getPrivatizeOp(privateLocalOp, computeRegion);
if (mlir::acc::GPUParallelDimsAttr parDimsAttr =
privatizeOp.getParDimsAttr()) {
for (auto parDim : parDimsAttr.getArray())
@@ -1423,13 +1407,22 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
return WalkResult::advance();
});
- if (threadYRequirement->active || threadYRequirement->inactive) {
- mlir::acc::GPUParallelDimAttr threadY =
- mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
- if (threadYRequirement->active)
- mlir::acc::insertParDim(ancestorParDims, threadY);
- else
- mlir::acc::removeParDim(ancestorParDims, threadY);
+ mlir::acc::GPUParallelDimAttr threadY =
+ mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
+ bool baselineThreadYActive = llvm::is_contained(ancestorParDims, threadY);
+ if (threadYInfo.hasActiveWorkerCombine && !baselineThreadYActive) {
+ if (threadYInfo.hasExplicitInactiveCombine ||
+ threadYInfo.hasBroadeningConflict) {
+ Operation *diagnosticOp =
+ threadYInfo.diagnosticOp ? threadYInfo.diagnosticOp : op;
+ (void)accSupport.emitNYI(
+ diagnosticOp->getLoc(),
+ "operations in the same predicate region require incompatible "
+ "ThreadY predication");
+ hasFailed = true;
+ return {};
+ }
+ mlir::acc::insertParDim(ancestorParDims, threadY);
}
}
@@ -1886,22 +1879,18 @@ ACCCGToGPULowering::getPrivateMemScope(acc::PrivatizeOp privatizeOp) {
/// Walks back from a memref use to its defining `acc.private_local`, if any.
static acc::PrivateLocalOp getPrivateLocalForMemref(Value memref) {
- Value current = memref;
- while (Operation *def = current.getDefiningOp()) {
- if (acc::PrivateLocalOp privateLocal = dyn_cast<acc::PrivateLocalOp>(def))
- return privateLocal;
- if (ViewLikeOpInterface viewLike = dyn_cast<ViewLikeOpInterface>(def)) {
- current = viewLike.getViewSource();
+ llvm::SmallVector<Value, 8> worklist{memref};
+ llvm::SmallPtrSet<Value, 8> seen;
+ while (!worklist.empty()) {
+ Value v = worklist.pop_back_val();
+ if (!seen.insert(v).second)
continue;
- }
- if (acc::PartialEntityAccessOpInterface partialAccess =
- dyn_cast<acc::PartialEntityAccessOpInterface>(def)) {
- current = partialAccess.getBaseEntity();
+ Operation *def = v.getDefiningOp();
+ if (!def)
continue;
- }
- // Do not follow arbitrary operands: multi-source operations such as
- // arith.select do not prove that their result aliases private storage.
- return nullptr;
+ if (acc::PrivateLocalOp privateLocal = dyn_cast<acc::PrivateLocalOp>(def))
+ return privateLocal;
+ worklist.append(def->getOperands().begin(), def->getOperands().end());
}
return nullptr;
}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index ba8f16d3966df..0de030eda1e59 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -35,6 +35,7 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
+ %unused = memref.load %result_arg[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
@@ -184,80 +185,3 @@ func.func @worker_combine_in_scf_if(%result: memref<i32>) {
}
return
}
-
-// CHECK-LABEL: func.func @aliased_worker_store
-// CHECK: gpu.launch {{.*}} threads([[ALIAS_TX:%[^,]+]], [[ALIAS_TY:%[^,]+]],
-// CHECK-NOT: arith.cmpi eq, [[ALIAS_TY]]
-// CHECK: %[[ALIAS_TX_ZERO:.*]] = arith.cmpi eq, [[ALIAS_TX]],
-// CHECK-NOT: arith.andi
-// CHECK: scf.if %[[ALIAS_TX_ZERO]]
-// CHECK: memref.store
-// CHECK: acc.atomic.update
-func.func @aliased_worker_store(%result: memref<i32>) {
- %c1 = arith.constant 1 : index
- %c4 = arith.constant 4 : index
- %c32 = arith.constant 32 : index
- %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
- %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
- %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
- acc.kernel_environment {
- %private = acc.privatize [#acc<par_dims[thread_y]>]
- : () -> !acc.private_type<memref<i32>>
- acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
- ins(%private_arg = %private, %result_arg = %result)
- : (!acc.private_type<memref<i32>>, memref<i32>) {
- %c0 = arith.constant 0 : index
- %c1_inner = arith.constant 1 : index
- %c7_i32 = arith.constant 7 : i32
- scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
- %local = acc.private_local %private_arg
- : (!acc.private_type<memref<i32>>) -> memref<i32>
- %cast = memref.cast %local : memref<i32> to memref<i32>
- acc.predicate_region {
- memref.store %c7_i32, %cast[] : memref<i32>
- acc.reduction_combine %local into %result_arg <add> : memref<i32>
- {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
- }
- scf.reduce
- } {acc.par_dims = #acc<par_dims[block_y]>}
- acc.yield
- } {origin = "acc.parallel"}
- }
- return
-}
-
-// CHECK-LABEL: func.func @aliased_thread_store
-// CHECK: gpu.launch {{.*}} threads([[THREAD_TX:%[^,]+]], [[THREAD_TY:%[^,]+]],
-// CHECK-NOT: arith.cmpi eq, [[THREAD_TX]]
-// CHECK-NOT: arith.cmpi eq, [[THREAD_TY]]
-// CHECK: memref.store
-func.func @aliased_thread_store() {
- %c1 = arith.constant 1 : index
- %c4 = arith.constant 4 : index
- %c32 = arith.constant 32 : index
- %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
- %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
- %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
- acc.kernel_environment {
- %private = acc.privatize [#acc<par_dims[thread_y, thread_x]>]
- : () -> !acc.private_type<memref<i32>>
- acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
- ins(%private_arg = %private)
- : (!acc.private_type<memref<i32>>) {
- %c0 = arith.constant 0 : index
- %c1_inner = arith.constant 1 : index
- %c7_i32 = arith.constant 7 : i32
- scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
- %local = acc.private_local %private_arg
- : (!acc.private_type<memref<i32>>) -> memref<i32>
- %cast = memref.cast %local : memref<i32> to memref<i32>
- acc.predicate_region {
- memref.store %c7_i32, %cast[] : memref<i32>
- }
- scf.reduce
- } {acc.par_dims = #acc<par_dims[block_y]>}
- acc.yield
- } {origin = "acc.parallel"}
- }
- return
-}
>From 30926fc3c3e3a6d6b9fe8e3fcb8442319849edf1 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 17:43:35 -0700
Subject: [PATCH 08/22] [OpenACC][test] Match broadening conflict locations
Expect diagnostics on the side effect that makes worker broadening unsafe.
---
.../acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index 1275e4c77890a..7f22fd20aa89e 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -69,8 +69,8 @@ func.func @worker_combine_with_single_store(%result: memref<i32>) {
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
- memref.store %c7_i32, %selected[] : memref<i32>
// expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
+ memref.store %c7_i32, %selected[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
@@ -103,12 +103,12 @@ func.func @worker_combine_with_atomic_update(%result: memref<i32>) {
%local = acc.private_local %private_arg
: (!acc.private_type<memref<i32>>) -> memref<i32>
acc.predicate_region {
+ // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
acc.atomic.update %result_arg : memref<i32> {
^bb0(%current: i32):
%next = arith.addi %current, %c1_i32 : i32
acc.yield %next : i32
}
- // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
>From 55293fcdfa6ce6f49e61e43293e80df04bfec368 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 17:51:30 -0700
Subject: [PATCH 09/22] [OpenACC] Require safe predicate speculation
Treat reads and non-speculatable operations as broadening conflicts so enabling worker rows cannot introduce invalid memory accesses.
---
mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 13 +++++--------
...to-gpu-worker-reduction-combine-mixed-scope.mlir | 6 ++----
.../acc-cg-to-gpu-worker-reduction-combine.mlir | 1 -
3 files changed, 7 insertions(+), 13 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 50635dde561db..250c2be09df85 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1276,13 +1276,9 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
return false;
};
- auto hasUnsafeOwnEffects = [](Operation *op) {
+ auto hasOwnEffects = [](Operation *op) {
if (auto effectOp = dyn_cast<MemoryEffectOpInterface>(op)) {
- SmallVector<MemoryEffects::EffectInstance> effects;
- effectOp.getEffects(effects);
- return llvm::any_of(effects, [](const auto &effect) {
- return !isa<MemoryEffects::Read>(effect.getEffect());
- });
+ return !effectOp.hasNoEffect();
}
return !op->hasTrait<OpTrait::HasRecursiveMemoryEffects>();
};
@@ -1325,7 +1321,7 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
continue;
}
if (nestedOp.getNumRegions() != 0) {
- if (hasUnsafeOwnEffects(&nestedOp)) {
+ if (hasOwnEffects(&nestedOp)) {
info.hasBroadeningConflict = true;
if (!info.diagnosticOp)
info.diagnosticOp = &nestedOp;
@@ -1336,7 +1332,8 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
continue;
}
- if (hasUnsafeOwnEffects(&nestedOp)) {
+ if (!nestedOp.mightHaveTrait<OpTrait::IsTerminator>() &&
+ !isPure(&nestedOp)) {
info.hasBroadeningConflict = true;
if (!info.diagnosticOp)
info.diagnosticOp = &nestedOp;
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index 7f22fd20aa89e..290786ab7c1ce 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -42,7 +42,7 @@ func.func @mixed_scope_worker_reduction_combine(
return
}
-func.func @worker_combine_with_single_store(%result: memref<i32>) {
+func.func @worker_combine_with_read(%result: memref<i32>) {
%c1 = arith.constant 1 : index
%c4 = arith.constant 4 : index
%c32 = arith.constant 32 : index
@@ -59,18 +59,16 @@ func.func @worker_combine_with_single_store(%result: memref<i32>) {
%c0 = arith.constant 0 : index
%c1_inner = arith.constant 1 : index
%c7_i32 = arith.constant 7 : i32
- %false = arith.constant false
scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
%local = acc.private_local %private_arg
: (!acc.private_type<memref<i32>>) -> memref<i32>
- %selected = arith.select %false, %local, %result_arg : memref<i32>
scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
memref.store %c7_i32, %local[] : memref<i32>
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
// expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
- memref.store %c7_i32, %selected[] : memref<i32>
+ %unused = memref.load %result_arg[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index 0de030eda1e59..fb1a0256d2a38 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -35,7 +35,6 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
- %unused = memref.load %result_arg[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
>From 300b3dc6a9334c5bd8b47daf1c1371bb6f5db806 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 17:59:06 -0700
Subject: [PATCH 10/22] [OpenACC] Reject worker combine regions
Keep generic combine regions single-row because their non-atomic load/store lowering is unsafe across worker rows.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 14 ++++--
...-worker-reduction-combine-mixed-scope.mlir | 37 ++++++++++++++
...cc-cg-to-gpu-worker-reduction-combine.mlir | 48 -------------------
3 files changed, 48 insertions(+), 51 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 250c2be09df85..3697ae7405c9f 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1286,11 +1286,17 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
[&](auto &&self, Block &predicateBlock) -> ThreadYBroadeningInfo {
ThreadYBroadeningInfo info;
auto classifyCombine = [&](Operation *combineOp, Value src,
- ArrayRef<GPUParallelDimAttr> parDims) {
+ ArrayRef<GPUParallelDimAttr> parDims,
+ bool supportsWorkerBroadening) {
bool hasThreadY = llvm::any_of(
parDims, [](auto parDim) { return parDim.isThreadY(); });
if (hasThreadY && isProvenWorkerPrivate(src)) {
info.hasActiveWorkerCombine = true;
+ if (!supportsWorkerBroadening) {
+ info.hasBroadeningConflict = true;
+ if (!info.diagnosticOp)
+ info.diagnosticOp = combineOp;
+ }
} else {
info.hasExplicitInactiveCombine = true;
if (!info.diagnosticOp)
@@ -1311,13 +1317,15 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
if (acc::ReductionCombineOp combineOp =
dyn_cast<acc::ReductionCombineOp>(nestedOp)) {
classifyCombine(combineOp, combineOp.getSrcMemref(),
- getReductionCombineParDims(combineOp));
+ getReductionCombineParDims(combineOp),
+ /*supportsWorkerBroadening=*/true);
continue;
}
if (acc::ReductionCombineRegionOp combineRegionOp =
dyn_cast<acc::ReductionCombineRegionOp>(nestedOp)) {
classifyCombine(combineRegionOp, combineRegionOp.getSrcVar(),
- getReductionCombineParDims(combineRegionOp));
+ getReductionCombineParDims(combineRegionOp),
+ /*supportsWorkerBroadening=*/false);
continue;
}
if (nestedOp.getNumRegions() != 0) {
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index 290786ab7c1ce..fffbfcb461c56 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -117,3 +117,40 @@ func.func @worker_combine_with_atomic_update(%result: memref<i32>) {
}
return
}
+
+func.func @worker_reduction_combine_region(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ // expected-error at +1 {{failed to legalize operation 'acc.compute_region' that was explicitly marked illegal}}
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ acc.predicate_region {
+ // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
+ acc.reduction_combine_region %local into %result_arg : memref<i32> {
+ %lhs = memref.load %result_arg[] : memref<i32>
+ %rhs = memref.load %local[] : memref<i32>
+ %sum = arith.addi %lhs, %rhs : i32
+ memref.store %sum, %result_arg[] : memref<i32>
+ acc.yield
+ } {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index fb1a0256d2a38..686e8fd233980 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -46,54 +46,6 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
return
}
-// CHECK-LABEL: func.func @worker_reduction_combine_region
-// CHECK: gpu.launch {{.*}} threads([[REGION_TID_X:%[^,]+]], [[REGION_TID_Y:%[^,]+]],
-// CHECK-NOT: arith.cmpi eq, [[REGION_TID_Y]]
-// CHECK: %[[REGION_IS_X_ZERO:.*]] = arith.cmpi eq, [[REGION_TID_X]],
-// CHECK-NOT: arith.andi
-// CHECK: scf.if %[[REGION_IS_X_ZERO]]
-// CHECK: arith.addi
-
-func.func @worker_reduction_combine_region(%result: memref<i32>) {
- %c1 = arith.constant 1 : index
- %c4 = arith.constant 4 : index
- %c32 = arith.constant 32 : index
- %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
- %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
- %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
- acc.kernel_environment {
- %private = acc.privatize [#acc<par_dims[thread_y]>]
- : () -> !acc.private_type<memref<i32>>
- acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
- ins(%private_arg = %private, %result_arg = %result)
- : (!acc.private_type<memref<i32>>, memref<i32>) {
- %c0 = arith.constant 0 : index
- %c1_inner = arith.constant 1 : index
- %c0_i32 = arith.constant 0 : i32
- scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
- %local = acc.private_local %private_arg
- : (!acc.private_type<memref<i32>>) -> memref<i32>
- scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
- memref.store %c0_i32, %local[] : memref<i32>
- scf.reduce
- } {acc.par_dims = #acc<par_dims[thread_y]>}
- acc.predicate_region {
- acc.reduction_combine_region %local into %result_arg : memref<i32> {
- %lhs = memref.load %result_arg[] : memref<i32>
- %rhs = memref.load %local[] : memref<i32>
- %sum = arith.addi %lhs, %rhs : i32
- memref.store %sum, %result_arg[] : memref<i32>
- acc.yield
- } {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
- }
- scf.reduce
- } {acc.par_dims = #acc<par_dims[block_y]>}
- acc.yield
- } {origin = "acc.parallel"}
- }
- return
-}
-
// Nested predicate regions choose their ThreadY predicates independently. The
// outer region must keep ThreadY active so it does not exclude worker rows
// before the nested worker-private combine is reached.
>From cbc18adba26e38ff6bebd01bf252112da4412887 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 18:03:05 -0700
Subject: [PATCH 11/22] [OpenACC] Require atomic worker combines
Broaden ThreadY only when block dimensions select the atomic combine lowering; reject ThreadY-only shared updates.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 4 ++-
...-worker-reduction-combine-mixed-scope.mlir | 32 +++++++++++++++++++
2 files changed, 35 insertions(+), 1 deletion(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 3697ae7405c9f..76f942a2730f9 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1290,7 +1290,9 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
bool supportsWorkerBroadening) {
bool hasThreadY = llvm::any_of(
parDims, [](auto parDim) { return parDim.isThreadY(); });
- if (hasThreadY && isProvenWorkerPrivate(src)) {
+ bool hasBlock = llvm::any_of(
+ parDims, [](auto parDim) { return parDim.isAnyBlock(); });
+ if (hasThreadY && hasBlock && isProvenWorkerPrivate(src)) {
info.hasActiveWorkerCombine = true;
if (!supportsWorkerBroadening) {
info.hasBroadeningConflict = true;
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index fffbfcb461c56..5758702a3b92e 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -154,3 +154,35 @@ func.func @worker_reduction_combine_region(%result: memref<i32>) {
}
return
}
+
+func.func @thread_y_only_worker_combine(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ // expected-error at +1 {{failed to legalize operation 'acc.compute_region' that was explicitly marked illegal}}
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ acc.predicate_region {
+ // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
>From 2f6b6b2e74b62114ad20366b768685e6af8bc000 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 18:06:26 -0700
Subject: [PATCH 12/22] [OpenACC][test] Check non-atomic worker predication
Verify ThreadY-only combines retain single-row predication while unsafe combine regions remain rejected.
---
...-worker-reduction-combine-mixed-scope.mlir | 31 ---------------
...cc-cg-to-gpu-worker-reduction-combine.mlir | 39 +++++++++++++++++++
2 files changed, 39 insertions(+), 31 deletions(-)
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index 5758702a3b92e..3d15b4ca4a59f 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -155,34 +155,3 @@ func.func @worker_reduction_combine_region(%result: memref<i32>) {
return
}
-func.func @thread_y_only_worker_combine(%result: memref<i32>) {
- %c1 = arith.constant 1 : index
- %c4 = arith.constant 4 : index
- %c32 = arith.constant 32 : index
- %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
- %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
- %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
- acc.kernel_environment {
- %private = acc.privatize [#acc<par_dims[thread_y]>]
- : () -> !acc.private_type<memref<i32>>
- // expected-error at +1 {{failed to legalize operation 'acc.compute_region' that was explicitly marked illegal}}
- acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
- ins(%private_arg = %private, %result_arg = %result)
- : (!acc.private_type<memref<i32>>, memref<i32>) {
- %c0 = arith.constant 0 : index
- %c1_inner = arith.constant 1 : index
- scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
- %local = acc.private_local %private_arg
- : (!acc.private_type<memref<i32>>) -> memref<i32>
- acc.predicate_region {
- // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
- acc.reduction_combine %local into %result_arg <add> : memref<i32>
- {acc.par_dims = #acc<par_dims[thread_y]>}
- }
- scf.reduce
- } {acc.par_dims = #acc<par_dims[block_y]>}
- acc.yield
- } {origin = "acc.parallel"}
- }
- return
-}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index 686e8fd233980..901b23b25316f 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -136,3 +136,42 @@ func.func @worker_combine_in_scf_if(%result: memref<i32>) {
}
return
}
+
+// Without a block dimension the combine lowers to a non-atomic update, so
+// ThreadY must retain the legacy single-row predicate.
+// CHECK-LABEL: func.func @thread_y_only_worker_combine
+// CHECK: gpu.launch {{.*}} threads([[ONLY_TX:%[^,]+]], [[ONLY_TY:%[^,]+]],
+// CHECK: %[[ONLY_TY_ZERO:.*]] = arith.cmpi eq, [[ONLY_TY]],
+// CHECK: %[[ONLY_TX_ZERO:.*]] = arith.cmpi eq, [[ONLY_TX]],
+// CHECK: %[[ONLY_ROW_ZERO:.*]] = arith.andi %[[ONLY_TX_ZERO]], %[[ONLY_TY_ZERO]]
+// CHECK: scf.if %[[ONLY_ROW_ZERO]]
+// CHECK: memref.store
+func.func @thread_y_only_worker_combine(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ acc.predicate_region {
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
>From 88026c11dfb0376be0ad438de4a530d99d821119 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 18:10:52 -0700
Subject: [PATCH 13/22] [OpenACC][test] Remove trailing blank line
---
.../acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir | 1 -
1 file changed, 1 deletion(-)
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index 3d15b4ca4a59f..dc634cff25b07 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -41,7 +41,6 @@ func.func @mixed_scope_worker_reduction_combine(
}
return
}
-
func.func @worker_combine_with_read(%result: memref<i32>) {
%c1 = arith.constant 1 : index
%c4 = arith.constant 4 : index
>From 984e8f2ce4366cb23a6104b6aa965783f2380a49 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 18:11:22 -0700
Subject: [PATCH 14/22] [OpenACC][test] Normalize test file ending
---
.../acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir | 1 -
1 file changed, 1 deletion(-)
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index dc634cff25b07..dba8057d48e95 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -78,7 +78,6 @@ func.func @worker_combine_with_read(%result: memref<i32>) {
}
return
}
-
func.func @worker_combine_with_atomic_update(%result: memref<i32>) {
%c1 = arith.constant 1 : index
%c4 = arith.constant 4 : index
>From 8f6c707f4e58523929b8482505946acdd9dc9991 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 18:12:19 -0700
Subject: [PATCH 15/22] [OpenACC][test] Finish file ending cleanup
---
.../acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index dba8057d48e95..a641fe74867e4 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -151,5 +151,4 @@ func.func @worker_reduction_combine_region(%result: memref<i32>) {
} {origin = "acc.parallel"}
}
return
-}
-
+}
\ No newline at end of file
>From f629ee527f4b56ac33d85867d488a793af7cfc09 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 18:43:52 -0700
Subject: [PATCH 16/22] [OpenACC] Localize worker combine classification
Keep private-memory scope behavior unchanged and remove redundant failure plumbing from combine dimension handling.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 37 ++++++++++---------
1 file changed, 20 insertions(+), 17 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 76f942a2730f9..98258dc622e3c 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1265,8 +1265,18 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
while (Operation *def = current.getDefiningOp()) {
if (acc::PrivateLocalOp privateLocal =
dyn_cast<acc::PrivateLocalOp>(def)) {
- return getPrivateMemScope(getPrivatizeOp(
- privateLocal, computeRegion)) == PrivateMemScope::Worker;
+ acc::PrivatizeOp privatize =
+ getPrivatizeOp(privateLocal, computeRegion);
+ GPUParallelDimsAttr parDims = privatize.getParDimsAttr();
+ if (!parDims)
+ return false;
+ bool hasThreadX = llvm::any_of(parDims.getArray(), [](auto parDim) {
+ return parDim.isThreadX();
+ });
+ bool hasThreadY = llvm::any_of(parDims.getArray(), [](auto parDim) {
+ return parDim.isThreadY();
+ });
+ return hasThreadY && !hasThreadX;
}
if (ViewLikeOpInterface viewLike = dyn_cast<ViewLikeOpInterface>(def)) {
current = viewLike.getViewSource();
@@ -1355,13 +1365,7 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
ThreadYBroadeningInfo threadYInfo =
analyzeThreadYBroadening(analyzeThreadYBroadening, *block);
- auto applyCombineParDims =
- [&](ArrayRef<mlir::acc::GPUParallelDimAttr> combineParDims) {
- for (mlir::acc::GPUParallelDimAttr parDim : combineParDims)
- mlir::acc::removeParDim(ancestorParDims, parDim);
- return success();
- };
- block->walk([&](Operation *op) -> WalkResult {
+ block->walk([&](Operation *op) {
// Check stores to acc.private_local - add the privatize's par_dims
// as active dims so predication is correct for per-worker/gang memory.
if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(op)) {
@@ -1392,15 +1396,15 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
// kernel and loop in combined constructs.
if (acc::ReductionCombineOp reductionCombineOp =
dyn_cast<acc::ReductionCombineOp>(op)) {
- if (failed(applyCombineParDims(
- getReductionCombineParDims(reductionCombineOp))))
- return WalkResult::interrupt();
+ for (GPUParallelDimAttr parDim :
+ getReductionCombineParDims(reductionCombineOp))
+ mlir::acc::removeParDim(ancestorParDims, parDim);
}
if (acc::ReductionCombineRegionOp combineRegionOp =
dyn_cast<acc::ReductionCombineRegionOp>(op)) {
- if (failed(applyCombineParDims(
- getReductionCombineParDims(combineRegionOp))))
- return WalkResult::interrupt();
+ for (GPUParallelDimAttr parDim :
+ getReductionCombineParDims(combineRegionOp))
+ mlir::acc::removeParDim(ancestorParDims, parDim);
}
// An array accumulate reduces across its par_dims via gpu.all_reduce, so
// all those threads must execute it - treat them as active (unlike the
@@ -1412,7 +1416,6 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
mlir::acc::insertParDim(ancestorParDims, parDim);
}
}
- return WalkResult::advance();
});
mlir::acc::GPUParallelDimAttr threadY =
mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
@@ -1877,7 +1880,7 @@ ACCCGToGPULowering::getPrivateMemScope(acc::PrivatizeOp privatizeOp) {
}
if (hasThreadX)
return PrivateMemScope::Thread;
- if (hasThreadY)
+ if (hasBlock && hasThreadY)
return PrivateMemScope::Worker;
if (hasBlock)
return PrivateMemScope::Gang;
>From edc8abd53cd69917a03f21c72c975e4a26a4ed24 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 18:50:05 -0700
Subject: [PATCH 17/22] [OpenACC] Merge ThreadY conflict state
Use one flag for all operations that require the legacy inactive ThreadY predicate.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 17 +++++++----------
1 file changed, 7 insertions(+), 10 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 98258dc622e3c..6a5aebc4db323 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1248,15 +1248,13 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
if (block) {
struct ThreadYBroadeningInfo {
bool hasActiveWorkerCombine = false;
- bool hasExplicitInactiveCombine = false;
- bool hasBroadeningConflict = false;
+ bool requiresInactiveThreadY = false;
Operation *diagnosticOp = nullptr;
};
auto mergeBroadeningInfo = [](ThreadYBroadeningInfo &dst,
const ThreadYBroadeningInfo &src) {
dst.hasActiveWorkerCombine |= src.hasActiveWorkerCombine;
- dst.hasExplicitInactiveCombine |= src.hasExplicitInactiveCombine;
- dst.hasBroadeningConflict |= src.hasBroadeningConflict;
+ dst.requiresInactiveThreadY |= src.requiresInactiveThreadY;
if (!dst.diagnosticOp)
dst.diagnosticOp = src.diagnosticOp;
};
@@ -1305,12 +1303,12 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
if (hasThreadY && hasBlock && isProvenWorkerPrivate(src)) {
info.hasActiveWorkerCombine = true;
if (!supportsWorkerBroadening) {
- info.hasBroadeningConflict = true;
+ info.requiresInactiveThreadY = true;
if (!info.diagnosticOp)
info.diagnosticOp = combineOp;
}
} else {
- info.hasExplicitInactiveCombine = true;
+ info.requiresInactiveThreadY = true;
if (!info.diagnosticOp)
info.diagnosticOp = combineOp;
}
@@ -1342,7 +1340,7 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
if (nestedOp.getNumRegions() != 0) {
if (hasOwnEffects(&nestedOp)) {
- info.hasBroadeningConflict = true;
+ info.requiresInactiveThreadY = true;
if (!info.diagnosticOp)
info.diagnosticOp = &nestedOp;
}
@@ -1354,7 +1352,7 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
if (!nestedOp.mightHaveTrait<OpTrait::IsTerminator>() &&
!isPure(&nestedOp)) {
- info.hasBroadeningConflict = true;
+ info.requiresInactiveThreadY = true;
if (!info.diagnosticOp)
info.diagnosticOp = &nestedOp;
}
@@ -1421,8 +1419,7 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
bool baselineThreadYActive = llvm::is_contained(ancestorParDims, threadY);
if (threadYInfo.hasActiveWorkerCombine && !baselineThreadYActive) {
- if (threadYInfo.hasExplicitInactiveCombine ||
- threadYInfo.hasBroadeningConflict) {
+ if (threadYInfo.requiresInactiveThreadY) {
Operation *diagnosticOp =
threadYInfo.diagnosticOp ? threadYInfo.diagnosticOp : op;
(void)accSupport.emitNYI(
>From 90cbc61e638708c7e25e37e588e20329fbfb736e Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 19:06:50 -0700
Subject: [PATCH 18/22] [OpenACC] Stabilize ThreadY conflict analysis
Capture structural ThreadY activity before operation-order mutations and cover a trailing worker-private store.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 6 +++---
...pu-worker-reduction-combine-mixed-scope.mlir | 1 +
.../acc-cg-to-gpu-worker-reduction-combine.mlir | 17 +++++------------
3 files changed, 9 insertions(+), 15 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 6a5aebc4db323..09dbbf892ffbd 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1246,6 +1246,9 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
mlir::acc::GPUParallelDimAttr lowestParDim =
mlir::acc::GPUParallelDimAttr::threadXDim(ctx);
if (block) {
+ mlir::acc::GPUParallelDimAttr threadY =
+ mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
+ bool baselineThreadYActive = llvm::is_contained(ancestorParDims, threadY);
struct ThreadYBroadeningInfo {
bool hasActiveWorkerCombine = false;
bool requiresInactiveThreadY = false;
@@ -1415,9 +1418,6 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
}
});
- mlir::acc::GPUParallelDimAttr threadY =
- mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
- bool baselineThreadYActive = llvm::is_contained(ancestorParDims, threadY);
if (threadYInfo.hasActiveWorkerCombine && !baselineThreadYActive) {
if (threadYInfo.requiresInactiveThreadY) {
Operation *diagnosticOp =
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index a641fe74867e4..a5ffb4bfb59f8 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -33,6 +33,7 @@ func.func @mixed_scope_worker_reduction_combine(
// expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
acc.reduction_combine %other_arg into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_x]>}
+ memref.store %c0_i32, %local[] : memref<i32>
}
scf.reduce
} {acc.par_dims = #acc<par_dims[block_y]>}
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index 901b23b25316f..4466ad74d1a91 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -46,17 +46,15 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
return
}
-// Nested predicate regions choose their ThreadY predicates independently. The
-// outer region must keep ThreadY active so it does not exclude worker rows
-// before the nested worker-private combine is reached.
+// The outer predicate must keep ThreadY active so it does not exclude worker
+// rows before the nested worker-private combine is reached.
// CHECK-LABEL: func.func @nested_worker_reduction_combines
// CHECK: gpu.launch {{.*}} threads([[NESTED_TX:%[^,]+]], [[NESTED_TY:%[^,]+]],
// CHECK-NOT: arith.cmpi eq, [[NESTED_TY]]
// CHECK: %[[NESTED_TX_ZERO:.*]] = arith.cmpi eq, [[NESTED_TX]],
// CHECK-NOT: arith.andi
// CHECK: scf.if %[[NESTED_TX_ZERO]]
-func.func @nested_worker_reduction_combines(
- %other: memref<i32>, %result: memref<i32>) {
+func.func @nested_worker_reduction_combines(%result: memref<i32>) {
%c1 = arith.constant 1 : index
%c4 = arith.constant 4 : index
%c32 = arith.constant 32 : index
@@ -67,9 +65,8 @@ func.func @nested_worker_reduction_combines(
%private = acc.privatize [#acc<par_dims[thread_y]>]
: () -> !acc.private_type<memref<i32>>
acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
- ins(%private_arg = %private, %other_arg = %other,
- %result_arg = %result)
- : (!acc.private_type<memref<i32>>, memref<i32>, memref<i32>) {
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
%c0 = arith.constant 0 : index
%c1_inner = arith.constant 1 : index
%c0_i32 = arith.constant 0 : i32
@@ -85,10 +82,6 @@ func.func @nested_worker_reduction_combines(
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
- acc.predicate_region {
- acc.reduction_combine %other_arg into %result_arg <add> : memref<i32>
- {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
- }
}
scf.reduce
} {acc.par_dims = #acc<par_dims[block_y]>}
>From 9e6fcfb6dd7ff5d68e1e7e6a34b14a7a69763286 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 19:10:18 -0700
Subject: [PATCH 19/22] [OpenACC] Allow worker-private stores during combine
Treat proven worker-private stores as row-local so they can share the broadened ThreadY predicate safely.
---
mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 4 ++++
.../OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir | 1 +
2 files changed, 5 insertions(+)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 09dbbf892ffbd..f78a973da4e5a 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1341,6 +1341,10 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
/*supportsWorkerBroadening=*/false);
continue;
}
+ if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(nestedOp)) {
+ if (isProvenWorkerPrivate(storeOp.getMemref()))
+ continue;
+ }
if (nestedOp.getNumRegions() != 0) {
if (hasOwnEffects(&nestedOp)) {
info.requiresInactiveThreadY = true;
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index 4466ad74d1a91..cbfabf28ea891 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -35,6 +35,7 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
+ memref.store %c0_i32, %local[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
>From fcc26cd10cf22d3fb30b46154828ecd7e4af7604 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 20:43:59 -0700
Subject: [PATCH 20/22] [OpenACC] Diagnose unproven worker aliases
Restrict ThreadY broadening to directly sized worker storage and conservatively reject aliases that may originate from it.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 73 +++++++++++--------
...-worker-reduction-combine-mixed-scope.mlir | 39 ++++++++++
...cc-cg-to-gpu-worker-reduction-combine.mlir | 41 ++++++++++-
3 files changed, 123 insertions(+), 30 deletions(-)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index f78a973da4e5a..f46a891f08652 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1248,7 +1248,6 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
if (block) {
mlir::acc::GPUParallelDimAttr threadY =
mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
- bool baselineThreadYActive = llvm::is_contained(ancestorParDims, threadY);
struct ThreadYBroadeningInfo {
bool hasActiveWorkerCombine = false;
bool requiresInactiveThreadY = false;
@@ -1262,30 +1261,46 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
dst.diagnosticOp = src.diagnosticOp;
};
auto isProvenWorkerPrivate = [&](Value memref) {
- Value current = memref;
- while (Operation *def = current.getDefiningOp()) {
- if (acc::PrivateLocalOp privateLocal =
- dyn_cast<acc::PrivateLocalOp>(def)) {
- acc::PrivatizeOp privatize =
- getPrivatizeOp(privateLocal, computeRegion);
- GPUParallelDimsAttr parDims = privatize.getParDimsAttr();
- if (!parDims)
- return false;
- bool hasThreadX = llvm::any_of(parDims.getArray(), [](auto parDim) {
- return parDim.isThreadX();
- });
- bool hasThreadY = llvm::any_of(parDims.getArray(), [](auto parDim) {
- return parDim.isThreadY();
- });
- return hasThreadY && !hasThreadX;
- }
- if (ViewLikeOpInterface viewLike = dyn_cast<ViewLikeOpInterface>(def)) {
- current = viewLike.getViewSource();
- continue;
- }
- break;
+ acc::PrivateLocalOp privateLocal =
+ memref.getDefiningOp<acc::PrivateLocalOp>();
+ if (!privateLocal)
+ return false;
+ acc::PrivatizeOp privatize = getPrivatizeOp(privateLocal, computeRegion);
+ GPUParallelDimsAttr parDims = privatize.getParDimsAttr();
+ if (!parDims)
+ return false;
+ bool hasThreadX = llvm::any_of(
+ parDims.getArray(), [](auto parDim) { return parDim.isThreadX(); });
+ bool hasThreadY = llvm::any_of(
+ parDims.getArray(), [](auto parDim) { return parDim.isThreadY(); });
+ return hasThreadY && !hasThreadX;
+ };
+ auto hasWorkerPrivateOrigin = [&](auto &&self, Value value,
+ DenseSet<Value> &visited) -> bool {
+ if (!visited.insert(value).second)
+ return false;
+ if (isProvenWorkerPrivate(value))
+ return true;
+ Operation *def = value.getDefiningOp();
+ if (!def) {
+ auto blockArg = dyn_cast<BlockArgument>(value);
+ def = blockArg ? blockArg.getOwner()->getParentOp() : nullptr;
}
- return false;
+ if (!def)
+ return false;
+ if (llvm::any_of(def->getOperands(), [&](Value operand) {
+ return self(self, operand, visited);
+ }))
+ return true;
+ return llvm::any_of(def->getRegions(), [&](Region ®ion) {
+ return llvm::any_of(region, [&](Block &block) {
+ Operation *terminator = block.getTerminator();
+ return terminator &&
+ llvm::any_of(terminator->getOperands(), [&](Value operand) {
+ return self(self, operand, visited);
+ });
+ });
+ });
};
auto hasOwnEffects = [](Operation *op) {
if (auto effectOp = dyn_cast<MemoryEffectOpInterface>(op)) {
@@ -1312,6 +1327,10 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
} else {
info.requiresInactiveThreadY = true;
+ DenseSet<Value> visited;
+ if (hasThreadY && hasBlock &&
+ hasWorkerPrivateOrigin(hasWorkerPrivateOrigin, src, visited))
+ info.hasActiveWorkerCombine = true;
if (!info.diagnosticOp)
info.diagnosticOp = combineOp;
}
@@ -1341,10 +1360,6 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
/*supportsWorkerBroadening=*/false);
continue;
}
- if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(nestedOp)) {
- if (isProvenWorkerPrivate(storeOp.getMemref()))
- continue;
- }
if (nestedOp.getNumRegions() != 0) {
if (hasOwnEffects(&nestedOp)) {
info.requiresInactiveThreadY = true;
@@ -1422,7 +1437,7 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
}
}
});
- if (threadYInfo.hasActiveWorkerCombine && !baselineThreadYActive) {
+ if (threadYInfo.hasActiveWorkerCombine) {
if (threadYInfo.requiresInactiveThreadY) {
Operation *diagnosticOp =
threadYInfo.diagnosticOp ? threadYInfo.diagnosticOp : op;
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
index a5ffb4bfb59f8..8724f1f57bc9b 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine-mixed-scope.mlir
@@ -42,6 +42,45 @@ func.func @mixed_scope_worker_reduction_combine(
}
return
}
+
+func.func @worker_combine_with_unproven_alias(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ // expected-error at +1 {{failed to legalize operation 'acc.compute_region' that was explicitly marked illegal}}
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %true = arith.constant true
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ %alias = scf.if %true -> (memref<i32>) {
+ scf.yield %local : memref<i32>
+ } else {
+ scf.yield %local : memref<i32>
+ }
+ acc.predicate_region {
+ // expected-error at +1 {{operations in the same predicate region require incompatible ThreadY predication}}
+ acc.reduction_combine %alias into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
+
func.func @worker_combine_with_read(%result: memref<i32>) {
%c1 = arith.constant 1 : index
%c4 = arith.constant 4 : index
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index cbfabf28ea891..fd4e12fd039a8 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -35,7 +35,6 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
- memref.store %c0_i32, %local[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
@@ -47,6 +46,46 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
return
}
+// CHECK-LABEL: func.func @worker_combine_in_thread_y_scope
+// CHECK: gpu.launch {{.*}} threads([[SCOPE_TX:%[^,]+]], [[SCOPE_TY:%[^,]+]],
+// CHECK-NOT: arith.cmpi eq, [[SCOPE_TY]]
+// CHECK: %[[SCOPE_TX_ZERO:.*]] = arith.cmpi eq, [[SCOPE_TX]],
+// CHECK: scf.if %[[SCOPE_TX_ZERO]]
+func.func @worker_combine_in_thread_y_scope(%result: memref<i32>) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ %block_y = acc.par_width %c1 {par_dim = #acc.par_dim<block_y>}
+ %thread_y = acc.par_width %c4 {par_dim = #acc.par_dim<thread_y>}
+ %thread_x = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
+ acc.kernel_environment {
+ %private = acc.privatize [#acc<par_dims[thread_y]>]
+ : () -> !acc.private_type<memref<i32>>
+ acc.compute_region launch(%by = %block_y, %ty = %thread_y, %tx = %thread_x)
+ ins(%private_arg = %private, %result_arg = %result)
+ : (!acc.private_type<memref<i32>>, memref<i32>) {
+ %c0 = arith.constant 0 : index
+ %c1_inner = arith.constant 1 : index
+ %c0_i32 = arith.constant 0 : i32
+ scf.parallel (%block_iv) = (%c0) to (%by) step (%c1_inner) {
+ %local = acc.private_local %private_arg
+ : (!acc.private_type<memref<i32>>) -> memref<i32>
+ scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
+ memref.store %c0_i32, %local[] : memref<i32>
+ acc.predicate_region {
+ acc.reduction_combine %local into %result_arg <add> : memref<i32>
+ {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ }
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[thread_y]>}
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y]>}
+ acc.yield
+ } {origin = "acc.parallel"}
+ }
+ return
+}
+
// The outer predicate must keep ThreadY active so it does not exclude worker
// rows before the nested worker-private combine is reached.
// CHECK-LABEL: func.func @nested_worker_reduction_combines
>From ca846e5a59a0e4ce0ec567a391cb4ab23f777640 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 20:57:39 -0700
Subject: [PATCH 21/22] [OpenACC] Allow stores to active worker sources
Permit row-local stores only when they target the exact private storage already sized for a supported worker combine.
---
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 16 ++++++++++++++++
.../acc-cg-to-gpu-worker-reduction-combine.mlir | 2 ++
2 files changed, 18 insertions(+)
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index f46a891f08652..bf90b51a2d05b 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -1311,6 +1311,18 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
auto analyzeThreadYBroadening =
[&](auto &&self, Block &predicateBlock) -> ThreadYBroadeningInfo {
ThreadYBroadeningInfo info;
+ DenseSet<Value> activeWorkerSources;
+ predicateBlock.walk([&](acc::ReductionCombineOp combineOp) {
+ ArrayRef<GPUParallelDimAttr> parDims =
+ getReductionCombineParDims(combineOp);
+ bool hasThreadY = llvm::any_of(
+ parDims, [](auto parDim) { return parDim.isThreadY(); });
+ bool hasBlock = llvm::any_of(
+ parDims, [](auto parDim) { return parDim.isAnyBlock(); });
+ if (hasThreadY && hasBlock &&
+ isProvenWorkerPrivate(combineOp.getSrcMemref()))
+ activeWorkerSources.insert(combineOp.getSrcMemref());
+ });
auto classifyCombine = [&](Operation *combineOp, Value src,
ArrayRef<GPUParallelDimAttr> parDims,
bool supportsWorkerBroadening) {
@@ -1360,6 +1372,10 @@ ACCCGToGPULowering::computeActiveAndInactiveParDims(Operation *op,
/*supportsWorkerBroadening=*/false);
continue;
}
+ if (memref::StoreOp storeOp = dyn_cast<memref::StoreOp>(nestedOp)) {
+ if (activeWorkerSources.contains(storeOp.getMemref()))
+ continue;
+ }
if (nestedOp.getNumRegions() != 0) {
if (hasOwnEffects(&nestedOp)) {
info.requiresInactiveThreadY = true;
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index fd4e12fd039a8..7063239fbc09d 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -9,6 +9,7 @@
// CHECK: %[[IS_X_ZERO:.*]] = arith.cmpi eq, [[TID_X]],
// CHECK-NOT: arith.andi
// CHECK: scf.if %[[IS_X_ZERO]]
+// CHECK: memref.store
// CHECK: acc.atomic.update
func.func @worker_reduction_combine(%result: memref<i32>) {
@@ -35,6 +36,7 @@ func.func @worker_reduction_combine(%result: memref<i32>) {
scf.reduce
} {acc.par_dims = #acc<par_dims[thread_y]>}
acc.predicate_region {
+ memref.store %c0_i32, %local[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
>From b67a8deeddbaaf77d0ef5677dc6f8a1ef274137a Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 20 Jul 2026 21:00:18 -0700
Subject: [PATCH 22/22] [OpenACC][test] Cover active-scope worker store
Exercise a same-source store and worker combine inside an already active ThreadY scope.
---
.../Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir | 1 +
1 file changed, 1 insertion(+)
diff --git a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
index 7063239fbc09d..03da2186987a6 100644
--- a/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-cg-to-gpu-worker-reduction-combine.mlir
@@ -75,6 +75,7 @@ func.func @worker_combine_in_thread_y_scope(%result: memref<i32>) {
scf.parallel (%worker_iv) = (%c0) to (%ty) step (%c1_inner) {
memref.store %c0_i32, %local[] : memref<i32>
acc.predicate_region {
+ memref.store %c0_i32, %local[] : memref<i32>
acc.reduction_combine %local into %result_arg <add> : memref<i32>
{acc.par_dims = #acc<par_dims[block_y, thread_y]>}
}
More information about the Mlir-commits
mailing list