[Bug] HIP: INT8 tensorwise matmul (GGML_TYPE_I8 weights) falls back to CPU — sd.cpp FP8/INT8 models run ~50-100× slower on ROCm

[akhenakh]

Git commit

6b3edaa

Operating System & Version

MANJARO 7.1

GGML backends

HIP

Command-line arguments used

using the library

Steps to reproduce

this is an issue on the gllm fork leejet/ggml

What you expected to happen

sd.cpp's FP8→INT8 weight conversion produces GGML_TYPE_I8 weights with f32 scale/bias tensors.
On CUDA builds these run on the GPU via ggml_cuda_mul_mat_i8 (cuBLAS CUBLAS_COMPUTE_32I GEMM + custom quantize/dequantize kernels). On HIP builds the entire path is compiled out:

// ggml-cuda.cu, ggml_cuda_mul_mat():
if (src0->type == GGML_TYPE_I8) {
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
    ggml_cuda_mul_mat_i8(ctx, src0, src1, dst);
#else
    GGML_ABORT("INT8 tensorwise matmul is only implemented for CUDA");
#endif

ggml_backend_cuda_supports_op() correspondingly returns false for MUL_MAT with GGML_TYPE_I8 src0 and for GGML_OP_QUANTIZE_I8_CONVROT under GGML_USE_HIP.
As a result, sd.cpp's compute-buffer planner appends a trailing CPU fallback backend, and every INT8 GEMM of the diffusion model executes on the CPU — model weights sit in VRAM, but the GPU idles.

• GPU: AMD Radeon RX 7900 XTX (gfx1100), ROCm 7.2 (hipBLAS/rocBLAS)

• Build: -DSD_HIPBLAS=ON -DSD_BUILD_SHARED_LIBS=ON, amdclang++ --offload-arch=gfx1100

• Model: any FP8/INT8 checkpoint (weight stats: f32: 231 | i8: 224 | bf16: 1106)

• total params memory size = 22323 MB (VRAM 22323 MB, RAM 0 MB) — weights fully in VRAM

• CPU pegged at 100%, GPU at 0% during sampling; a 4-step 1024×1024 generation does not complete in >60s (text-encoder conditioning alone takes ~8.5s)

The HIP exclusions are conservative, not fundamental:

1. The custom kernels (quantize_rowwise_i8_cuda, ggml_cuda_quantize_i8_convrot, dequantize_i32_rows_cuda) are plain CUDA and compile cleanly with hipcc.

2. cublasGemmEx(..., CUDA_R_8I, ..., CUBLAS_COMPUTE_32I, ...) maps to hipblasGemmEx(..., HIPBLAS_R_8I, ..., HIPBLAS_COMPUTE_32I, ...) — rocBLAS provides INT8 GEMM on RDNA3 (WMMA) and CDNA (MFMA).

3. Two small compat gaps in vendors/hip.h: CUDA_R_8I/CUDA_R_32I/CUBLAS_COMPUTE_32I are not defined, and __shfl_down_sync is not remapped (the CUDA variants allow a 3-arg call with implicit width, and HIP 7 requires 64-bit warp masks, so a direct call fails to compile).

What actually happened

llgm wrongly use the cpu compute unit

Logs / error messages / stack trace

No response

Additional context / environment details

AI generated:

Proposed fix (working patch, ~15 lines)

1. In vendors/hip.h, add:
   #define CUDA_R_8I HIPBLAS_R_8I
   #define CUDA_R_32I HIPBLAS_R_32I
   // HIP >= 6.5: #define CUBLAS_COMPUTE_32I HIPBLAS_COMPUTE_32I
   // older: #define CUBLAS_COMPUTE_32I HIPBLAS_R_32I
   #define __shfl_down_sync(mask, var, laneMask, ...) __shfl_down(var, laneMask, ##VA_ARGS)
   (variadic to match the CUDA 3-arg form)
2. In ggml-cuda.cu, change the four #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) guards guarding the INT8 path (int8 kernel/impl section, dispatch in ggml_cuda_mul_mat, supports_op cases for MUL_MAT/QUANTIZE_I8_CONVROT, forward-dispatch of QUANTIZE_I8_CONVROT) to #if !defined(GGML_USE_MUSA).
3. In supports_op for I8 MUL_MAT, replace the NVIDIA-only turing_mma_available(cc) requirement with an arch-aware gate, since HIP devices carry AMD-offset cc values (GGML_CUDA_CC_IS_NVIDIA is false for them):
   const int i8_dev_cc = ggml_cuda_info().devices[dev_ctx->device].cc;
   // ...
   ((GGML_CUDA_CC_IS_NVIDIA(i8_dev_cc) && turing_mma_available(i8_dev_cc)) ||
   GGML_CUDA_CC_IS_RDNA3(i8_dev_cc) || GGML_CUDA_CC_IS_CDNA(i8_dev_cc))
   Keeping an explicit capability gate means older AMD cards (pre-RDNA3, no int8 WMMA/MFMA) still get the safe CPU fallback instead of a runtime hipblasGemmEx failure.