From 26723931f017a7448014d9864f4f716c3b168a65 Mon Sep 17 00:00:00 2001 From: sunagent Date: Sat, 22 Aug 2026 14:36:13 +0800 Subject: [PATCH] cuda: add Blackwell (SM120) MMVQ parameter table and tune FATTN config for Blackwell - mmvq: new MMVQ_PARAMETERS_BLACKWELL table (SM120). nwarps: NVFP4 decode=1, other types decode=4, batch 2-4=8, batch 5-8=2 (+4-5% tg measured, MTP verify path); rows_per_block=2 at decode. - fattn: (256,256,8,128,2,...,false) Q_in_reg=false, +4-7% measured at 20K ctx. Measured on 2x RTX 5060 Ti (sm_120) with Qwen3.8-27B NVFP4 + MTP spec decode. Long-context (100-200K) sustained 40-50 t/s. --- ggml/src/ggml-cuda/fattn-mma-f16.cuh | 2 +- ggml/src/ggml-cuda/mmvq.cu | 44 +++++++++++++++++++++++++++- 2 files changed, 44 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index 7f4cfd5511f..0ea8d760b46 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -66,7 +66,7 @@ static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_co GGML_CUDA_FATTN_MMA_CONFIG_CASE(192, 128, 32, 128, 2, 32, 96, 64, 64, 2, true); GGML_CUDA_FATTN_MMA_CONFIG_CASE(192, 128, 64, 128, 2, 32, 96, 64, 64, 2, true); - GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 8, 64, 4, 64, 128, 128, 128, 2, true); + GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 8, 128, 2, 64, 128, 128, 128, 2, false); // Q_in_reg=false: faster on SM120 (measured +4-7% at 20K ctx) GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 16, 64, 4, 32, 128, 128, 128, 2, true); GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 32, 128, 2, 32, 128, 128, 128, 2, true); GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 64, 128, 2, 32, 128, 128, 128, 2, true); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index c9992380452..3c0a539ca9d 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -71,7 +71,8 @@ enum mmvq_parameter_table_id { MMVQ_PARAMETERS_RDNA2, MMVQ_PARAMETERS_RDNA3_0, MMVQ_PARAMETERS_RDNA4, - MMVQ_PARAMETERS_GB10 + MMVQ_PARAMETERS_GB10, + MMVQ_PARAMETERS_BLACKWELL }; static constexpr __device__ mmvq_parameter_table_id get_device_table_id() { @@ -85,6 +86,8 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() { return MMVQ_PARAMETERS_GCN; #elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING && __CUDA_ARCH__ < GGML_CUDA_CC_AMPERE return MMVQ_PARAMETERS_TURING; +#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_BLACKWELL + return MMVQ_PARAMETERS_BLACKWELL; #elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK return MMVQ_PARAMETERS_GB10; #else @@ -108,6 +111,9 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) { if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING && ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) { return MMVQ_PARAMETERS_TURING; } + if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_BLACKWELL) { + return MMVQ_PARAMETERS_BLACKWELL; + } if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARK) { return MMVQ_PARAMETERS_GB10; } @@ -389,6 +395,33 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d return 1; } } + if (table_id == MMVQ_PARAMETERS_BLACKWELL) { + // SM120 (RTX 50 series): 256 KB register file, 128 KB shared memory. + // nwarps=1 for NVFP4 at ncols_dst=1 (decode): VDR=8 fits a single warp + // and avoids the shared-memory reduction entirely (measured +5% tg). + // nwarps=4 for other types at ncols_dst=1 (e.g. Q8_0 output layer). + // Batch paths (ncols_dst>=2): nwarps=8 for 2-4 columns, nwarps=2 for 5-8 + // (measured 4-5% faster than nwarps=4 on the ncols_dst=5 decode path). + if (ncols_dst == 1) { + if (type == GGML_TYPE_NVFP4) { + return 1; + } + return 4; + } + switch (ncols_dst) { + case 2: + case 3: + case 4: + return 8; + case 5: + case 6: + case 7: + case 8: + return 2; + default: + return 1; + } + } if (table_id == MMVQ_PARAMETERS_RDNA4) { // nwarps=8 benefits types with simple vec_dot on RDNA4 (ncols_dst=1). // Types with complex vec_dot (Q3_K, IQ2_*, IQ3_*) regress due to register @@ -486,6 +519,15 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d } static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int table_id, bool small_k = false, int nwarps = 1) { + if (table_id == MMVQ_PARAMETERS_BLACKWELL) { + // rows_per_block=2 at decode improves SM occupancy vs 1; 4 was measured neutral. + switch (ncols_dst) { + case 1: + return small_k ? nwarps : 2; + default: + return 2; + } + } if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING || table_id == MMVQ_PARAMETERS_GB10) { switch (ncols_dst) { case 1: