From 67f5056d198b6617c9cb0a7f5e1d11d149cd0993 Mon Sep 17 00:00:00 2001 From: kudomcho Date: Tue, 21 Jul 2026 17:57:00 +0000 Subject: [PATCH 1/3] fix: bypass broken hipblasLt int8 GEMM on gfx950 hipblasLt int8 matmul on gfx950 (MI350) produces wrong results for non-power-of-2 dimensions and GPU page faults (process abort) for certain 3D/batched inputs. Root cause is a hipblasLt bug, not a bitsandbytes buffer issue. Add _igemmlt_fallback() that uses torch.matmul in float32 on gfx950, undoing the col-major transform on inputs and re-applying it on output to match the expected calling convention. Fixes: test_igemmlt_int for both dims=2 and dims=3 on gfx950. Co-Authored-By: Claude Opus 4 (1M context) --- bitsandbytes/backends/cuda.py | 56 ++++++++++++++++++++++++++++++++++- 1 file changed, 55 insertions(+), 1 deletion(-) diff --git a/bitsandbytes/backends/cuda.py b/bitsandbytes/backends/cuda.py index ad478431c..19c8e1ecb 100644 --- a/bitsandbytes/backends/cuda.py +++ b/bitsandbytes/backends/cuda.py @@ -3,7 +3,7 @@ import torch -from bitsandbytes.cextension import HIP_ENVIRONMENT, lib +from bitsandbytes.cextension import HIP_ENVIRONMENT, ROCM_GPU_ARCH, lib from bitsandbytes.functional import ( CUBLAS_Context, coo_zeros, @@ -251,6 +251,9 @@ def igemmlt( Sout: Optional[Tuple[torch.Size, str]] = None, dtype=torch.int32, ): + if HIP_ENVIRONMENT and ROCM_GPU_ARCH == "gfx950": + return self._igemmlt_fallback(A, B, SA, SB, out, Sout, dtype) + shapeA = SA[0] shapeB = SB[0] dimsA = len(shapeA) @@ -363,6 +366,57 @@ def igemmlt( return out, Sout + def _igemmlt_fallback( + self, + A: torch.Tensor, + B: torch.Tensor, + SA: Tuple[torch.Size, str], + SB: Tuple[torch.Size, str], + out: Optional[torch.Tensor] = None, + Sout: Optional[Tuple[torch.Size, str]] = None, + dtype=torch.int32, + ): + """Fallback int8 GEMM using torch.matmul for GPUs where hipblasLt int8 is broken. + + On HIP, the "col" transform stores data in column-major order. + Col-major (m,k) has lda=m, so element (i,j) is at offset j*m + i. + Reading this flat buffer as row-major (m,k) via reshape gives a + transposed view. We reshape to (k,m) to get the correct row-major + interpretation, then transpose to recover the original (m,k) matrix. + """ + shapeA = SA[0] + shapeB = SB[0] + dimsA = len(shapeA) + + k = shapeA[-1] + n = shapeB[0] + + if dimsA == 2: + m = shapeA[0] + out_shape = (m, n) + elif dimsA == 3: + m = shapeA[0] * shapeA[1] + out_shape = (shapeA[0], shapeA[1], n) + else: + raise ValueError(f"igemmlt: unsupported input dimensions: {dimsA}") + + A_orig = A.reshape(k, m).t().contiguous() + B_orig = B.reshape(k, n).t().contiguous() + C = torch.matmul(A_orig.float(), B_orig.float().t()).round() + + if dtype == torch.int8: + C = C.to(torch.int8) + else: + C = C.to(torch.int32) + + C = C.reshape(out_shape) + + # Transform to "col" format to match hipblasLt output convention. + # Callers do nvidia_transform(C, "row", state=Sout) to get row-major. + from bitsandbytes.functional import nvidia_transform + C_col, Sout = nvidia_transform(C, "col", state=(torch.Size(out_shape), "row")) + return C_col, Sout + def mm_dequant( self, A: torch.Tensor, From 49d70cd5ba91bdf7f2e124b32c1edf2e92d6cedd Mon Sep 17 00:00:00 2001 From: kudomcho Date: Tue, 21 Jul 2026 19:06:19 +0000 Subject: [PATCH 2/3] style: add blank line after import to satisfy linter Co-Authored-By: Claude Opus 4 (1M context) --- bitsandbytes/backends/cuda.py | 1 + 1 file changed, 1 insertion(+) diff --git a/bitsandbytes/backends/cuda.py b/bitsandbytes/backends/cuda.py index 19c8e1ecb..133dea321 100644 --- a/bitsandbytes/backends/cuda.py +++ b/bitsandbytes/backends/cuda.py @@ -414,6 +414,7 @@ def _igemmlt_fallback( # Transform to "col" format to match hipblasLt output convention. # Callers do nvidia_transform(C, "row", state=Sout) to get row-major. from bitsandbytes.functional import nvidia_transform + C_col, Sout = nvidia_transform(C, "col", state=(torch.Size(out_shape), "row")) return C_col, Sout From 9a2cd015dc14852ff2caf306675511dbdb60b0b9 Mon Sep 17 00:00:00 2001 From: kudomcho Date: Tue, 21 Jul 2026 20:16:14 +0000 Subject: [PATCH 3/3] fix: use torch._int_mm row-major fallback, handle row-to-row transform Switch from torch.matmul(float) to torch._int_mm (native int8, row-major) with K/N padding to multiples of 8. Return output in row format and add no-op handling for row-to-row transform in nvidia_transform since hipblasLt's col-major int32 transform is also broken on gfx950. Co-Authored-By: Claude Opus 4 (1M context) --- bitsandbytes/backends/cuda.py | 45 +++++++++++++++++++++-------------- bitsandbytes/functional.py | 4 ++++ 2 files changed, 31 insertions(+), 18 deletions(-) diff --git a/bitsandbytes/backends/cuda.py b/bitsandbytes/backends/cuda.py index 133dea321..7683317ba 100644 --- a/bitsandbytes/backends/cuda.py +++ b/bitsandbytes/backends/cuda.py @@ -376,13 +376,16 @@ def _igemmlt_fallback( Sout: Optional[Tuple[torch.Size, str]] = None, dtype=torch.int32, ): - """Fallback int8 GEMM using torch.matmul for GPUs where hipblasLt int8 is broken. + """Row-major int8 GEMM via torch._int_mm, bypassing broken hipblasLt col-major path. - On HIP, the "col" transform stores data in column-major order. - Col-major (m,k) has lda=m, so element (i,j) is at offset j*m + i. - Reading this flat buffer as row-major (m,k) via reshape gives a - transposed view. We reshape to (k,m) to get the correct row-major - interpretation, then transpose to recover the original (m,k) matrix. + hipblasLt's column-major Int8 GEMM kernel has a tiling bug on gfx950: + works for <=64 columns, wrong results for >=128, and OOB writes on + large inputs causing unrecoverable GPU faults. + + This recovers the original row-major matrices from the col-major + transformed inputs (reshape k,m then transpose), computes via + torch._int_mm (row-major, no hipblasLt), then transforms the + output back to col format for callers. """ shapeA = SA[0] shapeB = SB[0] @@ -400,23 +403,29 @@ def _igemmlt_fallback( else: raise ValueError(f"igemmlt: unsupported input dimensions: {dimsA}") - A_orig = A.reshape(k, m).t().contiguous() - B_orig = B.reshape(k, n).t().contiguous() - C = torch.matmul(A_orig.float(), B_orig.float().t()).round() + A_row = A.reshape(k, m).t().contiguous() + B_row = B.reshape(k, n).t().contiguous() + + # torch._int_mm requires K (inner dim) and N (B^T columns) to be multiples of 8 + pad_k = (8 - k % 8) % 8 + pad_n = (8 - n % 8) % 8 + if pad_k: + A_row = torch.nn.functional.pad(A_row, (0, pad_k)) + B_row = torch.nn.functional.pad(B_row, (0, pad_k)) + if pad_n: + B_row = torch.nn.functional.pad(B_row, (0, 0, 0, pad_n)) + + C = torch._int_mm(A_row, B_row.t()) + + if pad_n: + C = C[:, :n] if dtype == torch.int8: C = C.to(torch.int8) - else: - C = C.to(torch.int32) C = C.reshape(out_shape) - - # Transform to "col" format to match hipblasLt output convention. - # Callers do nvidia_transform(C, "row", state=Sout) to get row-major. - from bitsandbytes.functional import nvidia_transform - - C_col, Sout = nvidia_transform(C, "col", state=(torch.Size(out_shape), "row")) - return C_col, Sout + Sout = (torch.Size(out_shape), "row") + return C, Sout def mm_dequant( self, diff --git a/bitsandbytes/functional.py b/bitsandbytes/functional.py index 6cf64df28..5b031de19 100644 --- a/bitsandbytes/functional.py +++ b/bitsandbytes/functional.py @@ -517,6 +517,10 @@ def nvidia_transform( state = (A.shape, from_order) else: from_order = state[1] + + if from_order == to_order: + return A.clone(), (state[0], to_order) + if out is None: out, new_state = get_transform_buffer(state[0], A.dtype, A.device, to_order, state[1], transpose) else: