Skip to content

HIP: RDNA 4 MUL_MAT optimizations #25940

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
itterative wants to merge 3 commits into
base: master
Choose a base branch
from
+27 −4
Open

HIP: RDNA 4 MUL_MAT optimizations #25940

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
0a71b4f
fix: q6_k mma
itterative Jul 20, 2026
140878b
fix: q2_k mma
itterative Jul 20, 2026
2ba97e5
perf: sweep quants for best performance
itterative Jul 20, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions ggml/src/ggml-cuda/mmq-vec-dot.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -693,6 +693,7 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_

const int i0 = (threadIdx.y / ntx) * rows_per_warp;

#pragma unroll 1
for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) {
const int k0 = k00 + k01;

Expand Down Expand Up @@ -1055,7 +1056,7 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
for (int l = 0; l < tile_C::ne; ++l) {
const int i = i0 + n*tile_C::I + tile_C::get_i(l);
const int8_t * sc = (const int8_t *) (x_sc + i*sram_stride + k00/16);
sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l] * sc[k01/4] * x_df[i*sram_stride] * dB;
sum[(j0/tile_C::J + n)*tile_C::ne + l] += ((float) C.x[l]) * sc[k01/4] * x_df[i*sram_stride] * dB;
}
}
}
Expand Down Expand Up @@ -1237,4 +1238,3 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
}
}
}

27 changes: 25 additions & 2 deletions ggml/src/ggml-cuda/mmq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -377,8 +377,31 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
}
}

// For RDNA4 MMQ is consistently faster than dequantization + hipBLAS:
// https://github.com/ggml-org/llama.cpp/pull/18537#issuecomment-3706422301
if (GGML_CUDA_CC_IS_RDNA4(cc)) {
// For RDNA4 MMQ is consistently faster than dequantization + hipBLAS:
// https://github.com/ggml-org/llama.cpp/pull/18537#issuecomment-3706422301
if (n_experts > 0) {
return true;
}

switch (type) {
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
return ne11 <= 512;
case GGML_TYPE_Q5_0:
return ne11 <= 336;
case GGML_TYPE_Q6_K:
case GGML_TYPE_Q5_1:
return ne11 <= 256;
case GGML_TYPE_Q2_K:
return ne11 <= 96;
default:
break;
}
}

return true;
}

Expand Down