fix: bypass broken hipblasLt int8 GEMM on gfx950#94
Open
kudomcho wants to merge 3 commits into
Open
Conversation
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) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
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) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Root cause
hipBLASLt's column-major Int8 GEMM kernel has a tiling bug on gfx950 — works for ≤64 columns, produces wrong results for ≥128 columns, and writes out of bounds on large inputs causing unrecoverable GPU page faults. The col-major int32 matrix transform is also broken.
Fix
Bypass the broken hipBLASLt col-major path on gfx950 by routing Int8 GEMM through
torch._int_mm(row-major), which works correctly at all sizes. Python-only patch, no rebuild needed.Changes:
bitsandbytes/backends/cuda.py: Add_igemmlt_fallback()gated onROCM_GPU_ARCH == "gfx950". Recovers row-major matrices from col-major transformed inputs, computes viatorch._int_mmwith K/N padding to multiples of 8, returns result in row formatbitsandbytes/functional.py: Add no-op short-circuit innvidia_transform()whenfrom_order == to_order(needed because fallback returns row-format output)Test plan
pytest -xvs tests/test_functional.py -k "igemmlt and cuda"— 2/2 passed (dims=2 and dims=3)pytest -xvs tests/test_functional.py -k "igemmlt"— 3 passed, 12 skipped, 0 failed