[Mlir-commits] [mlir] [mlir][gpu] Derive has-redux from the chip in the NVVM pipeline (PR #221393)

[Alessandro Potenza]

https://github.com/alepot55 updated https://github.com/llvm/llvm-project/pull/221393

>From 3113ede2c2c2cb5b35458aa3b70eedbce4c70277 Mon Sep 17 00:00:00 2001
From: Alessandro Potenza <ap.alessandro.potenza at gmail.com>
Date: Thu, 3 Sep 2026 00:21:31 +0200
Subject: [PATCH 1/2] [mlir][gpu] Derive has-redux from the chip in the NVVM
 pipeline

gpu-lower-to-nvvm-pipeline cannot lower a gpu.subgroup_reduce. The only
pattern for it in convert-gpu-to-nvvm is the one that emits nvvm.redux.sync,
and that pattern is added only when the has-redux option is set, which the
pipeline does not expose and therefore never sets. The pipeline runs no shuffle
expansion either, so the operation reaches the conversion target as an illegal
operation and the pipeline fails.

Nothing about that option is a decision of its own: redux.sync exists from
sm_80, and the pipeline already knows the chip it is serializing for. Derive it
from cubin-chip. Below sm_80 the pipeline still cannot lower the operation, for
the second reason above.
---
 .../GPU/Pipelines/GPUToNVVMPipeline.cpp       |  9 +++++++++
 .../GPU/CUDA/subgroup-reduce-redux.mlir       | 20 +++++++++++++++++++
 2 files changed, 29 insertions(+)
 create mode 100644 mlir/test/Integration/GPU/CUDA/subgroup-reduce-redux.mlir

diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
index 5462cddd44718d..6629d8b065b71f 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
@@ -30,6 +30,7 @@
 #include "mlir/Pass/PassManager.h"
 #include "mlir/Pass/PassOptions.h"
 #include "mlir/Transforms/Passes.h"
+#include "llvm/ADT/StringExtras.h"
 
 using namespace mlir;
 
@@ -73,6 +74,14 @@ void buildGpuPassPipeline(OpPassManager &pm,
   opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv;
   opt.indexBitwidth = options.indexBitWidth;
   opt.allowPatternRollback = options.allowPatternRollback;
+  // redux.sync exists from sm_80, and the chip this pipeline serializes for is
+  // the only thing that decides whether it may be used, so there is no reason
+  // to ask for it as an option of its own.
+  StringRef chip(options.cubinChip);
+  unsigned smVersion = 0;
+  if (chip.consume_front("sm_"))
+    chip.take_while(llvm::isDigit).getAsInteger(10, smVersion);
+  opt.hasRedux = smVersion >= 80;
   pm.addNestedPass<gpu::GPUModuleOp>(createConvertGpuOpsToNVVMOps(opt));
   pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass());
   pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
diff --git a/mlir/test/Integration/GPU/CUDA/subgroup-reduce-redux.mlir b/mlir/test/Integration/GPU/CUDA/subgroup-reduce-redux.mlir
new file mode 100644
index 00000000000000..c3d0770c162c1f
--- /dev/null
+++ b/mlir/test/Integration/GPU/CUDA/subgroup-reduce-redux.mlir
@@ -0,0 +1,20 @@
+// RUN: mlir-opt %s \
+// RUN:  | mlir-opt -gpu-lower-to-nvvm-pipeline="cubin-chip=sm_80 cubin-format=isa" \
+// RUN:  | FileCheck %s
+
+// A subgroup reduction the source states is executed by the whole subgroup
+// becomes one redux.sync instruction, once the chip the pipeline serializes
+// for is one that has it. Below sm_80 the pipeline still cannot lower the
+// operation at all, since it runs no shuffle expansion either.
+
+// CHECK: redux.sync.add.s32
+
+gpu.module @kernels {
+  gpu.func @reduce(%out: memref<32xi32>) kernel {
+    %tid = gpu.thread_id x
+    %v = arith.index_cast %tid : index to i32
+    %r = gpu.subgroup_reduce add %v uniform : (i32) -> (i32)
+    memref.store %r, %out[%tid] : memref<32xi32>
+    gpu.return
+  }
+}

>From 37ceadb53eb7b1894112b6c7a646e7408388467c Mon Sep 17 00:00:00 2001
From: Alessandro Potenza <ap.alessandro.potenza at gmail.com>
Date: Mon, 14 Sep 2026 09:10:43 +0200
Subject: [PATCH 2/2] Outline the chip check into chipHasRedux

---
 .../GPU/Pipelines/GPUToNVVMPipeline.cpp       | 20 +++++++++++--------
 1 file changed, 12 insertions(+), 8 deletions(-)

diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
index 6629d8b065b71f..dd465a1a1c51c2 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
@@ -68,20 +68,24 @@ void buildCommonPassPipeline(
 //===----------------------------------------------------------------------===//
 // GPUModule-specific stuff.
 //===----------------------------------------------------------------------===//
+/// Whether `chip`, the NVPTX target this pipeline serializes for (an "sm_NN"
+/// name, possibly with a suffix such as "sm_90a"), has redux.sync, which exists
+/// from sm_80. The chip is the only thing that decides it, so the pipeline does
+/// not ask for it as an option of its own.
+static bool chipHasRedux(StringRef chip) {
+  unsigned smVersion = 0;
+  if (chip.consume_front("sm_"))
+    chip.take_while(llvm::isDigit).getAsInteger(10, smVersion);
+  return smVersion >= 80;
+}
+
 void buildGpuPassPipeline(OpPassManager &pm,
                           const mlir::gpu::GPUToNVVMPipelineOptions &options) {
   ConvertGpuOpsToNVVMOpsOptions opt;
   opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv;
   opt.indexBitwidth = options.indexBitWidth;
   opt.allowPatternRollback = options.allowPatternRollback;
-  // redux.sync exists from sm_80, and the chip this pipeline serializes for is
-  // the only thing that decides whether it may be used, so there is no reason
-  // to ask for it as an option of its own.
-  StringRef chip(options.cubinChip);
-  unsigned smVersion = 0;
-  if (chip.consume_front("sm_"))
-    chip.take_while(llvm::isDigit).getAsInteger(10, smVersion);
-  opt.hasRedux = smVersion >= 80;
+  opt.hasRedux = chipHasRedux(options.cubinChip);
   pm.addNestedPass<gpu::GPUModuleOp>(createConvertGpuOpsToNVVMOps(opt));
   pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass());
   pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());