CUDA/HIP: INT8 convrot support on RDNA2 - #9
Open
noctrex wants to merge 4 commits into
Open
Conversation
Remove the blanket GGML_USE_HIP exclusion from the INT8 tensorwise convrot ops and scope the capability check to GPUs with working integer GEMM support (turing_mma_available, amd_wmma_available, amd_mfma_available). RDNA2 stays excluded: hipBLASLt INT8 is broken there. vendors/hip.h: map CUDA_R_8I, CUDA_R_32I, CUBLAS_COMPUTE_32I and __shfl_down_sync to their HIP equivalents. HIP's masked shuffle templates require a 64-bit mask, so route to the maskless intrinsic like the other __shfl_*_sync mappings. Verified on gfx1100 (Windows ROCm 7.1): ggml tests/test-int8-convrot passes, and CPU-vs-GPU output is bit-exact at real model shapes (K=3840/10240, N up to 11520, 8104 rows).
Template the INT8 convrot quantization kernels on group_size and instantiate 64 alongside 256. The Hadamard normalization is now derived from the group size (2^-log4(group_size): 1/8 for G=64, 1/16 for G=256). The small-kernel max reduction uses a shared-memory tree instead of warp shuffles so it stays correct for 16-thread blocks. Capability gates accept group size 64 and 256. Z-Image int8_tensorwise exports in the wild use convrot_groupsize 64; without this they fall back to CPU on every GPU backend. Verified on gfx1100 (Windows ROCm 7.1): exact-reference tests pass for H64 (single-group and multi-group packed GEMM), H256 unchanged, and CPU-vs-GPU output is bit-exact at real Z-Image shapes.
hipblasGemmEx computes CUDA_R_8I x CUDA_R_8I -> CUDA_R_32I correctly on gfx1030 with ROCm >= 7.14 (verified first-hand on a V620, Ubuntu Server 24.04.4) and is ~5.7x faster than the DP4A MMQ kernel at Z-Image shapes. Extend ggml_cuda_i8_blas_available so both the supports_op gate and the dispatch select BLAS on RDNA2; GGML_CUDA_FORCE_I8_MMQ still forces the kernel for stacks with a broken rocBLAS INT8 route.
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.
Enables the INT8 convrot tensorwise mulmat on RDNA2 (gfx1030-class), previously rejected by supports_op.
Stacked on top of the open HIP enablement / H64 PRs (b688cbe, 0ec2d6a).
Once those merge, this PR reduces to the two RDNA2 commits.
Verification (V620 gfx1030, ROCm 7.14.0~pre3, ROCk 6.19.14, Ubuntu 24.04.4)
Correctness:
Performance (full quantize+GEMM+epilogue, 5 iters):
Known issue
If the op is not claimed by a GPU backend, the scheduler falls it back to CPU, where the padded (rows+1) GPU-quantized src1 trips GGML_ASSERT(ne1 == ne11) in ggml-cpu.c.
Unreachable with this gate in place, but a cross-backend fallback should tolerate the padding.
LLM Disclose: YES, used GLM-5.3 for the grunt work