From d590d224254a48a6835a54b24ee288994d1e49f8 Mon Sep 17 00:00:00 2001 From: ravel7524 <58877666+ravel7524@users.noreply.github.com> Date: Fri, 12 Jun 2026 22:05:29 +0200 Subject: [PATCH 1/5] adjust ncols_picker for routed MoE in mul_mat_q_case function --- ggml/src/ggml-cuda/mmq.cuh | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 2eb15fdfad93..9d00e4f69f1f 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1472,6 +1472,18 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, const int cc = ggml_cuda_info().devices[id].cc; const size_t smpbo = ggml_cuda_info().devices[id].smpbo; + int64_t ncols_picker = args.ncols_max; + if (args.expert_bounds != nullptr && GGML_CUDA_CC_IS_RDNA3(cc) && args.nchannels_x > 0) { + // In routed MoE, ncols_max is the worst-case per-expert width. Size the + // MMQ N-tile from the typical routed width while it is below the RDNA3 + // max tile width. The launch grid still uses args.ncols_max. + const int J_max = ggml_cuda_mmq_get_J_max(type, fallback, cc, 128); + const int64_t ncols_typical = (args.ncols_dst + args.nchannels_x - 1) / args.nchannels_x; + if (ncols_typical >= 1 && ncols_typical < J_max && ncols_typical < ncols_picker) { + ncols_picker = ncols_typical; + } + } + int J_best = 0; int ntiles_J_best = INT_MAX; @@ -1485,7 +1497,7 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, continue; } - const int ntiles_x = (args.ncols_max + config.J - 1) / config.J; + const int ntiles_x = (ncols_picker + config.J - 1) / config.J; if (ntiles_x < ntiles_J_best) { J_best = J; From 3087cea54754976e7f417df42158a5cf9ae4b381 Mon Sep 17 00:00:00 2001 From: ravel7524 <58877666+ravel7524@users.noreply.github.com> Date: Fri, 19 Jun 2026 14:43:12 +0200 Subject: [PATCH 2/5] Adding CDNA, RDNA2 and RDNA4 --- ggml/src/ggml-cuda/mmq.cuh | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 9d00e4f69f1f..025be70a17bd 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1466,6 +1466,13 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a ntx_fd); } +static bool mmq_use_routed_moe_ncols_picker(const int cc) { + return GGML_CUDA_CC_IS_CDNA(cc) || + GGML_CUDA_CC_IS_RDNA2(cc) || + GGML_CUDA_CC_IS_RDNA3(cc) || + GGML_CUDA_CC_IS_RDNA4(cc); +} + template void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { const int id = ggml_cuda_get_device(); @@ -1473,10 +1480,10 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, const size_t smpbo = ggml_cuda_info().devices[id].smpbo; int64_t ncols_picker = args.ncols_max; - if (args.expert_bounds != nullptr && GGML_CUDA_CC_IS_RDNA3(cc) && args.nchannels_x > 0) { + if (args.expert_bounds != nullptr && mmq_use_routed_moe_ncols_picker(cc) && args.nchannels_x > 0) { // In routed MoE, ncols_max is the worst-case per-expert width. Size the - // MMQ N-tile from the typical routed width while it is below the RDNA3 - // max tile width. The launch grid still uses args.ncols_max. + // MMQ N-tile from the typical routed width while it is below the current + // architecture's max tile width. The launch grid still uses args.ncols_max. const int J_max = ggml_cuda_mmq_get_J_max(type, fallback, cc, 128); const int64_t ncols_typical = (args.ncols_dst + args.nchannels_x - 1) / args.nchannels_x; if (ncols_typical >= 1 && ncols_typical < J_max && ncols_typical < ncols_picker) { From a57b8eac36e3eb01bb302bed24ed7f52c2e3065a Mon Sep 17 00:00:00 2001 From: ravel7524 <58877666+ravel7524@users.noreply.github.com> Date: Wed, 5 Aug 2026 15:00:23 -0400 Subject: [PATCH 3/5] fix: update mmq_use_routed_moe_ncols_picker to include NVIDIA + Volta support --- ggml/src/ggml-cuda/mmq.cuh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 025be70a17bd..665a3934acf6 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1467,7 +1467,8 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a } static bool mmq_use_routed_moe_ncols_picker(const int cc) { - return GGML_CUDA_CC_IS_CDNA(cc) || + return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_VOLTA) || + GGML_CUDA_CC_IS_CDNA(cc) || GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc); From f3a7a80ccfcfbb54df3e98c6af60aba872a5ba80 Mon Sep 17 00:00:00 2001 From: ravel7524 <58877666+ravel7524@users.noreply.github.com> Date: Wed, 5 Aug 2026 15:22:44 -0400 Subject: [PATCH 4/5] feat: enhance mmq configuration for various architectures with moe_ncols_min_cc support --- ggml/src/ggml-cuda/mmq-config-ampere.cuh | 3 +- ggml/src/ggml-cuda/mmq-config-blackwell.cuh | 1 + ggml/src/ggml-cuda/mmq-config-cdna.cuh | 3 +- ggml/src/ggml-cuda/mmq-config-pascal.cuh | 3 +- ggml/src/ggml-cuda/mmq-config-rdna2.cuh | 3 +- ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh | 3 +- ggml/src/ggml-cuda/mmq-config-rdna3.cuh | 3 +- ggml/src/ggml-cuda/mmq-config-rdna4.cuh | 3 +- ggml/src/ggml-cuda/mmq.cuh | 37 ++++++++++----------- 9 files changed, 33 insertions(+), 26 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-config-ampere.cuh b/ggml/src/ggml-cuda/mmq-config-ampere.cuh index 9f9fd197382f..7731d2242e9b 100644 --- a/ggml/src/ggml-cuda/mmq-config-ampere.cuh +++ b/ggml/src/ggml-cuda/mmq-config-ampere.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_ampere(ggml_type type, int J, bool fallback) { + constexpr int moe_ncols_min_cc = GGML_CUDA_CC_VOLTA; CASE(GGML_TYPE_Q1_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); @@ -379,5 +380,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh index 9fbe32b6972b..5e9dd340b326 100644 --- a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh +++ b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh @@ -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 int moe_ncols_min_cc = GGML_CUDA_CC_BLACKWELL; 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); diff --git a/ggml/src/ggml-cuda/mmq-config-cdna.cuh b/ggml/src/ggml-cuda/mmq-config-cdna.cuh index 4a8d89f72019..4330f54bdfcc 100644 --- a/ggml/src/ggml-cuda/mmq-config-cdna.cuh +++ b/ggml/src/ggml-cuda/mmq-config-cdna.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_cdna(ggml_type type, int J, bool fallback) { + constexpr int moe_ncols_min_cc = GGML_CUDA_CC_CDNA1; CASE(GGML_TYPE_Q1_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); @@ -181,5 +182,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); CASE(GGML_TYPE_NVFP4, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-pascal.cuh b/ggml/src/ggml-cuda/mmq-config-pascal.cuh index e7d4a9a3fcb5..b9976dc241ea 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_pascal(ggml_type type, int J, bool fallback) { + constexpr int moe_ncols_min_cc = 0; CASE(GGML_TYPE_Q1_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -269,5 +270,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh index 8324d9e1a830..59446fbf68bd 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna2(ggml_type type, int J, bool fallback) { + constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA2; CASE(GGML_TYPE_Q1_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -269,5 +270,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh index 180b2d9370d1..df60462b5c7e 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna3_5(ggml_type type, int J, bool fallback) { + constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA3_5; CASE(GGML_TYPE_Q1_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -286,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh index 676f27fea4d9..11e46d8b15f6 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna3(ggml_type type, int J, bool fallback) { + constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA3; CASE(GGML_TYPE_Q1_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -286,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh index 9293d9d55885..91d17beb21c0 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna4(ggml_type type, int J, bool fallback) { + constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA4; CASE(GGML_TYPE_Q1_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -286,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 665a3934acf6..d603264de0c6 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -170,12 +170,17 @@ struct ggml_cuda_mmq_config { int J; // SRAM tile width in src1->ne[1]/dst->ne[1] direction. ggml_cuda_mmq_sram_layout sram_layout; // SRAM tile length in src0->ne[0]/src1->ne[0] direction (physical 32 bit elements). int K_vram; // VRAM tile length in src0->ne[0]/src1->ne[0] direction (logical elements). + int moe_ncols_min_cc; // Minimum architecture to use the typical routed expert width. bool stream_k; // Whether or not to use stream-k decomposition. bool fallback; // Whether a fallback for out-of-bounds check in src0->ne[1] direction is needed. constexpr __host__ __device__ ggml_cuda_mmq_config( - ggml_type type, int nthreads, int occupancy, int I, int J, ggml_cuda_mmq_sram_layout sram_layout, int K_vram, bool stream_k, bool fallback) : - type(type), nthreads(nthreads), occupancy(occupancy), I(I), J(J), sram_layout(sram_layout), K_vram(K_vram), stream_k(stream_k), fallback(fallback) {} + ggml_type type, int nthreads, int occupancy, int I, int J, ggml_cuda_mmq_sram_layout sram_layout, int K_vram, int moe_ncols_min_cc, bool stream_k, bool fallback) : + type(type), nthreads(nthreads), occupancy(occupancy), I(I), J(J), sram_layout(sram_layout), K_vram(K_vram), moe_ncols_min_cc(moe_ncols_min_cc), stream_k(stream_k), fallback(fallback) {} + + constexpr __host__ __device__ bool use_moe_ncols(const int cc) const { + return moe_ncols_min_cc != 0 && cc >= moe_ncols_min_cc; + } constexpr __device__ int rows_per_warp() const { #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) @@ -210,7 +215,7 @@ struct ggml_cuda_mmq_config { static_assert((I_) % 32 == 0, "bad I"); \ static_assert((J_) % 8 == 0, "bad J"); \ static_assert((K_vram_) % 256 == 0, "bad K_vram"); \ - return ggml_cuda_mmq_config((type_), (nthreads_), (occupancy_), (I_), (J_), (sram_layout_), (K_vram_), (stream_k_), (fallback_)); \ + return ggml_cuda_mmq_config((type_), (nthreads_), (occupancy_), (I_), (J_), (sram_layout_), (K_vram_), moe_ncols_min_cc, (stream_k_), (fallback_)); \ } \ #include "mmq-config-pascal.cuh" @@ -1466,14 +1471,6 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a ntx_fd); } -static bool mmq_use_routed_moe_ncols_picker(const int cc) { - return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_VOLTA) || - GGML_CUDA_CC_IS_CDNA(cc) || - GGML_CUDA_CC_IS_RDNA2(cc) || - GGML_CUDA_CC_IS_RDNA3(cc) || - GGML_CUDA_CC_IS_RDNA4(cc); -} - template void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { const int id = ggml_cuda_get_device(); @@ -1481,14 +1478,16 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, const size_t smpbo = ggml_cuda_info().devices[id].smpbo; int64_t ncols_picker = args.ncols_max; - if (args.expert_bounds != nullptr && mmq_use_routed_moe_ncols_picker(cc) && args.nchannels_x > 0) { - // In routed MoE, ncols_max is the worst-case per-expert width. Size the - // MMQ N-tile from the typical routed width while it is below the current - // architecture's max tile width. The launch grid still uses args.ncols_max. - const int J_max = ggml_cuda_mmq_get_J_max(type, fallback, cc, 128); - const int64_t ncols_typical = (args.ncols_dst + args.nchannels_x - 1) / args.nchannels_x; - if (ncols_typical >= 1 && ncols_typical < J_max && ncols_typical < ncols_picker) { - ncols_picker = ncols_typical; + if (args.expert_bounds != nullptr && args.nchannels_x > 0) { + const int J_max = ggml_cuda_mmq_get_J_max(type, fallback, cc, 128); + const ggml_cuda_mmq_config config_max = ggml_cuda_mmq_get_config(type, J_max, fallback, cc); + if (config_max.use_moe_ncols(cc)) { + // Use the typical expert width only for tile selection. + // The launch grid still uses args.ncols_max. + const int64_t ncols_typical = (args.ncols_dst + args.nchannels_x - 1) / args.nchannels_x; + if (ncols_typical >= 1 && ncols_typical < J_max && ncols_typical < ncols_picker) { + ncols_picker = ncols_typical; + } } } From 2adc35bd58ba498f4a19a74090d8ea403546494f Mon Sep 17 00:00:00 2001 From: ravel7524 <58877666+ravel7524@users.noreply.github.com> Date: Thu, 6 Aug 2026 16:51:57 -0400 Subject: [PATCH 5/5] refactor: replace moe_ncols_min_cc with use_typical_moe_ncols in mmq configuration files --- ggml/src/ggml-cuda/mmq-config-ampere.cuh | 4 ++-- ggml/src/ggml-cuda/mmq-config-blackwell.cuh | 2 +- ggml/src/ggml-cuda/mmq-config-cdna.cuh | 4 ++-- ggml/src/ggml-cuda/mmq-config-pascal.cuh | 4 ++-- ggml/src/ggml-cuda/mmq-config-rdna2.cuh | 4 ++-- ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh | 4 ++-- ggml/src/ggml-cuda/mmq-config-rdna3.cuh | 4 ++-- ggml/src/ggml-cuda/mmq-config-rdna4.cuh | 4 ++-- ggml/src/ggml-cuda/mmq.cuh | 14 +++++--------- 9 files changed, 20 insertions(+), 24 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-config-ampere.cuh b/ggml/src/ggml-cuda/mmq-config-ampere.cuh index 7731d2242e9b..2c00aef2ce14 100644 --- a/ggml/src/ggml-cuda/mmq-config-ampere.cuh +++ b/ggml/src/ggml-cuda/mmq-config-ampere.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_ampere(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = GGML_CUDA_CC_VOLTA; + constexpr bool use_typical_moe_ncols = false; CASE(GGML_TYPE_Q1_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); @@ -380,5 +380,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh index 5e9dd340b326..8f928e217f2c 100644 --- a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh +++ b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_blackwell(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = GGML_CUDA_CC_BLACKWELL; + 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); diff --git a/ggml/src/ggml-cuda/mmq-config-cdna.cuh b/ggml/src/ggml-cuda/mmq-config-cdna.cuh index 4330f54bdfcc..1d51a773b9f2 100644 --- a/ggml/src/ggml-cuda/mmq-config-cdna.cuh +++ b/ggml/src/ggml-cuda/mmq-config-cdna.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_cdna(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = GGML_CUDA_CC_CDNA1; + constexpr bool use_typical_moe_ncols = false; CASE(GGML_TYPE_Q1_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q1_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); @@ -182,5 +182,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); CASE(GGML_TYPE_NVFP4, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-pascal.cuh b/ggml/src/ggml-cuda/mmq-config-pascal.cuh index b9976dc241ea..1ecbc7913bf3 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_pascal(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = 0; + constexpr bool use_typical_moe_ncols = false; CASE(GGML_TYPE_Q1_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -270,5 +270,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh index 59446fbf68bd..c1efef56fb37 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna2(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA2; + constexpr bool use_typical_moe_ncols = false; CASE(GGML_TYPE_Q1_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -270,5 +270,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh index df60462b5c7e..10fdad663ac7 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna3_5(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA3_5; + constexpr bool use_typical_moe_ncols = false; CASE(GGML_TYPE_Q1_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -287,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh index 11e46d8b15f6..e6917fba3574 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna3(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA3; + constexpr bool use_typical_moe_ncols = true; CASE(GGML_TYPE_Q1_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -287,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh index 91d17beb21c0..83e06d9dec3c 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh @@ -1,5 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna4(ggml_type type, int J, bool fallback) { - constexpr int moe_ncols_min_cc = GGML_CUDA_CC_RDNA4; + constexpr bool use_typical_moe_ncols = false; CASE(GGML_TYPE_Q1_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q1_0, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -287,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false); - return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, moe_ncols_min_cc, false, true); + return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index d603264de0c6..90b6484bc852 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -170,17 +170,13 @@ struct ggml_cuda_mmq_config { int J; // SRAM tile width in src1->ne[1]/dst->ne[1] direction. ggml_cuda_mmq_sram_layout sram_layout; // SRAM tile length in src0->ne[0]/src1->ne[0] direction (physical 32 bit elements). int K_vram; // VRAM tile length in src0->ne[0]/src1->ne[0] direction (logical elements). - int moe_ncols_min_cc; // Minimum architecture to use the typical routed expert width. + bool use_typical_moe_ncols; bool stream_k; // Whether or not to use stream-k decomposition. bool fallback; // Whether a fallback for out-of-bounds check in src0->ne[1] direction is needed. constexpr __host__ __device__ ggml_cuda_mmq_config( - ggml_type type, int nthreads, int occupancy, int I, int J, ggml_cuda_mmq_sram_layout sram_layout, int K_vram, int moe_ncols_min_cc, bool stream_k, bool fallback) : - type(type), nthreads(nthreads), occupancy(occupancy), I(I), J(J), sram_layout(sram_layout), K_vram(K_vram), moe_ncols_min_cc(moe_ncols_min_cc), stream_k(stream_k), fallback(fallback) {} - - constexpr __host__ __device__ bool use_moe_ncols(const int cc) const { - return moe_ncols_min_cc != 0 && cc >= moe_ncols_min_cc; - } + ggml_type type, int nthreads, int occupancy, int I, int J, ggml_cuda_mmq_sram_layout sram_layout, int K_vram, bool use_typical_moe_ncols, bool stream_k, bool fallback) : + type(type), nthreads(nthreads), occupancy(occupancy), I(I), J(J), sram_layout(sram_layout), K_vram(K_vram), use_typical_moe_ncols(use_typical_moe_ncols), stream_k(stream_k), fallback(fallback) {} constexpr __device__ int rows_per_warp() const { #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) @@ -215,7 +211,7 @@ struct ggml_cuda_mmq_config { static_assert((I_) % 32 == 0, "bad I"); \ static_assert((J_) % 8 == 0, "bad J"); \ static_assert((K_vram_) % 256 == 0, "bad K_vram"); \ - return ggml_cuda_mmq_config((type_), (nthreads_), (occupancy_), (I_), (J_), (sram_layout_), (K_vram_), moe_ncols_min_cc, (stream_k_), (fallback_)); \ + return ggml_cuda_mmq_config((type_), (nthreads_), (occupancy_), (I_), (J_), (sram_layout_), (K_vram_), use_typical_moe_ncols, (stream_k_), (fallback_)); \ } \ #include "mmq-config-pascal.cuh" @@ -1481,7 +1477,7 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, if (args.expert_bounds != nullptr && args.nchannels_x > 0) { const int J_max = ggml_cuda_mmq_get_J_max(type, fallback, cc, 128); const ggml_cuda_mmq_config config_max = ggml_cuda_mmq_get_config(type, J_max, fallback, cc); - if (config_max.use_moe_ncols(cc)) { + if (config_max.use_typical_moe_ncols) { // Use the typical expert width only for tile selection. // The launch grid still uses args.ncols_max. const int64_t ncols_typical = (args.ncols_dst + args.nchannels_x - 1) / args.nchannels_x;