Skip to content

Commit da70146

Browse files
jlebarjlebar
committed
Add NVPTX results.
1 parent 231179d commit da70146

449 files changed

Lines changed: 13760 additions & 3 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎NVPTX/README.md‎

Lines changed: 182 additions & 0 deletions
Large diffs are not rendered by default.

‎NVPTX/bugs/001-fptoui-fptosi-to-i1-compares-eq-zero/NOTES.md‎

Lines changed: 109 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,109 @@
1+
# 001 — fptosi/fptoui to i1 lowered as (a == 0.0) instead of LSB(trunc(a)), miscompiling all float->i1 conversions
2+
3+
- **Kind:** miscompile
4+
- **Reachable via:** default llc -O0/-O2
5+
- **Component:** NVPTXInstrInfo.td 2109,2115,2121,2127,2137,2148,2154,2160 (region `T1-instrinfo-td`)
6+
- **Candidate id:** c008
7+
8+
## Summary
9+
10+
`fptoui`/`fptosi` of any float to `i1` is lowered as `(a == 0.0)` instead of `trunc(a)&1`
11+
12+
## Mechanism / root cause
13+
14+
The patterns for converting a floating-point value directly to i1 use an integer setp-equal-to-zero on the raw float bits:
15+
16+
def : Pat<(i1 (fp_to_sint f16:$a)), (SETP_i16ri $a, 0, CmpEQ)>;
17+
def : Pat<(i1 (fp_to_uint f16:$a)), (SETP_i16ri $a, 0, CmpEQ)>;
18+
def : Pat<(i1 (fp_to_sint f32:$a)), (SETP_i32ri $a, 0, CmpEQ)>;
19+
def : Pat<(i1 (fp_to_uint f32:$a)), (SETP_i32ri $a, 0, CmpEQ)>;
20+
def : Pat<(i1 (fp_to_sint f64:$a)), (SETP_i64ri $a, 0, CmpEQ)>;
21+
... (also bf16 and fp_to_uint f64)
22+
23+
SETP_iNN with CmpEQ (CmpMode 0 -> PTXCmpMode::EQ) prints as 'setp.eq.bNN $a, 0', an INTEGER/bitwise comparison of the float's bit pattern against 0. So the emitted i1 result is (a_bits == 0), i.e. (a == +0.0).
24+
25+
LLVM semantics of fptoui/fptosi to i1 require the result to be the low bit of the trunc-toward-zero integer value of a (with poison only when that integer does not fit in i1, i.e. trunc(a) outside {0,1} for unsigned or {0,-1} for signed). For example fptoui(1.5) to i1 = trunc(1.5)=1 -> i1 = 1 (a defined input). The pattern instead yields (1.5 == 0) = 0. And fptoui(0.0) to i1 = 0, but the pattern yields (0.0 == 0) = 1. Both are wrong.
26+
27+
Verified with llc: for 'fptoui float %a to i1' the backend emits 'setp.eq.b32 %p1, %r1, 0; selp.b32 %r2, -1/1, 0, %p1', i.e. result = (a==0). x86 reference for the identical IR emits cvttss2si and takes the low bit (AL), confirming the correct value is LSB(trunc(a)), not (a==0). The same defect was confirmed for f16 (setp.eq.b16), f32 (setp.eq.b32), and f64 (setp.eq.b64). No test in test/CodeGen/NVPTX covers float->i1, which is why this has gone unnoticed.
28+
29+
## Trigger
30+
31+
Any nvptx64 target; an fptoui or fptosi from f16/bf16/f32/f64 directly to i1 with a runtime float operand. Concrete defined inputs that differ: a=1.5 (correct 1, emitted 0) and a=0.0 (correct 0, emitted 1).
32+
33+
## Reproducer
34+
35+
See `repro.ll` / `cmd.sh`.
36+
37+
```
38+
target triple = "nvptx64-nvidia-cuda"
39+
40+
define i32 @f2u_i1_zext(float %a) {
41+
%r = fptoui float %a to i1
42+
%z = zext i1 %r to i32
43+
ret i32 %z
44+
}
45+
46+
define i32 @f2s_i1_zext(float %a) {
47+
%r = fptosi float %a to i1
48+
%z = zext i1 %r to i32
49+
ret i32 %z
50+
}
51+
```
52+
53+
Command:
54+
55+
```
56+
llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll
57+
```
58+
59+
## Observed (wrong) output
60+
61+
```
62+
.visible .func (.param .b32 func_retval0) f2u_i1_zext( .param .b32 f2u_i1_zext_param_0 )
63+
{
64+
.reg .pred %p<2>;
65+
.reg .b32 %r<3>;
66+
ld.param.b32 %r1, [f2u_i1_zext_param_0];
67+
setp.eq.b32 %p1, %r1, 0;
68+
selp.b32 %r2, 1, 0, %p1;
69+
st.param.b32 [func_retval0], %r2;
70+
ret;
71+
}
72+
73+
(f2s_i1_zext is identical: setp.eq.b32 %p1, %r1, 0; selp.b32 %r2, 1, 0, %p1)
74+
75+
So for a=1.5f the function returns selp(1.5_bits==0 ? 1 : 0) = 0 (correct is 1); for a=0.0f it returns 1 (correct is 0). The float bits are loaded with ld.param.b32 into an integer register and compared bitwise to 0 with no cvt.
76+
```
77+
78+
## Expected
79+
80+
Result should be LSB(trunc_toward_zero(a)). Reference x86_64 codegen for the same IR:
81+
f2u_i1_zext: cvttss2si %xmm0, %eax ; movzbl %al, %eax ; ret (=> a=1.5 -> 1, a=0.0 -> 0)
82+
f2s_i1_zext: cvttss2si %xmm0, %eax ; movzbl %al, %eax ; andl $1, %eax ; ret
83+
Correct NVPTX would convert with round-toward-zero to an integer (e.g. cvt.rzi.s32.f32) and take/AND the low bit (analogous to the existing trunc-to-i1 pattern at line 2209: SETP_i32ri (AND_b32ri $a, 1), 0, CmpNE), NOT compare the float's bit pattern to 0. The current code computes (a == +0.0).
84+
85+
## Verification
86+
87+
Verified empirically with the built llc (independent verify + adversarial
88+
refute both `confirmed`, finder confidence 0.85, verify confidence 0.98).
89+
90+
> Confirmed real miscompile. NVPTXInstrInfo.td lines 2109/2115/2121/2127/2137/2148/2154/2160 match (i1 (fp_to_sint/fp_to_uint fXX:$a)) to (SETP_iNNri $a, 0, CmpEQ). The ISETP multiclass (lines 1634-1657) prints "setp.${cmp:ICmp}.${cmp:IType}NN", so CmpEQ -> "setp.eq.bNN $a, 0", a BITWISE integer comparison of the float register's raw bits against 0 — the float operand is consumed directly with no cvt. The emitted i1 is therefore (bitcast(a) == 0) == (a == +0.0).
91+
92+
LangRef requires fptoui/fptosi to i1 = LSB of trunc-toward-zero(a), with poison only when that integer doesn't fit i1 (outside {0,1} unsigned, {0,-1} signed). I exhibited two fully well-defined inputs (no poison/UB) where NVPTX is wrong:
93+
- a=0.0f: trunc=0 fits -> correct result 0; NVPTX computes (0x00000000==0)=1. WRONG.
94+
- a=1.5f: trunc=1 fits in unsigned i1 -> correct result 1; NVPTX computes (0x3FC00000==0)=0. WRONG.
95+
Both verified numerically (bit patterns 0x00000000 and 0x3FC00000) and against an x86_64 reference for the identical IR, which emits cvttss2si + low-bit extract (movzbl / andl $1) — LSB(trunc_toward_zero(a)) — and disagrees with NVPTX on both inputs. Note -0.0 (0x80000000) happens to coincidentally agree (NVPTX=0, correct=0) because its bits are nonzero. The defect spans f16/bf16/f32/f64, signed and unsigned. No NVPTX test covers float->i1 (only i16/i32/i64), which is why it went unnoticed. I considered and ruled out: the inputs are not UB/poison; nothing masks/normalizes the result (it is the direct r
96+
97+
## Independent cross-check
98+
99+
`opt -O2` constant-folds the truth:
100+
- `fptoui float 1.5 to i1` → **1** (trunc(1.5)=1, low bit 1)
101+
- `fptoui float 0.0 to i1` → **0**
102+
103+
NVPTX emits `setp.eq.b32 %p, a, 0` ⇒ result `(a == 0.0)`:
104+
- a=1.5 → 0 (should be 1) ✗
105+
- a=0.0 → 1 (should be 0) ✗
106+
107+
x86 reference codegen for the same IR is `cvttss2si %xmm0, %eax` (convert toward
108+
zero, take the low bit) — i.e. `LSB(trunc(a))`, confirming the correct lowering
109+
is the truncated-integer low bit, not an equality-with-zero test.

‎NVPTX/bugs/001-fptoui-fptosi-to-i1-compares-eq-zero/cmd.sh‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
#!/usr/bin/env bash
2+
set -uo pipefail
3+
TOOL="${TOOL:-/Users/justinlebar/code/llvm2/build/bin/llc}"
4+
cd "$(dirname "$0")"
5+
echo "===== llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll ====="
6+
"$TOOL" -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll
7+
echo "exit=$?"

‎NVPTX/bugs/001-fptoui-fptosi-to-i1-compares-eq-zero/repro.ll‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
target triple = "nvptx64-nvidia-cuda"
2+
3+
define i32 @f2u_i1_zext(float %a) {
4+
%r = fptoui float %a to i1
5+
%z = zext i1 %r to i32
6+
ret i32 %z
7+
}
8+
9+
define i32 @f2s_i1_zext(float %a) {
10+
%r = fptosi float %a to i1
11+
%z = zext i1 %r to i32
12+
ret i32 %z
13+
}

‎NVPTX/bugs/002-combinemulwide-sext-shl-topbit-negates/NOTES.md‎

Lines changed: 87 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,87 @@
1+
# 002 — combineMulWide: sext(shl nsw x, topbit) -> mul.wide.s with negative constant miscompiles
2+
3+
- **Kind:** miscompile
4+
- **Reachable via:** default llc -O1+
5+
- **Component:** NVPTXISelLowering.cpp 6379-6385 (region `L01-combines-mul-shl`)
6+
- **Candidate id:** c003
7+
8+
## Summary
9+
10+
`sext(shl nsw x, bitwidth-1)` folded to `mul.wide.s` with a negative constant, negating the result
11+
12+
## Mechanism / root cause
13+
14+
combineMulWide rewrites (sign_extend (shl nsw x, ShiftAmt)) into MUL_WIDE_SIGNED(x, 1<<ShiftAmt). For the SHL case it builds the multiplier constant in the *narrow* type FromVT:
15+
16+
if (Op.getOpcode() == ISD::SHL) {
17+
const auto ShiftAmt = Op.getConstantOperandVal(1);
18+
const auto MulVal = APInt(FromVT.getSizeInBits(), 1) << ShiftAmt; // narrow width!
19+
RHS = DCI.DAG.getConstant(MulVal, DL, FromVT);
20+
}
21+
return DCI.DAG.getNode(Opcode, DL, ToVT, Op.getOperand(0), RHS);
22+
23+
When ShiftAmt == FromVT.bits-1 (15 for i16, 31 for i32), MulVal sets the high (sign) bit of the narrow constant: for i16, 1<<15 = 0x8000. MUL_WIDE_SIGNED (mul.wide.s16) sign-extends BOTH operands, so the 0x8000 constant is interpreted as the *signed* value -32768, not the intended positive multiplier 2^15 = +32768. The product therefore has the wrong sign.
24+
25+
Concretely, the well-defined (non-poison) input x = -1 (0xFFFF) for `shl nsw i16 %x, 15`: original semantics give shl = 0x8000 = -32768, sext to i32 = -32768 (0xFFFF8000). opt constant-folds `sext(shl nsw i16 -1, 15)` to exactly -32768, confirming the input is defined and the correct result is -32768. The emitted PTX is `mul.wide.s16 %r, %rs, -32768`, which for %rs = -1 computes (-1)*(-32768) = +32768 (0x00008000). -32768 != +32768 -> miscompile. The i32->i64 case (ShiftAmt 31) is identical: emitted `mul.wide.s32 %rd, %r, -2147483648`. The ZERO_EXTEND path is fine (mul.wide.u treats the constant as unsigned). The older TryMULWIDECombine avoids this because it builds the constant at full width and rejects it via Val.isSignedIntN(OptSize) (32768 doesn't fit signed i16). The fix is to bail out (or use the unsigned widening, but signedness is fixed by the extend) when 1<<ShiftAmt does not fit as a signed value in DemotedVT, i.e. when ShiftAmt >= FromVT.bits-1 for the signed case.
26+
27+
## Trigger
28+
29+
nvptx64, any sm/ptx, -O1+ (OptLevel != None). Pattern: sext i32 (shl nsw i16 x, 15) to i32, or sext i64 (shl nsw i32 x, 31) to i64. Triggers for the defined input x = -1 (and any x where the shl-nsw is non-poison, e.g. x=-1). Reached via PerformDAGCombine ISD::SIGN_EXTEND -> combineMulWide (line 7100).
30+
31+
## Reproducer
32+
33+
See `repro.ll` / `cmd.sh`.
34+
35+
```
36+
define i32 @sext_shl15(i16 %x) {
37+
%s = shl nsw i16 %x, 15
38+
%e = sext i16 %s to i32
39+
ret i32 %e
40+
}
41+
42+
define i64 @sext_shl31(i32 %x) {
43+
%s = shl nsw i32 %x, 31
44+
%e = sext i32 %s to i64
45+
ret i64 %e
46+
}
47+
```
48+
49+
Command:
50+
51+
```
52+
llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll
53+
```
54+
55+
## Observed (wrong) output
56+
57+
```
58+
// sext_shl15:
59+
ld.param.b16 %rs1, [sext_shl15_param_0];
60+
mul.wide.s16 %r1, %rs1, -32768;
61+
st.param.b32 [func_retval0], %r1;
62+
ret;
63+
// sext_shl31:
64+
ld.param.b32 %r1, [sext_shl31_param_0];
65+
mul.wide.s32 %rd1, %r1, -2147483648;
66+
st.param.b64 [func_retval0], %rd1;
67+
ret;
68+
69+
For input x = -1: mul.wide.s16(-1, -32768) = (-1)*(-32768) = +32768 (0x00008000), and mul.wide.s32(-1, -2147483648) = +2147483648.
70+
```
71+
72+
## Expected
73+
74+
sext_shl15(-1) should return -32768 (0xFFFF8000); opt constant-folds sext(shl nsw i16 -1, 15) to -32768. sext_shl31(-1) should return -2147483648. A correct lowering is the multiplier as a POSITIVE wide constant, e.g. equivalent to sext(x) then shl by 15, which llc emits correctly for the reference IR `mul i32 (sext x), 32768` as `ld.param.s16; shl.b32 %r, %r, 15`. The emitted negative immediates (-32768, -2147483648) into mul.wide.s flip the product's sign for negative x, so the PTX produces +32768 / +2147483648 instead of -32768 / -2147483648.
75+
76+
## Verification
77+
78+
Verified empirically with the built llc (independent verify + adversarial
79+
refute both `confirmed`, finder confidence 0.92, verify confidence 0.98).
80+
81+
> Confirmed real miscompile. Mechanism verified at NVPTXISelLowering.cpp:6380-6385: combineMulWide rewrites (sign_extend (shl nsw x, ShiftAmt)) into NVPTXISD::MUL_WIDE_SIGNED(x, 1<<ShiftAmt), building the multiplier constant in the narrow type FromVT via APInt(FromVT.getSizeInBits(), 1) << ShiftAmt. When ShiftAmt == FromVT.bits-1 (15 for i16, 31 for i32), the constant's sign bit is set: 1<<15 = 0x8000, emitted into PTX as the signed value -32768. mul.wide.s16/s32 sign-extend BOTH operands, so the immediate is interpreted as -32768, not the intended +32768.
82+
83+
Concrete defined input x = -1 (i16 0xFFFF): IR semantics of sext(shl nsw i16 -1, 15) = -32768 (0xFFFF8000). This input is genuinely non-poison: -1<<15 mathematically = -32768, which fits in signed i16 so nsw is satisfied; opt -passes=sccp,instcombine,instsimplify folds it to -32768. The emitted PTX `mul.wide.s16 %r1, %rs1, -32768` computes (-1)*(-32768) = +32768 (0x00008000). -32768 != +32768 -> wrong sign, wrong value. The i32->i64 case (shl 31 -> mul.wide.s32 with -2147483648) is identical: for x=-1 it yields +2147483648 vs correct -2147483648.
84+
85+
PTX mul.wide.s16/s32 semantics confirmed (NVIDIA PTX ISA section 9.7.1.3 / general spec): signed n-bit operands sign-extended to full 2n-bit product. The immediate -32768 is the signed 16-bit value 0x8000, sign-extended to 0xFFFF8000.
86+
87+
Adversarial checks all pass: (1) input is defined/non-poison (no UB, no poison); (2) the ZERO_EXTEND path is fine since mul.wide.u treats the consta

‎NVPTX/bugs/002-combinemulwide-sext-shl-topbit-negates/cmd.sh‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
#!/usr/bin/env bash
2+
set -uo pipefail
3+
TOOL="${TOOL:-/Users/justinlebar/code/llvm2/build/bin/llc}"
4+
cd "$(dirname "$0")"
5+
echo "===== llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll ====="
6+
"$TOOL" -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll
7+
echo "exit=$?"

‎NVPTX/bugs/002-combinemulwide-sext-shl-topbit-negates/repro.ll‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
define i32 @sext_shl15(i16 %x) {
2+
%s = shl nsw i16 %x, 15
3+
%e = sext i16 %s to i32
4+
ret i32 %e
5+
}
6+
7+
define i64 @sext_shl31(i32 %x) {
8+
%s = shl nsw i32 %x, 31
9+
%e = sext i32 %s to i64
10+
ret i64 %e
11+
}

‎NVPTX/bugs/003-selectshiftcombine-i64-clamp-drops-high32/NOTES.md‎

Lines changed: 92 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,92 @@
1+
# 003 — PerformSELECTShiftCombine folds guarded i64 shift to clamp shift, dropping the high 32 bits of the guard's shift amount
2+
3+
- **Kind:** miscompile
4+
- **Reachable via:** default llc (after legalize)
5+
- **Component:** NVPTXISelLowering.cpp 6732-6770 (region `L03b-combines-setcc-extract`)
6+
- **Candidate id:** c006
7+
8+
## Summary
9+
10+
guarded i64 shift folded to PTX clamp shift using only the low 32 bits of the shift amount
11+
12+
## Mechanism / root cause
13+
14+
PerformSELECTShiftCombine rewrites `select (icmp ult shift_amt, BW), (x << shift_amt), 0` (and the ugt variant) into `NVPTXISD::SHL_CLAMP/SRL_CLAMP(x, ShiftOp.getOperand(1))`, relying on PTX shl/shr clamping amounts >= BitWidth to 0. The bug: PTX shift instructions take a 32-bit `.u32` shift-amount operand. For an i64 shift, ShiftOp.getOperand(1) is the shift amount TRUNCATED to i32 (the matcher even uses m_TruncOrSelf(m_Deferred(ShiftAmt)) at lines 6745-6746 to peek through that truncate). The guard `icmp ult shift_amt, 64` however compares the FULL i64 value. The combine emits the clamp using only the truncated low-32-bit amount (lines 6766-6769: `ShiftOp.getOperand(1)`), so the high 32 bits that the guard tested are silently discarded.
15+
16+
Concretely, with shift_amt = 0x0000000100000005 (i64): the guard `icmp ult i64 0x100000005, 64` is FALSE, so the IR `select` returns its `0` operand -> result 0 (the poison shl is in the unselected arm, so the select is well-defined). The generated PTX does `ld.param.b32 %r1` (low 32 bits = 5), then `shl.b64 %rd2, %rd1, %r1`, which clamps 5 < 64 and shifts: x << 5. For x=1 the IR yields 0 but the PTX yields 32.
17+
18+
Verified with the built llc: `guarded_shl_i64` and `guarded_srl_i64_ugt` both emit an unguarded `shl.b64/shr.u64` with the shift amount loaded as `.b32`, dropping the high half. The existing test test_guarded_i64_ult_shl in shift-opt.ll exhibits the same unguarded codegen, confirming the combine fires for i64. The combine is reached for ISD::SELECT (line 7141) with only an isAfterLegalizeDAG gate (no OptLevel gate), so it fires even at -O0. The i32/i64-with-i32-amount cases are safe; only shift-amount types wider than 32 bits (i64, i128) where the low 32 bits are < BitWidth but the high bits are nonzero are miscompiled. The fix needs to ensure the guard's compared value and the clamp's 32-bit amount agree on the dropped high bits (e.g. require ShiftAmt to be <= 32 bits, or that its high bits are known zero).
19+
20+
## Trigger
21+
22+
nvptx64, any sm/ptx, integer type with shift-amount wider than 32 bits (i64 is the common case). Pattern `shift >= 64 ? 0 : x >> shift` or `shift < 64 ? x << shift : 0` with an i64 shift amount whose high 32 bits are nonzero but low 32 bits are < 64. Fires at all opt levels (only gated on isAfterLegalizeDAG).
23+
24+
## Reproducer
25+
26+
See `repro.ll` / `cmd.sh`.
27+
28+
```
29+
target triple = "nvptx64-nvidia-cuda"
30+
31+
; IR semantics: when %shift >= 64 the icmp is false, so select returns 0
32+
; (the poison shl is in the unselected arm). For %shift=0x100000005 -> result 0.
33+
define i64 @guarded_shl_i64(i64 %x, i64 %shift) {
34+
%cmp = icmp ult i64 %shift, 64
35+
%shl = shl i64 %x, %shift
36+
%sel = select i1 %cmp, i64 %shl, i64 0
37+
ret i64 %sel
38+
}
39+
40+
; IR semantics: when %shift > 63 the icmp is true, so select returns 0.
41+
define i64 @guarded_srl_i64_ugt(i64 %x, i64 %shift) {
42+
%cmp = icmp ugt i64 %shift, 63
43+
%shr = lshr i64 %x, %shift
44+
%sel = select i1 %cmp, i64 0, i64 %shr
45+
ret i64 %sel
46+
}
47+
```
48+
49+
Command:
50+
51+
```
52+
llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll
53+
```
54+
55+
## Observed (wrong) output
56+
57+
```
58+
// guarded_shl_i64:
59+
ld.param.b64 %rd1, [guarded_shl_i64_param_0];
60+
ld.param.b32 %r1, [guarded_shl_i64_param_1]; // loads only low 32 bits of the i64 shift amount
61+
shl.b64 %rd2, %rd1, %r1; // unguarded shift; select/guard gone
62+
st.param.b64 [func_retval0], %rd2;
63+
ret;
64+
65+
// guarded_srl_i64_ugt:
66+
ld.param.b64 %rd1, [guarded_srl_i64_ugt_param_0];
67+
ld.param.b32 %r1, [guarded_srl_i64_ugt_param_1]; // low 32 bits only
68+
shr.u64 %rd2, %rd1, %r1; // unguarded shift
69+
st.param.b64 [func_retval0], %rd2;
70+
ret;
71+
72+
// ISel debug confirms: t18: i64 = NVPTXISD::SHL_CLAMP t4, t20, where
73+
// t20: i32,ch = load<(... load (s32), align 8 ...)> of guarded_shl_i64_param_1
74+
// i.e. the clamp amount is the low 32 bits of the i64 guard value.
75+
76+
// For x=1, shift=0x0000000100000005: PTX yields 1 << 5 = 32.
77+
```
78+
79+
## Expected
80+
81+
For x=1, shift=0x0000000100000005 the IR select returns 0 (icmp ult 4294967301, 64 is false), so a correct compiler must produce 0, not 32. The combine must only fire when the guard's compared value and the 32-bit clamp amount agree on the dropped high bits — e.g. require the shift-amount type to be <= 32 bits, or that the high bits above bit 31 are known zero (knownBits), before rewriting to SHL_CLAMP/SRL_CLAMP. Equivalent correct PTX would keep the guard, e.g. compare the full i64 amount against 64 (setp.ge.u64 over the 64-bit amount) and selp 0, or otherwise not drop the high bits. For the i32-amount case the combine is already correct.
82+
83+
## Verification
84+
85+
Verified empirically with the built llc (independent verify + adversarial
86+
refute both `confirmed`, finder confidence 0.85, verify confidence 0.97).
87+
88+
> Confirmed real miscompile. PerformSELECTShiftCombine (NVPTXISelLowering.cpp:6732-6770) rewrites `select (icmp ult shift_amt, BW), (x<<shift_amt), 0` (and ugt variant) into NVPTXISD::SHL_CLAMP/SRL_CLAMP(x, ShiftOp.getOperand(1)). The guard's icmp compares the FULL i64 shift amount (ShiftAmt bound in m_SetCC), but the clamp's amount operand is the shift's operand, which for an i64 shift is the amount truncated to i32 (the matcher peeks through that trunc via m_TruncOrSelf(m_Deferred(ShiftAmt)) at 6745-6746). The SHIFT multiclass (NVPTXInstrInfo.td:1393-1408) takes the amount as B32:$b ("shift is always int32"), so PTX shl.b64/shr.u64 only see the low 32 bits.
89+
90+
Verified codegen with the built llc: both guarded_shl_i64 and guarded_srl_i64_ugt emit an unguarded shl.b64/shr.u64 with the amount loaded via ld.param.b32 (low 32 bits of the i64 param), and the select/guard is entirely gone. ISel debug shows SHL_CLAMP64_rr / SRL_CLAMP64_rr consuming a 32-bit load (s32) of the i64 shift-amount parameter, dropping the high 32 bits.
91+
92+
Concrete defined-input discrepancy for guarded_shl_i64 with x=1, shift=0x0000000100000005 (=4294967301): IR `icmp ult 4294967301, 64` is FALSE, so select returns its 0 operand (the poison shl is in the UNSELECTED arm, so the select is well-defined) => IR result 0. PTX loads low 32 bits = 5, computes shl.b64 1,5 = 32. 0 != 32. Adversarial checks: (a) the poison shl is harmless because select picks the other arm; (b) the i32-amount variant is safe because guard

‎NVPTX/bugs/003-selectshiftcombine-i64-clamp-drops-high32/cmd.sh‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
#!/usr/bin/env bash
2+
set -uo pipefail
3+
TOOL="${TOOL:-/Users/justinlebar/code/llvm2/build/bin/llc}"
4+
cd "$(dirname "$0")"
5+
echo "===== llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll ====="
6+
"$TOOL" -mtriple=nvptx64 -mcpu=sm_90 -O2 -o - repro.ll
7+
echo "exit=$?"

0 commit comments

Comments
 (0)