[Mlir-commits] [mlir] [mlir][NVVM] Fix signed and unsigned inference of register ranges (PR #237116)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sun Nov 1 19:52:47 PST 2026
https://github.com/frokaikan updated https://github.com/llvm/llvm-project/pull/237116
>From 21f26bf7a860aba3903078c6953421830966c2f7 Mon Sep 17 00:00:00 2001
From: mingruiz <mingruiz at nvidia.com>
Date: Thu, 29 Oct 2026 23:58:16 -0700
Subject: [PATCH 1/2] [mlir][NVVM] Fix signed and unsigned inference of
register ranges
NVVM special-register ops store their range attribute as an LLVM
ConstantRange: a half-open interval of bit patterns that can wrap.
Integer range analysis consumes ConstantIntRanges instead, whose
signed and unsigned endpoints are inclusive.
The existing inference copied the raw lower and upper endpoints into
both interpretations. This included the excluded upper endpoint and
could form invalid signed or unsigned bounds. In particular, the valid
positive range [1, 0x80000000) gave a signed maximum of INT32_MIN and
allowed gridDim.x < 0 to fold to true, incorrectly selecting a
negative-step loop.
Construct an llvm::ConstantRange and use its unsigned and signed extrema
to populate inclusive bounds. This handles exclusive upper endpoints
and wrapping separately for each interpretation. For an empty set,
use the maximum range as the conservative fallback, since
ConstantIntRanges cannot represent empty sets. Keep the same fallback
when the range attribute is absent.
Extend nvvm-test-range.mlir with regressions for exclusive upper bounds,
the positive grid-dimension limit, singleton ranges, unsigned wrapping,
ranges that cross the sign bit, full sets, and empty sets.
Validation: staged whitespace checks and clang-format 23.1.3 on the
changed C++ lines pass. This llvm-project checkout has not been rebuilt
or its tests rerun after the port.
Assisted-by: OpenAI Codex
---
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 21 +++-
mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir | 95 +++++++++++++++++++
2 files changed, 112 insertions(+), 4 deletions(-)
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index bd3e18946393b2..92bb99e789a14d 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -33,6 +33,7 @@
#include "mlir/Transforms/InliningUtils.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/TypeSwitch.h"
+#include "llvm/IR/ConstantRange.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/NVVMIntrinsicUtils.h"
#include "llvm/Support/Casting.h"
@@ -5950,11 +5951,23 @@ void mlir::NVVM::nvvmInferResultRanges(
ArrayRef<::mlir::ConstantIntRanges> argRanges,
SetIntRangeFn setResultRanges) {
if (range) {
- setResultRanges(result, {range->getLower(), range->getUpper(),
- range->getLower(), range->getUpper()});
- } else {
- setResultRanges(result, IntegerValueRange::getMaxRange(result).getValue());
+ // ConstantRange uses a half-open interval of bit patterns, whereas
+ // ConstantIntRanges uses inclusive signed and unsigned bounds. Derive
+ // each interpretation separately to handle wrapping and sign boundaries.
+ llvm::ConstantRange valueRange(range->getLower(), range->getUpper());
+ if (!valueRange.isEmptySet()) {
+ setResultRanges(result,
+ {valueRange.getUnsignedMin(), valueRange.getUnsignedMax(),
+ valueRange.getSignedMin(), valueRange.getSignedMax()});
+ return;
+ }
}
+ // An empty ConstantRange has no values and cannot be represented by the
+ // inclusive bounds in ConstantIntRanges. Its endpoints must not be treated as
+ // valid minima/maxima. Fall back to the full signed and unsigned ranges, as
+ // when no range attribute is present, to record an unknown value explicitly
+ // and avoid folding comparisons based on inconsistent empty-set bounds.
+ setResultRanges(result, IntegerValueRange::getMaxRange(result).getValue());
}
/// Verify the range attribute satisfies LLVM ConstantRange constructor
diff --git a/mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir b/mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir
index 68a3b72653b055..c61f1edc231ee4 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir
@@ -39,3 +39,98 @@ gpu.module @module{
// CHECK: scf.if %[[false]] {
// CHECK: gpu.printf "threadidz"
// CHECK: arith.select
+
+// The exclusive upper bound is not part of either the signed or unsigned range.
+// CHECK-LABEL: func.func @exclusive_upper_bound
+// CHECK: %[[TRUE:.*]] = arith.constant true
+// CHECK: return %[[TRUE]], %[[TRUE]] : i1, i1
+func.func @exclusive_upper_bound() -> (i1, i1) {
+ %tid = nvvm.read.ptx.sreg.tid.x range <i32, 0, 32> : i32
+ %c32 = arith.constant 32 : i32
+ %signed = arith.cmpi slt, %tid, %c32 : i32
+ %unsigned = arith.cmpi ult, %tid, %c32 : i32
+ return %signed, %unsigned : i1, i1
+}
+
+// [1, 0x80000000) contains only positive signed i32 values. In particular,
+// the exclusive upper bound must not become a negative signed maximum.
+// CHECK-LABEL: func.func @positive_grid_dim
+// CHECK: %[[FALSE:.*]] = arith.constant false
+// CHECK: return %[[FALSE]] : i1
+func.func @positive_grid_dim() -> i1 {
+ %grid = nvvm.read.ptx.sreg.nctaid.x range <i32, 1, -2147483648> : i32
+ %zero = arith.constant 0 : i32
+ %negative = arith.cmpi slt, %grid, %zero : i32
+ return %negative : i1
+}
+
+// CHECK-LABEL: func.func @singleton_range
+// CHECK: %[[ONE:.*]] = arith.constant 1 : i32
+// CHECK: return %[[ONE]] : i32
+func.func @singleton_range() -> i32 {
+ %dim = nvvm.read.ptx.sreg.ntid.y range <i32, 1, 2> : i32
+ return %dim : i32
+}
+
+// This interval wraps in unsigned order but has signed bounds [-4, 3].
+// CHECK-LABEL: func.func @unsigned_wrapped_range
+// CHECK: %[[TRUE:.*]] = arith.constant true
+// CHECK: %[[FOUR:.*]] = arith.constant 4 : i32
+// CHECK: %[[ID:.*]] = nvvm.read.ptx.sreg.ctaid.x range <i32, -4, 4> : i32
+// CHECK: %[[UNSIGNED:.*]] = arith.cmpi ult, %[[ID]], %[[FOUR]] : i32
+// CHECK: return %[[TRUE]], %[[TRUE]], %[[UNSIGNED]] : i1, i1, i1
+func.func @unsigned_wrapped_range() -> (i1, i1, i1) {
+ %id = nvvm.read.ptx.sreg.ctaid.x range <i32, -4, 4> : i32
+ %lower = arith.constant -4 : i32
+ %upper = arith.constant 4 : i32
+ %signed_lower = arith.cmpi sge, %id, %lower : i32
+ %signed_upper = arith.cmpi slt, %id, %upper : i32
+ %unsigned = arith.cmpi ult, %id, %upper : i32
+ return %signed_lower, %signed_upper, %unsigned : i1, i1, i1
+}
+
+// This interval crosses the sign bit. Its unsigned bounds are narrow, but
+// neither sign can be excluded from its signed interpretation.
+// CHECK-LABEL: func.func @signed_wrapped_range
+// CHECK: %[[TRUE:.*]] = arith.constant true
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i32
+// CHECK: %[[ID:.*]] = nvvm.read.ptx.sreg.ctaid.x range <i32, 2147483646, -2147483646> : i32
+// CHECK: %[[NEGATIVE:.*]] = arith.cmpi slt, %[[ID]], %[[ZERO]] : i32
+// CHECK: return %[[TRUE]], %[[TRUE]], %[[NEGATIVE]] : i1, i1, i1
+func.func @signed_wrapped_range() -> (i1, i1, i1) {
+ %id = nvvm.read.ptx.sreg.ctaid.x range <i32, 2147483646, -2147483646> : i32
+ %lower = arith.constant 2147483646 : i32
+ %upper = arith.constant -2147483646 : i32
+ %zero = arith.constant 0 : i32
+ %unsigned_lower = arith.cmpi uge, %id, %lower : i32
+ %unsigned_upper = arith.cmpi ult, %id, %upper : i32
+ %negative = arith.cmpi slt, %id, %zero : i32
+ return %unsigned_lower, %unsigned_upper, %negative : i1, i1, i1
+}
+
+// A full range must not be mistaken for the constant -1.
+// CHECK-LABEL: func.func @full_range
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i32
+// CHECK: %[[ID:.*]] = nvvm.read.ptx.sreg.ctaid.x range <i32, -1, -1> : i32
+// CHECK: %[[NEGATIVE:.*]] = arith.cmpi slt, %[[ID]], %[[ZERO]] : i32
+// CHECK: return %[[NEGATIVE]] : i1
+func.func @full_range() -> i1 {
+ %id = nvvm.read.ptx.sreg.ctaid.x range <i32, -1, -1> : i32
+ %zero = arith.constant 0 : i32
+ %negative = arith.cmpi slt, %id, %zero : i32
+ return %negative : i1
+}
+
+// Empty sets have no ConstantIntRanges representation. Inference falls back
+// to the maximum range rather than constructing inconsistent signed bounds.
+// CHECK-LABEL: func.func @empty_range
+// CHECK: %[[ZERO:.*]] = arith.constant 0 : i32
+// CHECK: %[[ID:.*]] = nvvm.read.ptx.sreg.ctaid.x range <i32, 0, 0> : i32
+// CHECK: %[[NEGATIVE:.*]] = arith.cmpi slt, %[[ID]], %[[ZERO]] : i32
+// CHECK: return %[[NEGATIVE]] : i1
+func.func @empty_range() -> i1 {
+ %id = nvvm.read.ptx.sreg.ctaid.x range <i32, 0, 0> : i32
+ %zero = arith.constant 0 : i32
+ %negative = arith.cmpi slt, %id, %zero : i32
+ return %negative : i1
+}
>From 6209c43cb518129d3a5a2b41535e059c1bdaea92 Mon Sep 17 00:00:00 2001
From: mingruiz <mingruiz at nvidia.com>
Date: Sun, 1 Nov 2026 19:51:04 -0800
Subject: [PATCH 2/2] [mlir][NVVM] Shorten empty-range fallback comment
Keep the explanation of the conservative fallback concise. No functional changes.
Assisted-by: OpenAI Codex
---
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 7 ++-----
1 file changed, 2 insertions(+), 5 deletions(-)
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 92bb99e789a14d..c7adc6c95d3b8c 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5962,11 +5962,8 @@ void mlir::NVVM::nvvmInferResultRanges(
return;
}
}
- // An empty ConstantRange has no values and cannot be represented by the
- // inclusive bounds in ConstantIntRanges. Its endpoints must not be treated as
- // valid minima/maxima. Fall back to the full signed and unsigned ranges, as
- // when no range attribute is present, to record an unknown value explicitly
- // and avoid folding comparisons based on inconsistent empty-set bounds.
+ // ConstantIntRanges cannot represent an empty set. Conservatively fall back
+ // to the full range rather than deriving invalid inclusive bounds.
setResultRanges(result, IntegerValueRange::getMaxRange(result).getValue());
}
More information about the Mlir-commits
mailing list