From 3d65ff86514ea82e056b566f294e8b0ff11332f9 Mon Sep 17 00:00:00 2001 From: Tai An Date: Wed, 29 Jul 2026 00:04:12 -0700 Subject: [PATCH 1/3] Fix truncated FP4 constant in Triton dequantize_fp4_tree The magnitude of the smallest nonzero FP4 level (codes 0b0001/0b1001) was written as 0.00520833, which rounds to a different float32 (0x3BAAAAA4) than every other FP4 table in the repository. The CUDA kernel, CPU kernel, 4-bit GEMM LUT, and the get_4bit_type("fp4") Python table all use 1/192 = 0.005208333333 (0x3BAAAAAB). Use the full-precision literal so the Triton backend agrees with the other backends bit-for-bit. Fixes #2025 --- bitsandbytes/backends/triton/kernels_4bit.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/bitsandbytes/backends/triton/kernels_4bit.py b/bitsandbytes/backends/triton/kernels_4bit.py index bdd59fad2..f9026fecc 100644 --- a/bitsandbytes/backends/triton/kernels_4bit.py +++ b/bitsandbytes/backends/triton/kernels_4bit.py @@ -219,7 +219,7 @@ def dequantize_fp4_tree(val, absmax): branch2 = tl.where( second_bit, tl.where(first_bit, 1.0, 0.66666667), # 1011, 1010 - tl.where(first_bit, 0.00520833, 0.0), # 1001, 1000 + tl.where(first_bit, 0.005208333333, 0.0), # 1001, 1000 ) out = tl.where(third_bit, branch1, branch2) return out * sign * absmax @@ -574,4 +574,4 @@ def quantize_4bit_blockwise_kernel( packed_flat = tl.reshape(packed, (BLOCK_SIZE * SPLIT_NUM_BLOCKS,)) out_offsets = block_start_idx * BLOCK_SIZE // 2 + tl.arange(0, SPLIT_NUM_BLOCKS * BLOCK_SIZE) out_mask = out_offsets < n_elements // 2 - tl.store(out_ptr + out_offsets, packed_flat, mask=out_mask) + tl.store(out_ptr + out_offsets, packed_flat, mask=out_mask) \ No newline at end of file From 9647805a5ba429f31ae04c3381310982ca4901d0 Mon Sep 17 00:00:00 2001 From: Tai An Date: Wed, 29 Jul 2026 09:06:41 -0700 Subject: [PATCH 2/3] style: add missing newline at end of kernels_4bit.py (ruff W292) From 1c9d959f25ebcf904e7295062cfd9cd7a64e29f0 Mon Sep 17 00:00:00 2001 From: Tai An Date: Wed, 29 Jul 2026 09:07:07 -0700 Subject: [PATCH 3/3] style: add missing newline at end of kernels_4bit.py (ruff W292) --- bitsandbytes/backends/triton/kernels_4bit.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bitsandbytes/backends/triton/kernels_4bit.py b/bitsandbytes/backends/triton/kernels_4bit.py index f9026fecc..6e8775bc2 100644 --- a/bitsandbytes/backends/triton/kernels_4bit.py +++ b/bitsandbytes/backends/triton/kernels_4bit.py @@ -574,4 +574,4 @@ def quantize_4bit_blockwise_kernel( packed_flat = tl.reshape(packed, (BLOCK_SIZE * SPLIT_NUM_BLOCKS,)) out_offsets = block_start_idx * BLOCK_SIZE // 2 + tl.arange(0, SPLIT_NUM_BLOCKS * BLOCK_SIZE) out_mask = out_offsets < n_elements // 2 - tl.store(out_ptr + out_offsets, packed_flat, mask=out_mask) \ No newline at end of file + tl.store(out_ptr + out_offsets, packed_flat, mask=out_mask)