Lumpiasty/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#24546)

Browse Source 
* adjust ncols_picker for routed MoE in mul_mat_q_case function

* Adding CDNA, RDNA2 and RDNA4

* fix: update mmq_use_routed_moe_ncols_picker to include NVIDIA + Volta support

* feat: enhance mmq configuration for various architectures with moe_ncols_min_cc support

* refactor: replace moe_ncols_min_cc with use_typical_moe_ncols in mmq configuration files

* HIP: mmq: enable typical moe ncols on RDNA4

---------

Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
This commit is contained in:
ravel7524
2026-09-07 15:22:42 +02:00
committed by GitHub
co-authored by Carl Philipp Klemm
parent 4735997382
commit 0c963452ea
10 changed files with 36 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/mmq-config-blackwell.cuh
+1
View File
@@ -1,4 +1,5 @@
static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_blackwell(ggml_type type, int J, bool fallback) {
constexpr bool use_typical_moe_ncols = false;
CASE(GGML_TYPE_MXFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);
CASE(GGML_TYPE_MXFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);
CASE(GGML_TYPE_MXFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);