feat: enable hip INT8 tensorwise matmul and convrot

[leejet]

Summary

Enable HIP INT8 tensorwise matmul on CDNA, RDNA3/3.5, and RDNA4 devices, avoiding CPU fallback for supported configurations.

Related Issue / Discussion

Fix #1929.

Additional Information

N/A

Checklist

• I have read and confirmed this PR follows the contribution guidelines.

[kwilkins-82]

Thanks for this. Note I needed a couple of changes to get it working with Bedovyy/FLUX.2-klein-4B-INT8-Comfy, but it's now running very well on my AMD Radeon 860M.

diff --git a/src/ggml-cuda/ggml-cuda.cu b/src/ggml-cuda/ggml-cuda.cu
index 1012a457..99c619d8 100644
--- a/src/ggml-cuda/ggml-cuda.cu
+++ b/src/ggml-cuda/ggml-cuda.cu
@@ -1609,7 +1609,7 @@ global void convrot_group_quantize_i8_cuda(

global void dequantize_i32_rows_cuda(
const int32_t * src, const float * scales, const float * weight_scales,

•    const float * bias, float * dst, int64_t n, int64_t rows) {
  

•    const float * bias, float * dst, int64_t n, int64_t rows, bool single_weight_scale) {
  

  const int64_t output = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
  const int64_t row = blockIdx.y;
  if (output < n && row < rows) {
  @@ -1617,7 +1617,7 @@ global void dequantize_i32_rows_cuda(
  const float activation_scale = scales[row];
  float value = (float)src[index] * activation_scale;
  if (weight_scales != nullptr) {

•        value *= weight_scales[output];
  

•        value *= weight_scales[single_weight_scale ? 0 : output];
     }
     if (bias != nullptr) {
         value += bias[output];
  

@@ -1751,10 +1751,11 @@ static void ggml_cuda_mul_mat_i8(
&beta, accum.get(), CUDA_R_32I, (int)n,
CUBLAS_COMPUTE_32I, CUBLAS_GEMM_DEFAULT_TENSOR_OP));

• const bool single_weight_scale = dst->src[2] != nullptr && ggml_nelements(dst->src[2]) == 1;
  const dim3 dequant_grid((n + 255) / 256, rows, 1);
  dequantize_i32_rows_cuda<<<dequant_grid, 256, 0, stream>>>(
  accum.get(), scales_d, weight_scales, bias,

•    (float *)dst->data, n, rows);
  

•    (float *)dst->data, n, rows, single_weight_scale);
  

}
#endif

@@ -5852,7 +5853,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
ggml_is_contiguous(a) && ggml_is_contiguous(b) && ggml_is_contiguous(op) &&
(weight_scale == nullptr ||
(weight_scale->type == GGML_TYPE_F32 && ggml_is_contiguous(weight_scale) &&

•                         ggml_nelements(weight_scale) == a->ne[1])) &&
  

•                         (ggml_nelements(weight_scale) == a->ne[1] || ggml_nelements(weight_scale) == 1))) &&
                        (bias == nullptr ||
                         (bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias) &&
                          ggml_nelements(bias) == a->ne[1])) &&
  

diff --git a/src/ggml.c b/src/ggml.c
index 76b10130..dc1fbfe5 100644
--- a/src/ggml.c
+++ b/src/ggml.c
@@ -3388,7 +3388,7 @@ GGML_API struct ggml_tensor * ggml_mul_mat_i8_tensorwise(
GGML_ASSERT(input->type == GGML_TYPE_F32 || input->type == GGML_TYPE_I8);
GGML_ASSERT(weight_scale != NULL && weight_scale->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_contiguous(weight_scale));

• GGML_ASSERT(ggml_nelements(weight_scale) == weight->ne[1]);

• GGML_ASSERT(ggml_nelements(weight_scale) == weight->ne[1] || ggml_nelements(weight_scale) == 1);
  GGML_ASSERT(bias == NULL || (bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias)));
  GGML_ASSERT(bias == NULL || ggml_nelements(bias) == weight->ne[1]);
  GGML_ASSERT(convrot_group_size == 0 || weight->ne[0] % convrot_group_size == 0);