[Mlir-commits] [mlir] [mlir][NVVM] Fix signed and unsigned inference of register ranges (PR #237116)

[llvmlistbot at llvm.org]

================
@@ -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()) {
----------------
frokaikan wrote:

This fixes the shared range conversion for all special-register ops using this helper. The reproducer specifically involves nctaid.x (gridDim.x): its range [1, 0x80000000) contains only positive signed values, but the old conversion incorrectly treated the exclusive upper bound as a negative signed maximum.

Smaller ranges are also affected by the exclusive/inclusive mismatch. For example, tid.x with range [0, 32) should have an inclusive maximum of 31, whereas the old code used 32. The sign-boundary case exposed a miscompilation; the smaller-range cases lose precision.

https://github.com/llvm/llvm-project/pull/237116