[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:01:34 PDT 2026
https://github.com/frokaikan created https://github.com/llvm/llvm-project/pull/237116
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.
>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] [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
+}
More information about the Mlir-commits
mailing list