[BUG] NF4 Linear4bit wrong on gfx1100 (prebuilt ROCm 6.4 wheel); dequantize correct

System Info

• Radeon Pro W7800, gfx1100 (RDNA3); host ROCm 7.13; PyTorch 2.9.1+rocm6.4
• bitsandbytes 0.50.3.dev0 prebuilt wheel; loader uses libbitsandbytes_rocm64.so
• Linear4bit output is unrelated to the reference; quantize/dequantize is correct (~9% error). The fused 4-bit GEMM is wrong, not the quantizer.

Reproduction

import torch, bitsandbytes as bnb, bitsandbytes.functional as F
torch.manual_seed(0); dev="cuda"; dt=torch.bfloat16
x=torch.randn(8,512,dtype=dt,device=dev); w=torch.randn(256,512,dtype=dt,device=dev)
w4,qs=F.quantize_4bit(w,quant_type="nf4",compress_statistics=True)
print("dequant", ((F.dequantize_4bit(w4,qs)-w).abs().mean()/w.abs().mean()).item())   # ~0.09
q=bnb.nn.Linear4bit(512,256,bias=False,quant_type="nf4",compute_dtype=dt,device="meta")
q.weight=bnb.nn.Params4bit(w,quant_type="nf4",requires_grad=False); q=q.to(dev)
print("linear4bit", ((q(x)[email protected]()).abs().mean()/([email protected]()).abs().mean()).item())      # ~1.0

Expected behavior

Expect both errors small (quantization-only, ~0.09); Linear4bit is ~1.0 instead.

Narrowing: prebuilt rocm64 wrong; same tree built with -DCOMPUTE_BACKEND=hip -DBNB_ROCM_ARCH=gfx1100 (ROCm 7.13, clang 23) correct (verified on a 27B NF4 model); gfx1100 is already in the release matrix.

Likely cause: RDNA3-only SIMT path in csrc/gemm_4bit_simt.cu (IS_RDNA3). Disassembly of the prebuilt gfx1100 code object shows a float↔int conversion around the int-typed __builtin_amdgcn_mov_dpp in the warp reduction (v_cvt_i32_f32_dpp truncating); clang 23 emits bit-preserving v_add_f32_dpp. Fix: explicit __float_as_uint/__uint_as_float bitcast there.

CI gap: no ROCm GPU correctness tests.

Workaround: source-build with -DBNB_ROCM_ARCH=gfx1100, or bypass the fused op via dequantize_4bit + F.linear.

Confirm the RDNA3 SIMT reduce path; I'll submit the bitcast fix + a regression test.

Note: this issue was made and written with help of Deepseek V4.1