From 93247bcd5a11871ec72a601c165c8cff5454ee20 Mon Sep 17 00:00:00 2001 From: raufaser <72879802+raufaser@users.noreply.github.com> Date: Wed, 9 Sep 2026 18:30:27 +0200 Subject: [PATCH 1/2] hip: shared perf (MMQ MoE tiles + mmvdq, prefill plumbing) --- ggml/src/ggml-backend-meta.cpp | 7 +- ggml/src/ggml-cuda/ggml-cuda.cu | 54 +- 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-dp4a.cuh | 3 +- .../src/ggml-cuda/mmq-config-pascal-older.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-vec-dot.cuh | 55 +- ggml/src/ggml-cuda/mmq.cuh | 134 +++- ggml/src/ggml-cuda/mmvdq.cu | 713 ++++++++++++++++++ ggml/src/ggml-cuda/mmvdq.cuh | 29 + src/llama.cpp | 40 +- 16 files changed, 1010 insertions(+), 47 deletions(-) create mode 100644 ggml/src/ggml-cuda/mmvdq.cu create mode 100644 ggml/src/ggml-cuda/mmvdq.cuh diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index c4620eebb01d..81ca2abd5388 100644 --- a/ggml/src/ggml-backend-meta.cpp +++ b/ggml/src/ggml-backend-meta.cpp @@ -1895,7 +1895,12 @@ struct ggml_backend_buffer * ggml_backend_meta_alloc_ctx_tensors_from_buft(struc // Speculative graphs can create more than 16 transient views per source // tensor when a target or draft uses tensor-parallel Meta placement. - constexpr size_t compute_headroom = 32; + // Views of the static tensors that are created between graph evals are stored in the compute + // containers. The number of such views is proportional to the number of tensors in the graph + // that share the buffer, which for hybrid recurrent models with n_rs_seq snapshotting can be + // much larger than 16 per static tensor (e.g. Qwen35: ~2*(n_rs_seq+1) views per recurrent + // layer are created for the conv-state snapshot copies). Size the headroom accordingly. + constexpr size_t compute_headroom = 128; const ggml_init_params params_static = { /*.mem_size =*/ ggml_get_mem_size(ctx), /*.mem_buffer =*/ nullptr, diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 880db1a77ae7..4dca42fb95c0 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -35,6 +35,7 @@ #include "ggml-cuda/mmq.cuh" #include "ggml-cuda/mmvf.cuh" #include "ggml-cuda/mmvq.cuh" +#include "ggml-cuda/mmvdq.cuh" #include "ggml-cuda/moe-weighted-reduction.cuh" #include "ggml-cuda/norm.cuh" #include "ggml-cuda/opt-step-adamw.cuh" @@ -1893,6 +1894,22 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor ggml_cuda_mul_mat_f(ctx, src0, src1, nullptr, dst); return; } + // GGML_CUDA_DQ_Q6K (unset = arch default, 0 = off, non-zero = on). + const bool dq_default = GGML_CUDA_CC_IS_RDNA3_5(cc); + if (ggml_cuda_dq_mmv_enabled(dq_default) && ne11 == 1 + && (src0->type == GGML_TYPE_Q4_K || src0->type == GGML_TYPE_Q5_K + || (src0->type == GGML_TYPE_Q6_K && ggml_cuda_dq_q6k_enabled(dq_default))) + && ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst) + && ne02 == 1 && ne03 == 1 && ne12 == 1 && ne13 == 1 && src0->ne[0] % QK_K == 0) { + if (src0->type == GGML_TYPE_Q4_K) { + ggml_cuda_mul_mat_vec_dq_q4_K(ctx, src0, src1, dst); + } else if (src0->type == GGML_TYPE_Q5_K) { + ggml_cuda_mul_mat_vec_dq_q5_K(ctx, src0, src1, dst); + } else { + ggml_cuda_mul_mat_vec_dq_q6_K(ctx, src0, src1, dst); + } + return; + } if (ggml_cuda_should_use_mmvq(src0->type, cc, ne11)) { ggml_cuda_mul_mat_vec_q(ctx, src0, src1, nullptr, dst); return; @@ -4025,6 +4042,23 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph } if (ggml_cuda_should_fuse_mul_mat_vec_q(up)) { + const bool dq_default = GGML_CUDA_CC_IS_RDNA3_5(ggml_cuda_info().devices[cuda_ctx->device].cc); + if (ggml_cuda_dq_mmv_enabled(dq_default) && ids == nullptr + && ggml_get_glu_op(glu) == GGML_GLU_OP_SWIGLU + && (src0->type == GGML_TYPE_Q4_K || src0->type == GGML_TYPE_Q5_K + || (src0->type == GGML_TYPE_Q6_K && ggml_cuda_dq_q6k_enabled(dq_default))) + && gate->src[0]->type == src0->type && ggml_are_same_shape(src0, gate->src[0]) + && src1->type == GGML_TYPE_F32 && glu->type == GGML_TYPE_F32 + && src1->ne[1] == 1 && src0->ne[0] % QK_K == 0 + && ggml_is_contiguous(src0) && ggml_is_contiguous(gate->src[0]) + && ggml_is_contiguous(src1) && ggml_is_contiguous(glu) + && src0->ne[2] == 1 && src0->ne[3] == 1 && src1->ne[2] == 1 && src1->ne[3] == 1) { + ggml_cuda_mul_mat_vec_dq_glu(*cuda_ctx, src0, gate->src[0], src1, glu); + fused_mul_mat_vec = true; + fused_node_count = 3; + break; + } + ggml_cuda_mm_fusion_args_host fusion_data{}; fusion_data.gate = gate->src[0]; fusion_data.glu_op = ggml_get_glu_op(glu); @@ -4500,6 +4534,15 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, if (graph->is_enabled()) { const bool graph_compatible = ggml_cuda_graph_check_compability(cgraph); if (graph_compatible) { + // PRE-FILL graphs (the graph key node[0].ne[1] > 1) use varying ubatch + // sizes, so each is a separate graph key and CUDA-graph capture never + // amortizes: the per-call update_required probe + failed capture is pure + // overhead. Measured pp512 is ~6.7% faster with graphs OFF. Only decode + // (ne[1]==1, stable shape) benefits from graph replay. Skip the whole + // graph path (incl. the update_required probe) for multi-token graphs. + if (cgraph->n_nodes > 0 && cgraph->nodes[0]->ne[1] > 1) { + use_cuda_graph = false; + } else { const bool properties_changed = ggml_cuda_graph_update_required(cuda_ctx, cgraph); if (!graph->warmup_complete) { @@ -4522,6 +4565,7 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, cuda_graph_update_required = graph->instance == nullptr; } } + } // else: not prefill } } #endif // USE_CUDA_GRAPH @@ -4600,9 +4644,13 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph GGML_UNUSED(cgraph); #endif - static bool enable_graph_optimization = [] { - const char * env = getenv("GGML_CUDA_GRAPH_OPT"); - return env != nullptr && atoi(env) == 1; + static bool enable_graph_optimization = [cuda_ctx] { + const char * env = getenv("GGML_CUDA_GRAPH_OPT"); + if (env != nullptr) { + return atoi(env) == 1; + } + const int cc = ggml_cuda_info().devices[cuda_ctx->device].cc; + return GGML_CUDA_CC_IS_RDNA3_5(cc); }(); if (!enable_graph_optimization) { diff --git a/ggml/src/ggml-cuda/mmq-config-ampere.cuh b/ggml/src/ggml-cuda/mmq-config-ampere.cuh index 9f9fd197382f..2c00aef2ce14 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 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); @@ -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, 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 9fbe32b6972b..8f928e217f2c 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 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 4a8d89f72019..1d51a773b9f2 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 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); @@ -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, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh b/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh index 83eb7c146e11..557a04e1853a 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_pascal_dp4a(ggml_type type, int J, bool fallback) { + 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); @@ -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, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh b/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh index 2a8dc9e1a93e..751ce026d034 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh @@ -1,4 +1,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_pascal_older(ggml_type type, int J, bool fallback) { + 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); @@ -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, 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 8324d9e1a830..c1efef56fb37 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 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); @@ -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, 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 180b2d9370d1..10fdad663ac7 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 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); @@ -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, 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 3a3ef7bd9c09..ba569337b6b5 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 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, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); @@ -270,5 +271,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_NVFP4, 256, 2, 128, 96, 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, 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 9293d9d55885..6cce1d7e82f7 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 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, 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, use_typical_moe_ncols, false, true); } diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index 301e11c5a97a..c2245231eca5 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -339,6 +339,17 @@ template static __device__ __forceinline_ load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*sram_stride + k0, sram_stride); } + // Row scale pair (d, m) is invariant over the j0 loop; load it once per element. + float2 dmA_reg[ntx][tile_C::ne]; +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_A::I + tile_C::get_i(l); + dmA_reg[n][l] = __half22float2(x_dm[i*sram_stride + k0/QI8_1]); + } + } + #pragma unroll for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) { tile_B B; @@ -354,10 +365,8 @@ template static __device__ __forceinline_ #pragma unroll for (int l = 0; l < tile_C::ne; ++l) { - const int i = i0 + n*tile_A::I + tile_C::get_i(l); - float2 dmA = __half22float2(x_dm[i*sram_stride + k0/QI8_1]); - sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA.x*dsB.x*C.x[l]; - sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA.y*dsB.y; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA_reg[n][l].x*dsB.x*C.x[l]; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA_reg[n][l].y*dsB.y; } } } @@ -1028,6 +1037,18 @@ template static __device__ __forceinline_ const int i0 = (threadIdx.y / ntx) * rows_per_warp; + // Row base scales are invariant over the k01 and j0 loops; load them once. + // Each thread owns fixed elements of the C tile, so one value per element suffices. + float x_df_reg[ntx][tile_C::ne]; +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_C::I + tile_C::get_i(l); + x_df_reg[n][l] = x_df[i*sram_stride]; + } + } + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { const int k0 = k00 + k01; @@ -1037,6 +1058,28 @@ template static __device__ __forceinline_ load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*sram_stride + k0, sram_stride); } + // Sub-scales for this k01 chunk; invariant over the j0 loop. + int8_t x_sc_reg[ntx][tile_C::ne]; +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_C::I + tile_C::get_i(l); + x_sc_reg[n][l] = ((const int8_t *) (x_sc + i*sram_stride + k00/16))[k01/4]; + } + } + + // Fold the sub-scale and the row base scale into one f32 per element; + // saves one int-multiply and one convert per element in the j0 loop. + float x_s2_reg[ntx][tile_C::ne]; +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + x_s2_reg[n][l] = (float) x_sc_reg[n][l] * x_df_reg[n][l]; + } + } + #pragma unroll for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) { tile_B B; @@ -1052,9 +1095,7 @@ template static __device__ __forceinline_ #pragma unroll 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] * x_s2_reg[n][l] * dB; } } } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 6367ec88b8ea..281ae0cb594c 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -176,12 +176,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). + 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, 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, 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) @@ -216,7 +217,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_), use_typical_moe_ncols, (stream_k_), (fallback_)); \ } \ #include "mmq-config-pascal-older.cuh" @@ -253,7 +254,7 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_retag_co const ggml_cuda_mmq_config config, const ggml_type type) { return config.type == GGML_TYPE_COUNT ? config : ggml_cuda_mmq_config( type, config.nthreads, config.occupancy, config.I, config.J, config.sram_layout, - config.K_vram, config.stream_k, config.fallback); + config.K_vram, config.use_typical_moe_ncols, config.stream_k, config.fallback); } static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type type, const int J, const bool fallback, const int cc) { @@ -398,6 +399,10 @@ static constexpr __device__ int ggml_cuda_mmq_get_sram_stride(ggml_type type, in static __host__ int ggml_cuda_mmq_get_J_max(const ggml_type type, const bool fallback, const int cc, const int64_t ne11) { int ret = std::min(ne11, int64_t(512)); ret -= ret % 8; + const char * env = getenv("GGML_CUDA_MMQ_J_MAX"); + if (env != nullptr) { + ret = std::min(ret, std::atoi(env)); + } for (;ret > 0; ret -= 8) { if (ggml_cuda_mmq_get_config(type, ret, fallback, cc).type != GGML_TYPE_COUNT) { return ret; @@ -1063,7 +1068,8 @@ static __global__ void mul_mat_q( const uint3 blocks_per_ne00, const int nrows_x, const int ncols_dst, const int stride_row_x, const int ncols_y, const int stride_col_dst, const uint3 channel_ratio, const uint3 nchannels_y, const int stride_channel_x, const int stride_channel_y, const int stride_channel_dst, const uint3 sample_ratio, const uint3 nsamples_y, const int stride_sample_x, const int stride_sample_y, const int stride_sample_dst, - const uint3 ntx) { + const uint3 ntx, + const int32_t * __restrict__ block_expert, const int32_t * __restrict__ block_start, const int n_experts) { // Skip unused template specializations for faster compilation: if (ggml_cuda_mmq_get_config(type, J, fallback).type == GGML_TYPE_COUNT) { @@ -1076,6 +1082,8 @@ static __global__ void mul_mat_q( constexpr int qk = ggml_cuda_type_traits::qk; constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + GGML_UNUSED_VARS(block_expert, block_start, n_experts); + const uint32_t nty = (nrows_x + I - 1) / I; // Number of tiles y // Initialize the ids for writing back data with just the index. @@ -1095,11 +1103,24 @@ static __global__ void mul_mat_q( __syncthreads(); if constexpr (!ggml_cuda_mmq_get_stream_k(type, J, fallback)) { - const uint2 tmp2 = fast_div_modulo(blockIdx.z, nchannels_y); - const int wt = tmp2.x; - const int zt = tmp2.y; - const int jt = blockIdx.y; + int wt; + int zt; + int jt; const int it = blockIdx.x; + if (block_expert != nullptr) { + const int m_block = blockIdx.y; + if (m_block >= block_start[n_experts]) { + return; + } + zt = block_expert[m_block]; + jt = m_block - block_start[zt]; + wt = 0; + } else { + const uint2 tmp2 = fast_div_modulo(blockIdx.z, nchannels_y); + wt = tmp2.x; + zt = tmp2.y; + jt = blockIdx.y; + } // Defaults for regular matrix multiplication: int col_low = 0; @@ -1495,6 +1516,51 @@ static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const i return nbs_ids + nbs_x + GGML_PAD(nbs_y, config.nthreads*sizeof(int)); } +static constexpr int MMQ_MAX_GRIDDIM_Y = 65535; + +static __global__ void mmq_build_moe_block_map( + const int32_t * __restrict__ expert_bounds, const int n_experts, const int J, + int32_t * __restrict__ block_start, int32_t * __restrict__ block_expert) { + extern __shared__ int s_start[]; + const int tid = threadIdx.x; + + for (int e = tid; e < n_experts; e += blockDim.x) { + const int count = expert_bounds[e + 1] - expert_bounds[e]; + s_start[e] = (count + J - 1) / J; + } + __syncthreads(); + + if (tid == 0) { + int sum = 0; + for (int e = 0; e < n_experts; ++e) { + const int count = s_start[e]; + s_start[e] = sum; + sum += count; + } + s_start[n_experts] = sum; + } + __syncthreads(); + + for (int e = tid; e <= n_experts; e += blockDim.x) { + block_start[e] = s_start[e]; + } + + const int total = s_start[n_experts]; + for (int m = tid; m < total; m += blockDim.x) { + int lo = 0; + int hi = n_experts; + while (lo < hi) { + const int mid = (lo + hi) >> 1; + if (s_start[mid] <= m) { + lo = mid + 1; + } else { + hi = mid; + } + } + block_expert[m] = lo - 1; + } +} + template static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { const int id = ggml_cuda_get_device(); @@ -1530,12 +1596,35 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a const uint3 sample_ratio_fd = init_fastdiv_values(sample_ratio); if (!ggml_cuda_mmq_get_stream_k(type, J, fallback, cc)) { - mul_mat_q<<>> + const bool use_compact = args.expert_bounds != nullptr && GGML_CUDA_CC_IS_RDNA3(cc); + const int32_t * block_expert_ptr = nullptr; + const int32_t * block_start_ptr = nullptr; + int n_experts = 0; + dim3 block_nums = block_nums_xy_tiling; + + ggml_cuda_pool & pool = ctx.pool(id); + ggml_cuda_pool_alloc block_start(pool); + ggml_cuda_pool_alloc block_expert(pool); + if (use_compact) { + n_experts = args.nchannels_y; + const int64_t max_m_blocks = (args.ncols_dst + int64_t(n_experts)*(J - 1) + J - 1) / J; + if (max_m_blocks < MMQ_MAX_GRIDDIM_Y) { + block_start.alloc(n_experts + 1); + block_expert.alloc(max_m_blocks); + mmq_build_moe_block_map<<<1, 256, (n_experts + 1)*sizeof(int), stream>>>( + args.expert_bounds, n_experts, J, block_start.ptr, block_expert.ptr); + block_expert_ptr = block_expert.ptr; + block_start_ptr = block_start.ptr; + block_nums = dim3(nty, (unsigned) max_m_blocks, 1); + } + } + + mul_mat_q<<>> (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, args.y_scale, blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, channel_ratio_fd, nchannels_y_fd, args.stride_channel_x, args.stride_channel_y, args.stride_channel_dst, sample_ratio_fd, nsamples_y_fd, args.stride_sample_x, args.stride_sample_y, args.stride_sample_dst, - ntx_fd); + ntx_fd, block_expert_ptr, block_start_ptr, n_experts); return; } @@ -1564,7 +1653,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst, channel_ratio_fd, nchannels_y_fd, args.stride_channel_x, args.stride_channel_y, args.stride_channel_dst, sample_ratio_fd, nsamples_y_fd, args.stride_sample_x, args.stride_sample_y, args.stride_sample_dst, - ntx_fd); + ntx_fd, nullptr, nullptr, 0); if (!fixup_needed) { return; @@ -1583,6 +1672,25 @@ 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 && args.nchannels_x > 0) { + if (GGML_CUDA_CC_IS_RDNA3(cc)) { + const int64_t ncols_per_expert = (args.ncols_dst + args.nchannels_y - 1) / args.nchannels_y; + ncols_picker = std::min(2*ncols_per_expert, args.ncols_max); + } else { + 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_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; + 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; @@ -1596,7 +1704,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; diff --git a/ggml/src/ggml-cuda/mmvdq.cu b/ggml/src/ggml-cuda/mmvdq.cu new file mode 100644 index 000000000000..bd5de8dedad8 --- /dev/null +++ b/ggml/src/ggml-cuda/mmvdq.cu @@ -0,0 +1,713 @@ +#include "mmvdq.cuh" +#include "unary.cuh" + +#include + +// -1 = unset (use arch default), 0 = force off, 1 = force on. +static int dq_env_override(const char * name) { + const char * v = getenv(name); + if (!v) return -1; + return (v[0] == '0' && v[1] == '\0') ? 0 : 1; +} + +bool ggml_cuda_dq_mmv_enabled(bool arch_default) { + static const int ov = dq_env_override("GGML_CUDA_DQ_MMV"); + return ov < 0 ? arch_default : (bool) ov; +} + +bool ggml_cuda_dq_q6k_enabled(bool arch_default) { + static const int ov = dq_env_override("GGML_CUDA_DQ_Q6K"); + return ov < 0 ? arch_default : (bool) ov; +} + +// Rows-per-block tuning knob. Only 1/2/4/8 are instantiated; anything else +// warns once and falls back to 1. Cached so we don't re-read the env per matvec. +static int dq_num_rows_init() { + const char * v = getenv("GGML_CUDA_DQ_ROWS"); + if (!v) return 1; + const int r = atoi(v); + if (r == 1 || r == 2 || r == 4 || r == 8) return r; + GGML_LOG_WARN("%s: unsupported GGML_CUDA_DQ_ROWS=%s (expected 1/2/4/8), using 1\n", __func__, v); + return 1; +} + +static int dq_num_rows() { + static const int rows = dq_num_rows_init(); + return rows; +} + +// Q4_K scale/min extraction, identical to get_scale_min_k4 in convert.cu. +static __device__ __forceinline__ void dq_get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m) { + if (j < 4) { + d = q[j] & 63; m = q[j + 4] & 63; + } else { + d = (q[j+4] & 0xF) | ((q[j-4] >> 6) << 4); + m = (q[j+4] >> 4) | ((q[j-0] >> 6) << 4); + } +} + +// Per-super-block dot for one Q4_K block against pre-loaded (float4) activation. +// Shared by the plain and gate+up-fused kernels so the dequant math lives once. +static __device__ __forceinline__ float dq_dot_q4_K( + const block_q4_K * b, int q_offset, int v_im, + const float4 & by10, const float4 & by132, const float4 & by20, const float4 & by232, + float sum10, float sum32, float sum20, float sum42) { + const float dall = __low2float(b->dm); + const float dmin = __high2float(b->dm); + + uint8_t dx, mx, dy, my, dz, mz, dw, mw; + dq_get_scale_min_k4(2*v_im + 0, b->scales, dx, mx); + dq_get_scale_min_k4(2*v_im + 1, b->scales, dy, my); + dq_get_scale_min_k4(2*v_im + 4, b->scales, dz, mz); + dq_get_scale_min_k4(2*v_im + 5, b->scales, dw, mw); + + const uint32_t * qs32 = (const uint32_t *) b->qs; + const uint32_t qs0 = qs32[q_offset/4 ]; + const uint32_t qs64 = qs32[q_offset/4 + 16]; + + const float sx = + by10.x * (float) ((qs0 >> 0) & 0xF) + by10.y * (float) ((qs0 >> 8) & 0xF) + + by10.z * (float) ((qs0 >> 16) & 0xF) + by10.w * (float) ((qs0 >> 24) & 0xF); + const float sy = + by132.x * (float) ((qs0 >> 4) & 0xF) + by132.y * (float) ((qs0 >> 12) & 0xF) + + by132.z * (float) ((qs0 >> 20) & 0xF) + by132.w * (float) ((qs0 >> 28) & 0xF); + const float sz = + by20.x * (float) ((qs64 >> 0) & 0xF) + by20.y * (float) ((qs64 >> 8) & 0xF) + + by20.z * (float) ((qs64 >> 16) & 0xF) + by20.w * (float) ((qs64 >> 24) & 0xF); + const float sw = + by232.x * (float) ((qs64 >> 4) & 0xF) + by232.y * (float) ((qs64 >> 12) & 0xF) + + by232.z * (float) ((qs64 >> 20) & 0xF) + by232.w * (float) ((qs64 >> 28) & 0xF); + + const float smin = sum10*mx + sum32*my + sum20*mz + sum42*mw; + return dall * (sx*dx + sy*dy + sz*dz + sw*dw) - dmin*smin; +} + +// Per-super-block dot for one Q5_K block against pre-loaded (float2) activation. +// Scale unpacking (s04l/s04h/s8) and the qh bit-plane merges mirror the canonical +// vec_dot_q5_K_q8_1 in vecdotq.cuh ÔÇö keep in sync if that changes. +static __device__ __forceinline__ float dq_dot_q5_K( + const block_q5_K * b, int q_offset, int l0, int v_im, + const float2 & by10, const float2 & by116, const float2 & by132, const float2 & by148, + const float2 & by20, const float2 & by216, const float2 & by232, const float2 & by248, + float smin_x, float smin_y, float smin_z, float smin_w) { + const float dall = __low2float(b->dm); + const float dmin = __high2float(b->dm); + + const uint16_t * sc16 = (const uint16_t *) b->scales; + const uint32_t scale0 = sc16[v_im ]; + const uint32_t scale4 = sc16[v_im + 2]; + const uint32_t scale8 = sc16[v_im + 4]; + const uint32_t s04l = (scale4 << 16) | scale0; + const uint32_t s04h = (s04l & 0xC0C0C0C0u) >> 2; + const uint32_t s04m = s04l & 0x3F3F3F3Fu; + const uint32_t s8 = (((scale8 << 12) | scale8) & 0x0F0F0F0Fu) | s04h; + + const float sc0 = (float) ((s04m >> 0) & 0xFF); + const float sc1 = (float) ((s04m >> 8) & 0xFF); + const float sc2 = (float) ((s04m >> 16) & 0xFF); + const float sc3 = (float) ((s04m >> 24) & 0xFF); + const float sc4 = (float) ((s8 >> 0) & 0xFF); + const float sc5 = (float) ((s8 >> 8) & 0xFF); + const float sc6 = (float) ((s8 >> 16) & 0xFF); + const float sc7 = (float) ((s8 >> 24) & 0xFF); + + const uint16_t * qs16 = (const uint16_t *) b->qs; + const uint32_t qs0 = (uint32_t) qs16[q_offset/2 ] | ((uint32_t) qs16[q_offset/2 + 8] << 16); + const uint32_t qs64 = (uint32_t) qs16[q_offset/2 + 32] | ((uint32_t) qs16[q_offset/2 + 40] << 16); + + uint32_t qs0_lo = qs0 & 0x0F0F0F0Fu; + uint32_t qs0_hi = (qs0 >> 4) & 0x0F0F0F0Fu; + uint32_t qs64_lo = qs64 & 0x0F0F0F0Fu; + uint32_t qs64_hi = (qs64 >> 4) & 0x0F0F0F0Fu; + + const uint16_t * qh16 = (const uint16_t *) b->qh; + const uint32_t qh = (uint32_t) qh16[l0/2] | ((uint32_t) qh16[l0/2 + 8] << 16); + + qs0_lo += ((qh >> (2*v_im)) & 0x01010101u) << 4; + qs0_hi += ((qh >> (2*v_im)) & 0x02020202u) << 3; + qs64_lo += ((qh >> (2*v_im)) & 0x10101010u); + qs64_hi += ((qh >> (2*v_im)) & 0x20202020u) >> 1; + + const float sx = + by10.x * (float) ((qs0_lo >> 0) & 0xFF) + by10.y * (float) ((qs0_lo >> 8) & 0xFF) + + by116.x * (float) ((qs0_lo >> 16) & 0xFF) + by116.y * (float) ((qs0_lo >> 24) & 0xFF); + const float sy = + by132.x * (float) ((qs0_hi >> 0) & 0xFF) + by132.y * (float) ((qs0_hi >> 8) & 0xFF) + + by148.x * (float) ((qs0_hi >> 16) & 0xFF) + by148.y * (float) ((qs0_hi >> 24) & 0xFF); + const float sz = + by20.x * (float) ((qs64_lo >> 0) & 0xFF) + by20.y * (float) ((qs64_lo >> 8) & 0xFF) + + by216.x * (float) ((qs64_lo >> 16) & 0xFF) + by216.y * (float) ((qs64_lo >> 24) & 0xFF); + const float sw = + by232.x * (float) ((qs64_hi >> 0) & 0xFF) + by232.y * (float) ((qs64_hi >> 8) & 0xFF) + + by248.x * (float) ((qs64_hi >> 16) & 0xFF) + by248.y * (float) ((qs64_hi >> 24) & 0xFF); + + const float smin = smin_x*sc2 + smin_y*sc3 + smin_z*sc6 + smin_w*sc7; + return dall * (sx*sc0 + sy*sc1 + sz*sc4 + sw*sc5) - dmin*smin; +} + +// Per-super-block dot for one Q6_K block against pre-loaded (float4) activation. +// The ql/qh bit-plane merges (q0u..q3u) and the -32 bias mirror the canonical +// vec_dot_q6_K_q8_1 in vecdotq.cuh ÔÇö keep in sync if that changes. +static __device__ __forceinline__ float dq_dot_q6_K( + const block_q6_K * b, int ql_offset, int qh_offset, int s_offset, + const float4 & by0, const float4 & by32, const float4 & by64, const float4 & by96) { + const float d = __half2float(b->d); + + const uint16_t * ql16 = (const uint16_t *) b->ql; + const uint32_t ql0 = (uint32_t) ql16[ql_offset/2 ] | ((uint32_t) ql16[ql_offset/2 + 1] << 16); + const uint32_t ql32 = (uint32_t) ql16[ql_offset/2 + 16] | ((uint32_t) ql16[ql_offset/2 + 17] << 16); + + const uint32_t ql0_lo = ql0 & 0x0F0F0F0Fu; + const uint32_t ql0_hi = (ql0 >> 4) & 0x0F0F0F0Fu; + const uint32_t ql32_lo = ql32 & 0x0F0F0F0Fu; + const uint32_t ql32_hi = (ql32 >> 4) & 0x0F0F0F0Fu; + + const uint16_t * qh16 = (const uint16_t *) b->qh; + const uint32_t qh = (uint32_t) qh16[qh_offset/2] | ((uint32_t) qh16[qh_offset/2 + 1] << 16); + + const uint32_t q0u = ql0_lo | ((qh & 0x03030303u) << 4); + const uint32_t q1u = ql32_lo | ((qh & 0x0C0C0C0Cu) << 2); + const uint32_t q2u = ql0_hi | (qh & 0x30303030u); + const uint32_t q3u = ql32_hi | ((qh & 0xC0C0C0C0u) >> 2); + + const int8_t * sc = b->scales + s_offset; + const float sc0 = (float) sc[0]; + const float sc2 = (float) sc[2]; + const float sc4 = (float) sc[4]; + const float sc6 = (float) sc[6]; + + const float sum0 = + by0.x * (float) ((int) ((q0u >> 0) & 0xFF) - 32) + by0.y * (float) ((int) ((q0u >> 8) & 0xFF) - 32) + + by0.z * (float) ((int) ((q0u >> 16) & 0xFF) - 32) + by0.w * (float) ((int) ((q0u >> 24) & 0xFF) - 32); + const float sum1 = + by32.x * (float) ((int) ((q1u >> 0) & 0xFF) - 32) + by32.y * (float) ((int) ((q1u >> 8) & 0xFF) - 32) + + by32.z * (float) ((int) ((q1u >> 16) & 0xFF) - 32) + by32.w * (float) ((int) ((q1u >> 24) & 0xFF) - 32); + const float sum2 = + by64.x * (float) ((int) ((q2u >> 0) & 0xFF) - 32) + by64.y * (float) ((int) ((q2u >> 8) & 0xFF) - 32) + + by64.z * (float) ((int) ((q2u >> 16) & 0xFF) - 32) + by64.w * (float) ((int) ((q2u >> 24) & 0xFF) - 32); + const float sum3 = + by96.x * (float) ((int) ((q3u >> 0) & 0xFF) - 32) + by96.y * (float) ((int) ((q3u >> 8) & 0xFF) - 32) + + by96.z * (float) ((int) ((q3u >> 16) & 0xFF) - 32) + by96.w * (float) ((int) ((q3u >> 24) & 0xFF) - 32); + + return d * (sum0*sc0 + sum1*sc2 + sum2*sc4 + sum3*sc6); +} + +// ---- Q4_K geometry ---- +struct dq_geom_q4_K { + int itid, ix, v_im, q_offset, y_offset; +}; +static __device__ __forceinline__ dq_geom_q4_K dq_setup_q4_K(int tid) { + const int itid = tid % 16; + const int ix = tid / 16; + const int il = itid / 4; + const int ir = itid % 4; + const int v_im = il / 2; + const int v_in = il % 2; + const int l0 = 4 * (2 * ir + v_in); + return { itid, ix, v_im, 32 * v_im + l0, 64 * v_im + l0 }; +} + +// mul_mat_vec_q4_k shader: 16 threads process one super-block, each block +// computes NUM_ROWS output rows, activations are loaded once as float4 and +// reused across rows. No q8_1 activation pass (unlike mul_mat_vec_q). +template +static __global__ void mul_mat_vec_dq_q4_K( + const void * __restrict__ vx, const float * __restrict__ y, float * __restrict__ dst, + const int ncols_x, const int nrows_x) { + const int first_row = num_rows * blockIdx.x; + const int nblocks = ncols_x / QK_K; + const int it_size = warp_size / 16; + + const dq_geom_q4_K g = dq_setup_q4_K(threadIdx.x); + const block_q4_K * x = (const block_q4_K *) vx; + + float sumf[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) sumf[n] = 0.0f; + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + const float4 by10 = *(const float4 *) (yb + g.y_offset ); + const float4 by132 = *(const float4 *) (yb + g.y_offset + 32); + const float4 by20 = *(const float4 *) (yb + g.y_offset + 128); + const float4 by232 = *(const float4 *) (yb + g.y_offset + 160); + + const float sum10 = by10.x + by10.y + by10.z + by10.w; + const float sum32 = by132.x + by132.y + by132.z + by132.w; + const float sum20 = by20.x + by20.y + by20.z + by20.w; + const float sum42 = by232.x + by232.y + by232.z + by232.w; + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + const block_q4_K * b = &x[(int64_t) row * nblocks + i]; + sumf[n] += dq_dot_q4_K(b, g.q_offset, g.v_im, by10, by132, by20, by232, sum10, sum32, sum20, sum42); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float total = warp_reduce_sum(sumf[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) dst[first_row + n] = total; + } +} + +// Fused gate+up SwiGLU dequant matvec for Q4_K: computes up and gate matvecs +// from the shared activation in one pass, writes silu(gate)*up. +template +static __global__ void mul_mat_vec_dq_glu_q4_K( + const void * __restrict__ vx_up, const void * __restrict__ vx_gate, + const float * __restrict__ y, float * __restrict__ dst, + const int ncols_x, const int nrows_x) { + const int first_row = num_rows * blockIdx.x; + const int nblocks = ncols_x / QK_K; + const int it_size = warp_size / 16; + + const dq_geom_q4_K g = dq_setup_q4_K(threadIdx.x); + const block_q4_K * xu = (const block_q4_K *) vx_up; + const block_q4_K * xg = (const block_q4_K *) vx_gate; + + if (num_rows >= 8) { + // Two-pass: one accumulator array at a time keeps register pressure at + // plain-kernel levels. up is reduced and stashed in dst, then pass 2 + // computes gate and combines. Weights are still read once each; only the + // small activation vector y is re-read. Avoids the single-pass VGPR spill. + for (int pass = 0; pass < 2; ++pass) { + const block_q4_K * xw = pass == 0 ? xu : xg; + float acc[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) acc[n] = 0.0f; + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + const float4 by10 = *(const float4 *) (yb + g.y_offset ); + const float4 by132 = *(const float4 *) (yb + g.y_offset + 32); + const float4 by20 = *(const float4 *) (yb + g.y_offset + 128); + const float4 by232 = *(const float4 *) (yb + g.y_offset + 160); + + const float sum10 = by10.x + by10.y + by10.z + by10.w; + const float sum32 = by132.x + by132.y + by132.z + by132.w; + const float sum20 = by20.x + by20.y + by20.z + by20.w; + const float sum42 = by232.x + by232.y + by232.z + by232.w; + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + acc[n] += dq_dot_q4_K(&xw[(int64_t) row * nblocks + i], g.q_offset, g.v_im, by10, by132, by20, by232, sum10, sum32, sum20, sum42); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float r = warp_reduce_sum(acc[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) { + if (pass == 0) dst[first_row + n] = r; + else dst[first_row + n] = ggml_cuda_op_silu_single(r) * dst[first_row + n]; + } + } + } + return; + } + + float up[num_rows], gate[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) { up[n] = 0.0f; gate[n] = 0.0f; } + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + const float4 by10 = *(const float4 *) (yb + g.y_offset ); + const float4 by132 = *(const float4 *) (yb + g.y_offset + 32); + const float4 by20 = *(const float4 *) (yb + g.y_offset + 128); + const float4 by232 = *(const float4 *) (yb + g.y_offset + 160); + + const float sum10 = by10.x + by10.y + by10.z + by10.w; + const float sum32 = by132.x + by132.y + by132.z + by132.w; + const float sum20 = by20.x + by20.y + by20.z + by20.w; + const float sum42 = by232.x + by232.y + by232.z + by232.w; + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + const int64_t off = (int64_t) row * nblocks + i; + up[n] += dq_dot_q4_K(&xu[off], g.q_offset, g.v_im, by10, by132, by20, by232, sum10, sum32, sum20, sum42); + gate[n] += dq_dot_q4_K(&xg[off], g.q_offset, g.v_im, by10, by132, by20, by232, sum10, sum32, sum20, sum42); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float u = warp_reduce_sum(up[n]); + const float gt = warp_reduce_sum(gate[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) dst[first_row + n] = ggml_cuda_op_silu_single(gt) * u; + } +} + +// ---- Q5_K geometry ---- +struct dq_geom_q5_K { + int ix, l0, v_im, q_offset, y_offset; +}; +static __device__ __forceinline__ dq_geom_q5_K dq_setup_q5_K(int tid) { + const int itid = tid % 16; + const int ix = tid / 16; + const int il = itid / 4; + const int ir = itid % 4; + const int v_im = il / 2; + const int v_in = il % 2; + const int l0 = 4 * ir + 2 * v_in; + return { ix, l0, v_im, 32 * v_im + l0, 64 * v_im + l0 }; +} + +#define DQ_Q5_K_LOAD_ACT() \ + const float2 by10 = *(const float2 *) (yb + g.y_offset ); \ + const float2 by116 = *(const float2 *) (yb + g.y_offset + 16); \ + const float2 by132 = *(const float2 *) (yb + g.y_offset + 32); \ + const float2 by148 = *(const float2 *) (yb + g.y_offset + 48); \ + const float2 by20 = *(const float2 *) (yb + g.y_offset + 128); \ + const float2 by216 = *(const float2 *) (yb + g.y_offset + 144); \ + const float2 by232 = *(const float2 *) (yb + g.y_offset + 160); \ + const float2 by248 = *(const float2 *) (yb + g.y_offset + 176); \ + const float smin_x = by10.x + by10.y + by116.x + by116.y; \ + const float smin_y = by132.x + by132.y + by148.x + by148.y; \ + const float smin_z = by20.x + by20.y + by216.x + by216.y; \ + const float smin_w = by232.x + by232.y + by248.x + by248.y + +#define DQ_Q5_K_ARGS by10, by116, by132, by148, by20, by216, by232, by248, smin_x, smin_y, smin_z, smin_w + +template +static __global__ void mul_mat_vec_dq_q5_K( + const void * __restrict__ vx, const float * __restrict__ y, float * __restrict__ dst, + const int ncols_x, const int nrows_x) { + const int first_row = num_rows * blockIdx.x; + const int nblocks = ncols_x / QK_K; + const int it_size = warp_size / 16; + + const dq_geom_q5_K g = dq_setup_q5_K(threadIdx.x); + const block_q5_K * x = (const block_q5_K *) vx; + + float sumf[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) sumf[n] = 0.0f; + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + DQ_Q5_K_LOAD_ACT(); + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + const block_q5_K * b = &x[(int64_t) row * nblocks + i]; + sumf[n] += dq_dot_q5_K(b, g.q_offset, g.l0, g.v_im, DQ_Q5_K_ARGS); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float total = warp_reduce_sum(sumf[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) dst[first_row + n] = total; + } +} + +template +static __global__ void mul_mat_vec_dq_glu_q5_K( + const void * __restrict__ vx_up, const void * __restrict__ vx_gate, + const float * __restrict__ y, float * __restrict__ dst, + const int ncols_x, const int nrows_x) { + const int first_row = num_rows * blockIdx.x; + const int nblocks = ncols_x / QK_K; + const int it_size = warp_size / 16; + + const dq_geom_q5_K g = dq_setup_q5_K(threadIdx.x); + const block_q5_K * xu = (const block_q5_K *) vx_up; + const block_q5_K * xg = (const block_q5_K *) vx_gate; + + if (num_rows >= 8) { + // Two-pass: one accumulator array at a time keeps register pressure at + // plain-kernel levels. up is reduced and stashed in dst, then pass 2 + // computes gate and combines. Weights are still read once each; only the + // small activation vector y is re-read. Avoids the single-pass VGPR spill. + for (int pass = 0; pass < 2; ++pass) { + const block_q5_K * xw = pass == 0 ? xu : xg; + float acc[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) acc[n] = 0.0f; + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + DQ_Q5_K_LOAD_ACT(); + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + acc[n] += dq_dot_q5_K(&xw[(int64_t) row * nblocks + i], g.q_offset, g.l0, g.v_im, DQ_Q5_K_ARGS); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float r = warp_reduce_sum(acc[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) { + if (pass == 0) dst[first_row + n] = r; + else dst[first_row + n] = ggml_cuda_op_silu_single(r) * dst[first_row + n]; + } + } + } + return; + } + + float up[num_rows], gate[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) { up[n] = 0.0f; gate[n] = 0.0f; } + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + DQ_Q5_K_LOAD_ACT(); + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + const int64_t off = (int64_t) row * nblocks + i; + up[n] += dq_dot_q5_K(&xu[off], g.q_offset, g.l0, g.v_im, DQ_Q5_K_ARGS); + gate[n] += dq_dot_q5_K(&xg[off], g.q_offset, g.l0, g.v_im, DQ_Q5_K_ARGS); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float u = warp_reduce_sum(up[n]); + const float gt = warp_reduce_sum(gate[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) dst[first_row + n] = ggml_cuda_op_silu_single(gt) * u; + } +} + +// ---- Q6_K geometry ---- +struct dq_geom_q6_K { + int ix, ql_offset, qh_offset, s_offset, y_offset; +}; +static __device__ __forceinline__ dq_geom_q6_K dq_setup_q6_K(int tid) { + const int itid = tid % 16; + const int ix = tid / 16; + const int v_im = itid / 8; + const int v_in = itid % 8; + const int l0 = 4 * v_in; + const int is = v_in / 4; + return { ix, 64 * v_im + l0, 32 * v_im + l0, 8 * v_im + is, 128 * v_im + l0 }; +} + +template +static __global__ void mul_mat_vec_dq_q6_K( + const void * __restrict__ vx, const float * __restrict__ y, float * __restrict__ dst, + const int ncols_x, const int nrows_x) { + const int first_row = num_rows * blockIdx.x; + const int nblocks = ncols_x / QK_K; + const int it_size = warp_size / 16; + + const dq_geom_q6_K g = dq_setup_q6_K(threadIdx.x); + const block_q6_K * x = (const block_q6_K *) vx; + + float sumf[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) sumf[n] = 0.0f; + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + const float4 by0 = *(const float4 *) (yb + g.y_offset ); + const float4 by32 = *(const float4 *) (yb + g.y_offset + 32); + const float4 by64 = *(const float4 *) (yb + g.y_offset + 64); + const float4 by96 = *(const float4 *) (yb + g.y_offset + 96); + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + const block_q6_K * b = &x[(int64_t) row * nblocks + i]; + sumf[n] += dq_dot_q6_K(b, g.ql_offset, g.qh_offset, g.s_offset, by0, by32, by64, by96); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float total = warp_reduce_sum(sumf[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) dst[first_row + n] = total; + } +} + +template +static __global__ void mul_mat_vec_dq_glu_q6_K( + const void * __restrict__ vx_up, const void * __restrict__ vx_gate, + const float * __restrict__ y, float * __restrict__ dst, + const int ncols_x, const int nrows_x) { + const int first_row = num_rows * blockIdx.x; + const int nblocks = ncols_x / QK_K; + const int it_size = warp_size / 16; + + const dq_geom_q6_K g = dq_setup_q6_K(threadIdx.x); + const block_q6_K * xu = (const block_q6_K *) vx_up; + const block_q6_K * xg = (const block_q6_K *) vx_gate; + + if (num_rows >= 8) { + // Two-pass: one accumulator array at a time keeps register pressure at + // plain-kernel levels. up is reduced and stashed in dst, then pass 2 + // computes gate and combines. Weights are still read once each; only the + // small activation vector y is re-read. Avoids the single-pass VGPR spill. + for (int pass = 0; pass < 2; ++pass) { + const block_q6_K * xw = pass == 0 ? xu : xg; + float acc[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) acc[n] = 0.0f; + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + const float4 by0 = *(const float4 *) (yb + g.y_offset ); + const float4 by32 = *(const float4 *) (yb + g.y_offset + 32); + const float4 by64 = *(const float4 *) (yb + g.y_offset + 64); + const float4 by96 = *(const float4 *) (yb + g.y_offset + 96); + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + acc[n] += dq_dot_q6_K(&xw[(int64_t) row * nblocks + i], g.ql_offset, g.qh_offset, g.s_offset, by0, by32, by64, by96); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float r = warp_reduce_sum(acc[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) { + if (pass == 0) dst[first_row + n] = r; + else dst[first_row + n] = ggml_cuda_op_silu_single(r) * dst[first_row + n]; + } + } + } + return; + } + + float up[num_rows], gate[num_rows]; +#pragma unroll + for (int n = 0; n < num_rows; ++n) { up[n] = 0.0f; gate[n] = 0.0f; } + + for (int i = g.ix; i < nblocks; i += it_size) { + const float * yb = y + (int64_t) i * QK_K; + const float4 by0 = *(const float4 *) (yb + g.y_offset ); + const float4 by32 = *(const float4 *) (yb + g.y_offset + 32); + const float4 by64 = *(const float4 *) (yb + g.y_offset + 64); + const float4 by96 = *(const float4 *) (yb + g.y_offset + 96); + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const int row = min(first_row + n, nrows_x - 1); + const int64_t off = (int64_t) row * nblocks + i; + up[n] += dq_dot_q6_K(&xu[off], g.ql_offset, g.qh_offset, g.s_offset, by0, by32, by64, by96); + gate[n] += dq_dot_q6_K(&xg[off], g.ql_offset, g.qh_offset, g.s_offset, by0, by32, by64, by96); + } + } + +#pragma unroll + for (int n = 0; n < num_rows; ++n) { + const float u = warp_reduce_sum(up[n]); + const float gt = warp_reduce_sum(gate[n]); + if (threadIdx.x == 0 && first_row + n < nrows_x) dst[first_row + n] = ggml_cuda_op_silu_single(gt) * u; + } +} + +// ---- launchers ---- +#define DQ_LAUNCH_PLAIN(KERN, NR) \ + do { \ + const dim3 bn((nrows_x + (NR) - 1) / (NR), 1, 1); \ + const dim3 bd(warp_size, 1, 1); \ + if (warp_size == 64) KERN<64, NR><<>>(vx, y, d, ncols_x, nrows_x); \ + else KERN<32, NR><<>>(vx, y, d, ncols_x, nrows_x); \ + } while (0) + +#define DQ_LAUNCH_GLU(KERN, NR) \ + do { \ + const dim3 bn((nrows_x + (NR) - 1) / (NR), 1, 1); \ + const dim3 bd(warp_size, 1, 1); \ + if (warp_size == 64) KERN<64, NR><<>>(vx_up, vx_gate, y, d, ncols_x, nrows_x); \ + else KERN<32, NR><<>>(vx_up, vx_gate, y, d, ncols_x, nrows_x); \ + } while (0) + +template +static void launch_dq_q4_K(const void * vx, const float * y, float * d, int ncols_x, int nrows_x, int warp_size, cudaStream_t stream) { + DQ_LAUNCH_PLAIN(mul_mat_vec_dq_q4_K, num_rows); +} +template +static void launch_dq_q5_K(const void * vx, const float * y, float * d, int ncols_x, int nrows_x, int warp_size, cudaStream_t stream) { + DQ_LAUNCH_PLAIN(mul_mat_vec_dq_q5_K, num_rows); +} +template +static void launch_dq_q6_K(const void * vx, const float * y, float * d, int ncols_x, int nrows_x, int warp_size, cudaStream_t stream) { + DQ_LAUNCH_PLAIN(mul_mat_vec_dq_q6_K, num_rows); +} +template +static void launch_dq_glu_q4_K(const void * vx_up, const void * vx_gate, const float * y, float * d, int ncols_x, int nrows_x, int warp_size, cudaStream_t stream) { + DQ_LAUNCH_GLU(mul_mat_vec_dq_glu_q4_K, num_rows); +} +template +static void launch_dq_glu_q5_K(const void * vx_up, const void * vx_gate, const float * y, float * d, int ncols_x, int nrows_x, int warp_size, cudaStream_t stream) { + DQ_LAUNCH_GLU(mul_mat_vec_dq_glu_q5_K, num_rows); +} +template +static void launch_dq_glu_q6_K(const void * vx_up, const void * vx_gate, const float * y, float * d, int ncols_x, int nrows_x, int warp_size, cudaStream_t stream) { + DQ_LAUNCH_GLU(mul_mat_vec_dq_glu_q6_K, num_rows); +} + +#define DQ_DISPATCH_ROWS(LAUNCH, ...) \ + switch (dq_num_rows()) { \ + case 2: LAUNCH<2>(__VA_ARGS__); break; \ + case 4: LAUNCH<4>(__VA_ARGS__); break; \ + case 8: LAUNCH<8>(__VA_ARGS__); break; \ + default: LAUNCH<1>(__VA_ARGS__); break; \ + } + +void ggml_cuda_mul_mat_vec_dq_q4_K( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + const int ncols_x = src0->ne[0]; + const int nrows = src0->ne[1]; + cudaStream_t stream = ctx.stream(); + const int warp_size = ggml_cuda_info().devices[ctx.device].warp_size; + const void * vx = src0->data; + const float * y = (const float *) src1->data; + float * d = (float *) dst->data; + DQ_DISPATCH_ROWS(launch_dq_q4_K, vx, y, d, ncols_x, nrows, warp_size, stream); +} + +void ggml_cuda_mul_mat_vec_dq_q5_K( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + const int ncols_x = src0->ne[0]; + const int nrows = src0->ne[1]; + cudaStream_t stream = ctx.stream(); + const int warp_size = ggml_cuda_info().devices[ctx.device].warp_size; + const void * vx = src0->data; + const float * y = (const float *) src1->data; + float * d = (float *) dst->data; + DQ_DISPATCH_ROWS(launch_dq_q5_K, vx, y, d, ncols_x, nrows, warp_size, stream); +} + +void ggml_cuda_mul_mat_vec_dq_q6_K( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + const int ncols_x = src0->ne[0]; + const int nrows = src0->ne[1]; + cudaStream_t stream = ctx.stream(); + const int warp_size = ggml_cuda_info().devices[ctx.device].warp_size; + const void * vx = src0->data; + const float * y = (const float *) src1->data; + float * d = (float *) dst->data; + DQ_DISPATCH_ROWS(launch_dq_q6_K, vx, y, d, ncols_x, nrows, warp_size, stream); +} + +void ggml_cuda_mul_mat_vec_dq_glu( + ggml_backend_cuda_context & ctx, const ggml_tensor * up, const ggml_tensor * gate, + const ggml_tensor * src1, ggml_tensor * dst) { + const int ncols_x = up->ne[0]; + const int nrows = up->ne[1]; + cudaStream_t stream = ctx.stream(); + const int warp_size = ggml_cuda_info().devices[ctx.device].warp_size; + const void * vx_up = up->data; + const void * vx_gate = gate->data; + const float * y = (const float *) src1->data; + float * d = (float *) dst->data; + switch (up->type) { + case GGML_TYPE_Q4_K: DQ_DISPATCH_ROWS(launch_dq_glu_q4_K, vx_up, vx_gate, y, d, ncols_x, nrows, warp_size, stream); break; + case GGML_TYPE_Q5_K: DQ_DISPATCH_ROWS(launch_dq_glu_q5_K, vx_up, vx_gate, y, d, ncols_x, nrows, warp_size, stream); break; + default: DQ_DISPATCH_ROWS(launch_dq_glu_q6_K, vx_up, vx_gate, y, d, ncols_x, nrows, warp_size, stream); break; + } +} diff --git a/ggml/src/ggml-cuda/mmvdq.cuh b/ggml/src/ggml-cuda/mmvdq.cuh new file mode 100644 index 000000000000..bcc656810a77 --- /dev/null +++ b/ggml/src/ggml-cuda/mmvdq.cuh @@ -0,0 +1,29 @@ +#pragma once + +#include "common.cuh" + +// Prototype: dequantize-to-float matvec for K-quant weights (decode, n=1). +// Skips the q8_1 activation quantization pass that mul_mat_vec_q requires. +// On by default where arch_default is true (RDNA3.5); GGML_CUDA_DQ_MMV overrides +// (unset = arch default, 0 = force off, non-zero = force on). + +bool ggml_cuda_dq_mmv_enabled(bool arch_default); + +// Q6_K dequant-float matvec (bandwidth-bound, ~neutral vs mmvq). Same override +// semantics via GGML_CUDA_DQ_Q6K (unset = arch default, 0 = off, non-zero = on). +bool ggml_cuda_dq_q6k_enabled(bool arch_default); + +void ggml_cuda_mul_mat_vec_dq_q4_K( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst); + +void ggml_cuda_mul_mat_vec_dq_q5_K( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst); + +void ggml_cuda_mul_mat_vec_dq_q6_K( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst); + +// Fused gate+up SwiGLU dequant matvec (up = src0, gate = separate weight, shared +// activation src1). Writes silu(gate)*up. Handles Q4_K/Q5_K/Q6_K (dispatch on up->type). +void ggml_cuda_mul_mat_vec_dq_glu( + ggml_backend_cuda_context & ctx, const ggml_tensor * up, const ggml_tensor * gate, + const ggml_tensor * src1, ggml_tensor * dst); diff --git a/src/llama.cpp b/src/llama.cpp index ad8e443882ad..7e7eacbab2b7 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -167,16 +167,21 @@ static bool llama_prepare_model_devices(const llama_model_params & params, llama LLAMA_LOG_ERROR("%s: LLAMA_SPLIT_MODE_TENSOR needs >= 1 devices\n", __func__); return false; } - LLAMA_LOG_INFO("%s: creating a Meta device with %zu devices\n", __func__, n_devs); - for (size_t i = 0; i < n_devs; ++i) { - LLAMA_LOG_INFO("%s: - device %zu: %s\n", __func__, i, ggml_backend_dev_name(params.devices[i])); + if (n_devs == 1) { + // single device: no tensor splitting is possible, skip the Meta wrapper + model->devices.push_back({false, params.devices[0]}); + } else { + LLAMA_LOG_INFO("%s: creating a Meta device with %zu devices\n", __func__, n_devs); + for (size_t i = 0; i < n_devs; ++i) { + LLAMA_LOG_INFO("%s: - device %zu: %s\n", __func__, i, ggml_backend_dev_name(params.devices[i])); + } + model->get_split_state_ud.n_devices = n_devs; + model->get_split_state_ud.model = model; + model->devices.push_back({ + true, ggml_backend_meta_device( + params.devices, n_devs, llama_meta_device_get_split_state, &model->get_split_state_ud) + }); } - model->get_split_state_ud.n_devices = n_devs; - model->get_split_state_ud.model = model; - model->devices.push_back({ - true, ggml_backend_meta_device( - params.devices, n_devs, llama_meta_device_get_split_state, &model->get_split_state_ud) - }); } else { for (ggml_backend_dev_t * dev = params.devices; *dev; ++dev) { model->devices.push_back({false, *dev}); @@ -212,12 +217,17 @@ static bool llama_prepare_model_devices(const llama_model_params & params, llama } GGML_ASSERT(!devs.empty()); - model->get_split_state_ud.n_devices = devs.size(); - model->get_split_state_ud.model = model; - gpus.push_back({ - true, ggml_backend_meta_device( - devs.data(), devs.size(), llama_meta_device_get_split_state, &model->get_split_state_ud) - }); + if (devs.size() == 1) { + // single device: no tensor splitting is possible, skip the Meta wrapper + gpus.push_back({false, devs[0]}); + } else { + model->get_split_state_ud.n_devices = devs.size(); + model->get_split_state_ud.model = model; + gpus.push_back({ + true, ggml_backend_meta_device( + devs.data(), devs.size(), llama_meta_device_get_split_state, &model->get_split_state_ud) + }); + } } else { for (size_t i = 0; i < ggml_backend_dev_count(); ++i) { ggml_backend_dev_t dev = ggml_backend_dev_get(i); From 31a739649545975a656941e2c1cd4ab98a018fc4 Mon Sep 17 00:00:00 2001 From: raufaser <72879802+raufaser@users.noreply.github.com> Date: Fri, 11 Sep 2026 22:32:52 +0200 Subject: [PATCH 2/2] hip: address #154 review (HIP-gated graph skip, strict DQ parsing, J_MAX validation, split-state init) --- docs/beellama-args.md | 8 ++++++++ ggml/src/ggml-cuda/ggml-cuda.cu | 18 ++++++++++------- ggml/src/ggml-cuda/mmq.cuh | 34 ++++++++++++++++++++++++++++++--- ggml/src/ggml-cuda/mmvdq.cu | 7 +++++-- ggml/src/ggml-cuda/mmvdq.cuh | 6 ++++-- src/llama.cpp | 7 +++++++ 6 files changed, 66 insertions(+), 14 deletions(-) diff --git a/docs/beellama-args.md b/docs/beellama-args.md index 811ba6c57bfb..d375936e503d 100644 --- a/docs/beellama-args.md +++ b/docs/beellama-args.md @@ -301,6 +301,14 @@ build host cannot detect it. Pre-Turing support remains runtime-unqualified until matching real devices pass the KVarN parity, memory, and model-smoke tests. +## CUDA/HIP dequant matvec knobs + +| Env var | Default | Behavior | +|---|---|---| +| `GGML_CUDA_DQ_MMV` | Arch default (on for RDNA3.5) | `0` forces the K-quant dequant-float matvec off, `1` forces it on. Unset or anything else warns (when set) and keeps the arch default. | +| `GGML_CUDA_DQ_Q6K` | Arch default (on for RDNA3.5) | Same `0`/`1`/arch-default semantics for the Q6_K dequant-float matvec arm. | +| `GGML_CUDA_DQ_ROWS` | `1` | Rows per block for the dequant matvec kernels. Only `1`/`2`/`4`/`8` are instantiated; anything else warns and uses `1`. | + ## Migration from earlier versions | Earlier spelling or surface | v0.4.0 behavior | Replacement | diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 4dca42fb95c0..47111c0356fc 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1894,7 +1894,7 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor ggml_cuda_mul_mat_f(ctx, src0, src1, nullptr, dst); return; } - // GGML_CUDA_DQ_Q6K (unset = arch default, 0 = off, non-zero = on). + // GGML_CUDA_DQ_Q6K (unset/invalid = arch default, 0 = off, 1 = on). const bool dq_default = GGML_CUDA_CC_IS_RDNA3_5(cc); if (ggml_cuda_dq_mmv_enabled(dq_default) && ne11 == 1 && (src0->type == GGML_TYPE_Q4_K || src0->type == GGML_TYPE_Q5_K @@ -4534,15 +4534,19 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, if (graph->is_enabled()) { const bool graph_compatible = ggml_cuda_graph_check_compability(cgraph); if (graph_compatible) { - // PRE-FILL graphs (the graph key node[0].ne[1] > 1) use varying ubatch - // sizes, so each is a separate graph key and CUDA-graph capture never - // amortizes: the per-call update_required probe + failed capture is pure - // overhead. Measured pp512 is ~6.7% faster with graphs OFF. Only decode - // (ne[1]==1, stable shape) benefits from graph replay. Skip the whole - // graph path (incl. the update_required probe) for multi-token graphs. + // HIP-only: skip the graph path (incl. the update_required probe) for + // multi-token (prefill) graphs. Final-tree A/B (27B Q5_K_S, kvarn6, + // pp512, gfx1100: gate on 466.5 vs gate off 465.4-467.5) shows no + // measurable difference, so this is currently perf-neutral; the + // earlier ~6.7% dev-tree reading did not reproduce and is scheduled + // for investigation. Decode (ne[1]==1, stable shape) keeps replay. +#if defined(GGML_USE_HIP) if (cgraph->n_nodes > 0 && cgraph->nodes[0]->ne[1] > 1) { use_cuda_graph = false; } else { +#else + { +#endif const bool properties_changed = ggml_cuda_graph_update_required(cuda_ctx, cgraph); if (!graph->warmup_complete) { diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 281ae0cb594c..bc0956c2a275 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -396,12 +396,40 @@ static constexpr __device__ int ggml_cuda_mmq_get_sram_stride(ggml_type type, in return ggml_cuda_mmq_get_sram_stride(ggml_cuda_mmq_get_sram_layout(type, J, fallback)); } +static __host__ int ggml_cuda_mmq_get_J_max_env() { + // Cached override cap: -1 = unset/invalid (no cap). Only multiples of 8 in + // 8..512 are accepted; anything else warns once and is ignored (raw atoi + // would turn non-numeric input into 0 and 4/5/0 into a J outside the + // 8..128 config table used below). + static const int cached = []() -> int { + const char * env = getenv("GGML_CUDA_MMQ_J_MAX"); + if (env == nullptr) { + return -1; + } + bool digits = env[0] != '\0'; + for (const char * p = env; *p != '\0'; ++p) { + if (*p < '0' || *p > '9') { + digits = false; + break; + } + } + const int v = digits ? std::atoi(env) : -1; + if (!digits || v < 8 || v > 512 || v % 8 != 0) { + GGML_LOG_WARN("ggml_cuda_mmq_get_J_max_env: unsupported GGML_CUDA_MMQ_J_MAX=%s (expected multiple of 8 in 8..512), ignoring\n", + env); + return -1; + } + return v; + }(); + return cached; +} + static __host__ int ggml_cuda_mmq_get_J_max(const ggml_type type, const bool fallback, const int cc, const int64_t ne11) { int ret = std::min(ne11, int64_t(512)); ret -= ret % 8; - const char * env = getenv("GGML_CUDA_MMQ_J_MAX"); - if (env != nullptr) { - ret = std::min(ret, std::atoi(env)); + const int env_cap = ggml_cuda_mmq_get_J_max_env(); + if (env_cap > 0) { + ret = std::min(ret, env_cap); } for (;ret > 0; ret -= 8) { if (ggml_cuda_mmq_get_config(type, ret, fallback, cc).type != GGML_TYPE_COUNT) { diff --git a/ggml/src/ggml-cuda/mmvdq.cu b/ggml/src/ggml-cuda/mmvdq.cu index bd5de8dedad8..e5f4ef2b1e04 100644 --- a/ggml/src/ggml-cuda/mmvdq.cu +++ b/ggml/src/ggml-cuda/mmvdq.cu @@ -3,11 +3,14 @@ #include -// -1 = unset (use arch default), 0 = force off, 1 = force on. +// -1 = unset/invalid (use arch default), 0 = force off, 1 = force on. static int dq_env_override(const char * name) { const char * v = getenv(name); if (!v) return -1; - return (v[0] == '0' && v[1] == '\0') ? 0 : 1; + if (v[0] == '0' && v[1] == '\0') return 0; + if (v[0] == '1' && v[1] == '\0') return 1; + GGML_LOG_WARN("%s: unsupported %s=%s (expected 0/1), using arch default\n", __func__, name, v); + return -1; } bool ggml_cuda_dq_mmv_enabled(bool arch_default) { diff --git a/ggml/src/ggml-cuda/mmvdq.cuh b/ggml/src/ggml-cuda/mmvdq.cuh index bcc656810a77..3375b8c8b142 100644 --- a/ggml/src/ggml-cuda/mmvdq.cuh +++ b/ggml/src/ggml-cuda/mmvdq.cuh @@ -5,12 +5,14 @@ // Prototype: dequantize-to-float matvec for K-quant weights (decode, n=1). // Skips the q8_1 activation quantization pass that mul_mat_vec_q requires. // On by default where arch_default is true (RDNA3.5); GGML_CUDA_DQ_MMV overrides -// (unset = arch default, 0 = force off, non-zero = force on). +// (unset = arch default, 0 = force off, 1 = force on; anything else warns and +// keeps the arch default). bool ggml_cuda_dq_mmv_enabled(bool arch_default); // Q6_K dequant-float matvec (bandwidth-bound, ~neutral vs mmvq). Same override -// semantics via GGML_CUDA_DQ_Q6K (unset = arch default, 0 = off, non-zero = on). +// semantics via GGML_CUDA_DQ_Q6K (unset = arch default, 0 = off, 1 = on; +// anything else warns and keeps the arch default). bool ggml_cuda_dq_q6k_enabled(bool arch_default); void ggml_cuda_mul_mat_vec_dq_q4_K( diff --git a/src/llama.cpp b/src/llama.cpp index 7e7eacbab2b7..e476ed44d1c8 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -169,6 +169,10 @@ static bool llama_prepare_model_devices(const llama_model_params & params, llama } if (n_devs == 1) { // single device: no tensor splitting is possible, skip the Meta wrapper + // (still publish single-device split state: the callback is user-reachable + // via explicit single-device -ts/--tensor-split). + model->get_split_state_ud.n_devices = 1; + model->get_split_state_ud.model = model; model->devices.push_back({false, params.devices[0]}); } else { LLAMA_LOG_INFO("%s: creating a Meta device with %zu devices\n", __func__, n_devs); @@ -219,6 +223,9 @@ static bool llama_prepare_model_devices(const llama_model_params & params, llama GGML_ASSERT(!devs.empty()); if (devs.size() == 1) { // single device: no tensor splitting is possible, skip the Meta wrapper + // (still publish single-device split state, as above). + model->get_split_state_ud.n_devices = 1; + model->get_split_state_ud.model = model; gpus.push_back({false, devs[0]}); } else { model->get_split_state_ud.n_devices = devs.size();