[ROCm] Preserve float bits in SIMT DPP reduction

[liminfei-amd]

Summary

• Preserve the floating-point accumulator bits around AMD's
  __builtin_amdgcn_mov_dpp in the HIP SIMT warp reduction.
• Correct the DPP comments: control values 0x101-0x10f are row_shl, not
  row_shr.
• Add a regression test for the fused Linear4bit path. It runs on every ROCm
  architecture where _gemm_4bit_use_custom_rocm selects the fused SIMT
  kernel, for M = 1, 2, and 8 and fp16, bf16, and fp32, and compares against an
  fp32 matmul with the same dequantized weight.

The second commit is from @LaelaZorana, with her authorship preserved. It adds
the comment correction and the broader test.

Fixes #2087.

Root cause

The HIP warp reduction passed a float directly to
__builtin_amdgcn_mov_dpp. In clang 19, which ROCm 6.4.x uses, the builtin
takes an int, so each partial sum was converted to an integer and back
(v_cvt_i32_f32 + v_cvt_f32_i32_dpp) instead of moving its IEEE-754 bits.
LLVM 20 made the builtin overloaded and removed the implicit conversion
(llvm/llvm-project#112447), so newer compilers generate correct code for the
old source.

Explicit __float_as_uint / __uint_as_float conversions move the bits with
every compiler. ROCm 6.4.4 is still in the bitsandbytes build matrix.

Validation

The nine new test cases (fp16, bf16, fp32 x M = 1, 2, 8):

GPU	Toolchain	Without the fix	With the fix
gfx1100, 2x RX 7900 GRE (cuda:0 and cuda:1)	ROCm 6.4.4 / clang 19, torch 2.7.1	9 failed (relative L1 0.124-0.130)	9 passed (relative L1 <= 0.0027)
gfx1100	ROCm 10.1, torch 2.12	9 passed	9 passed
gfx1151, Radeon 8060S	ROCm 10.1, torch 2.12	9 passed	9 passed

• With ROCm 10.1 the old source also passes, as expected from the compiler
  change above, so the test catches this bug only in clang 19 builds.
• The ROCm 6.4.4 PyTorch image does not run on gfx1151, so for that target the
  clang 19 code objects were checked instead: without the fix, the 24 SIMT
  kernels contain 96 v_cvt_f32_i32_dpp instructions; with the fix they
  contain none and the reduction uses v_add_f32_dpp.
• Without a GPU, all nine cases are skipped.
• pre-commit run --all-files passed.

RDNA4 and CDNA hardware were not tested. The model workload from the original
report was not rerun; the test uses the minimal tensor reproducer.

AI assistance was used for reproduction, diagnosis, validation, and drafting
this description. I reviewed the changes and the validation results.

[LaelaZorana]

Hi liminfei-amd, I also confirmed that it stands on the codegen side for all targets under ROCm 6.4.4.

Compiling just those 4 DPP steps with clang for the amdgcn target generates a pair of cvt instructions in the unmodified version. With the fix it generates 1 instruction per DPP. Details below. It appears the issue is clang version specific.

With upstream clang 19.1.7 it will generate v_cvt_i32_f32 and v_cvt_f32_i32_dpp instructions on the unmodified code. So it converts the partial sums into integers before the shift, thus losing precision. With the modification it will generate a single instruction per DPP, v_add_f32_dpp, which is what we want.

With later versions of upstream clang it does not generate integer conversion on the unmodified code. Specifically,

clang 20.1.8, clang 21.1.8, clang 22.1.8, no int conversion on the old code.

clang 23.1.2 generates v_add_f32_dpp with or without the bitcast.

So this seems to be a special issue with clang19, the version used by ROCm 6.4.4. Also it happens on all 9 targets I've compiled, which are from the ROCm 6.4.4 list, gfx1100, gfx1101, gfx1102, gfx1150, gfx1151, gfx1200, gfx1201, gfx90a, gfx942.

Modeling lane split and reduction on CPU using the repro shape, M 8, N 256 and K 512, confirms converting the moved value will result in 0.1245 relative L1 error against the dequantized weight on bf16 and 0.1246 on fp16 and fp32, which is very close to 0.125295 as you've seen on GRE. Keeping the bits will result in 0.0 error.

Suggestions

I noticed _gemm_4bit_use_custom_rocm allows the use of this kernel on gfx12 for M up to 8, as well as gfx9 for M up to 4 on fp16 and bf16 or M up to 2 on fp32. So perhaps the newly added test should be enabled for all archs that invoke this kernel? For example simply check on _gemm_4bit_use_custom_rocm, and also include M 1 and 2 in the list. That way it will be tested by RDNA4 and MI300 runners.

Some comments are wrong. For example all those constants, 0x101 to 0x10f, are row_shl, row_shl:8 in the printout, but the comments call them row_shr. The instructions are correct for lane 0 though.

I created a branch that makes the same modification, fixes comments and extends test. Take any of it into this PR.

[LaelaZorana]

Branch is up at https://github.com/LaelaZorana/bitsandbytes/tree/fix/rocm-dpp-float-bits-2087 and the diff against main is at main...LaelaZorana:bitsandbytes:fix/rocm-dpp-float-bits-2087

[LaelaZorana]

To be exact about the compiler, the matrix above used upstream clang 19.1.7, the same major version ROCm 6.4.4 ships, not the ROCm toolchain itself. The 0.0 is the CPU model checked against its own dequantized reference, and the 0.002051 you measured on the GRE is the real hardware number.

[liminfei-amd]

Thanks @LaelaZorana, this was very helpful. Your commit is now part of this PR with your authorship preserved (eb5f933). It brings in the row_shl comment fix and your broader test (every ROCm arch that selects the kernel, M = 1, 2, 8).

I ran the new test on RX 7900 GRE (gfx1100):

• ROCm 6.4.4 / clang 19: without the fix all 9 cases fail (relative L1 0.124-0.130); with the fix all 9 pass (<= 0.0027).
• ROCm 10.1: all 9 pass. The old code passes there too, which matches your clang 20+ results.