[Mlir-commits] [mlir] [mlir][NVVM] Fix signed and unsigned inference of register ranges (PR #237116)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Oct 30 00:02:13 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir
Author: frokaikan
<details>
<summary>Changes</summary>
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.
AI assistance: This patch and its regression tests were developed with assistance from OpenAI Codex. I have reviewed and understand all submitted changes, and I take full responsibility for them.
---
Full diff: https://github.com/llvm/llvm-project/pull/237116.diff
2 Files Affected:
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+17-4)
- (modified) mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir (+95)
``````````diff
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index bd3e18946393b..92bb99e789a14 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 68a3b72653b05..c61f1edc231ee 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
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/237116
More information about the Mlir-commits
mailing list