From 25a1d63f4346b472e508c6dbd9ab2ed1d81ace2e Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Mon, 29 Jun 2026 15:24:44 +0200 Subject: [PATCH 1/6] vulkan: use flops instead of weight tensor size for submission heuristic (#25005) * vulkan: extract flops calculation into function * use flops instead of matmul src0 tensor size for submission threshold * use unsigned ints --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 97 +++++++++++++++++----------- 1 file changed, 58 insertions(+), 39 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index ea9191873..ac021c7a6 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -1907,6 +1907,38 @@ static bool vk_enable_sync_logger = false; static uint32_t vk_perf_logger_frequency = 1; static std::string vk_pipeline_stats_filter; +static uint64_t ggml_vk_get_node_flops(const ggml_tensor * node) { + if (node->op == GGML_OP_MUL_MAT || node->op == GGML_OP_MUL_MAT_ID) { + const uint64_t m = node->ne[0]; + const uint64_t n = node->ne[1]; + const uint64_t k = node->src[1]->ne[0]; + const uint64_t batch = node->ne[2] * node->ne[3]; + return m * n * (k + (k - 1)) * batch; + } + if (node->op == GGML_OP_CONV_2D || node->op == GGML_OP_CONV_TRANSPOSE_2D) { + const ggml_tensor * knl = node->src[0]; + const uint64_t Cout = node->ne[2]; + const uint64_t size_K = node->src[1]->ne[2] * knl->ne[0] * knl->ne[1]; + const uint64_t size_N = node->ne[3] * node->ne[0] * node->ne[1]; + return Cout * size_N * (size_K + (size_K - 1)); + } + if (node->op == GGML_OP_CONV_3D) { + const ggml_tensor * knl = node->src[0]; + const uint64_t OC = ggml_get_op_params_i32(node, 11); + const uint64_t IC = ggml_get_op_params_i32(node, 9); + const uint64_t size_K = IC * knl->ne[0] * knl->ne[1] * knl->ne[2]; + const uint64_t size_N = node->ne[3] / OC * node->ne[0] * node->ne[1] * node->ne[2]; + return OC * size_N * (size_K + (size_K - 1)); + } + if (node->op == GGML_OP_FLASH_ATTN_EXT) { + const ggml_tensor * q = node->src[0]; + const ggml_tensor * k = node->src[1]; + const ggml_tensor * v = node->src[2]; + return 2ull * q->ne[1] * q->ne[2] * (k->ne[0] + v->ne[0]) * k->ne[1] * q->ne[3]; + } + return 0; +} + class vk_perf_logger { public: void print_timings(bool force = false) { @@ -1955,7 +1987,7 @@ class vk_perf_logger { } std::string get_node_fusion_name(const ggml_tensor * node, const char *fusion_name, uint64_t *n_flops) { - *n_flops = 0; + *n_flops = ggml_vk_get_node_flops(node); std::string fusion_str; if (fusion_name) { fusion_str = fusion_name + std::string(" "); @@ -1982,35 +2014,22 @@ class vk_perf_logger { if (batch > 1) { name += " batch=" + std::to_string(batch); } - name = fusion_str + name; - *n_flops = m * n * (k + (k - 1)) * batch; - return name; + return fusion_str + name; } if (node->op == GGML_OP_CONV_2D || node->op == GGML_OP_CONV_TRANSPOSE_2D) { std::string name = ggml_op_name(node->op); - ggml_tensor * knl = node->src[0]; - uint64_t OW = node->ne[0]; - uint64_t OH = node->ne[1]; - uint64_t N = node->ne[3]; + const ggml_tensor * knl = node->src[0]; uint64_t Cout = node->ne[2]; - uint64_t KW = knl->ne[0]; - uint64_t KH = knl->ne[1]; - uint64_t Cin = node->src[1]->ne[2]; - // KxCRS @ CRSxNPQ = KxNPQ -> M=K, K=CRS, N=NPQ - uint64_t size_M = Cout; - uint64_t size_K = Cin * KW * KH; - uint64_t size_N = N * OW * OH; - *n_flops = size_M * size_N * (size_K + (size_K - 1)); - name += " M=Cout=" + std::to_string(size_M) + ", K=Cin*KW*KH=" + std::to_string(size_K) + + uint64_t size_K = node->src[1]->ne[2] * knl->ne[0] * knl->ne[1]; + uint64_t size_N = node->ne[3] * node->ne[0] * node->ne[1]; + name += " M=Cout=" + std::to_string(Cout) + ", K=Cin*KW*KH=" + std::to_string(size_K) + ", N=N*OW*OH=" + std::to_string(size_N); - name = fusion_str + name; - return name; + return fusion_str + name; } if (node->op == GGML_OP_RMS_NORM) { std::string name = ggml_op_name(node->op); name += "(" + std::to_string(node->ne[0]) + "," + std::to_string(node->ne[1]) + "," + std::to_string(node->ne[2]) + "," + std::to_string(node->ne[3]) + ")"; - name = fusion_str + name; - return name; + return fusion_str + name; } if (node->op == GGML_OP_FLASH_ATTN_EXT) { const ggml_tensor * dst = node; @@ -2026,7 +2045,6 @@ class vk_perf_logger { " k(" << k->ne[0] << "," << k->ne[1] << "," << k->ne[2] << "," << k->ne[3] << "), " << " v(" << v->ne[0] << "," << v->ne[1] << "," << v->ne[2] << "," << v->ne[3] << "), " << " m(" << (m?m->ne[0]:0) << "," << (m?m->ne[1]:0) << "," << (m?m->ne[2]:0) << "," << (m?m->ne[3]:0) << ")"; - *n_flops = 2ull * q->ne[1] * q->ne[2] * (k->ne[0] + v->ne[0]) * k->ne[1] * q->ne[3]; return name.str(); } if (node->op == GGML_OP_TOP_K) { @@ -2090,7 +2108,7 @@ struct ggml_backend_vk_context { bool do_add_rms_partials_offset_calculation; bool do_add_rms_partials; - uint64_t last_total_mul_mat_bytes {}; + uint64_t last_total_flops {UINT64_MAX}; // Cache most recent tensor that was converted into prealloc_y, and what pipeline it used to convert. vk_pipeline_struct * prealloc_y_last_pipeline_used {}; @@ -16188,22 +16206,23 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg } // Submit after enough work has accumulated, to overlap CPU cmdbuffer generation with GPU execution. - // Estimate the amount of matmul work by looking at the weight matrix size, and submit every 100MB - // (and scaled down based on model size, so smaller models submit earlier). - int submitted_nodes = 0; - int submit_count = 0; - uint64_t mul_mat_bytes = 0; - uint64_t total_mul_mat_bytes = 0; - uint64_t mul_mat_bytes_per_submit = std::min(uint64_t(100*1000*1000), ctx->last_total_mul_mat_bytes / 40u); + // Estimate the amount of compute work using flops, and submit every 200 GFLOP + // (and scaled down based on total graph flops, so smaller models submit earlier). + // Also submit at least every 100 nodes, in case there are workloads without heavy compute. + uint32_t submitted_nodes = 0; + uint32_t submit_count = 0; + uint64_t batch_flops = 0; + uint64_t total_flops = 0; + uint64_t flops_per_submit = std::min(uint64_t(200'000'000'000), ctx->last_total_flops / 40u); for (int i = 0; i < cgraph->n_nodes; i++) { if (first_node_in_batch) { submit_node_idx = i; } - if (cgraph->nodes[i]->op == GGML_OP_MUL_MAT || cgraph->nodes[i]->op == GGML_OP_MUL_MAT_ID) { - auto bytes = ggml_nbytes(cgraph->nodes[i]->src[0]); - mul_mat_bytes += bytes; - total_mul_mat_bytes += bytes; + { + auto node_flops = ggml_vk_get_node_flops(cgraph->nodes[i]); + batch_flops += node_flops; + total_flops += node_flops; } // op_srcs_fused_elementwise indicates whether an op's srcs all contribute to @@ -16415,8 +16434,8 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg // Signal the almost_ready fence when the graph is mostly complete (< 20% remaining) bool almost_ready = (cgraph->n_nodes - i) < cgraph->n_nodes / 5; - bool submit = ((uint32_t)submitted_nodes >= ctx->device->max_nodes_per_submit) || - (mul_mat_bytes_per_submit != 0 && mul_mat_bytes >= mul_mat_bytes_per_submit) || + bool submit = (submitted_nodes >= ctx->device->max_nodes_per_submit) || + (flops_per_submit != 0 && batch_flops >= flops_per_submit) || (i + ctx->num_additional_fused_ops >= last_node) || (almost_ready && !ctx->almost_ready_fence_pending); @@ -16450,9 +16469,9 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg if (submit && enqueued) { first_node_in_batch = true; submitted_nodes = 0; - mul_mat_bytes = 0; + batch_flops = 0; if (submit_count < 3) { - mul_mat_bytes_per_submit *= 2; + flops_per_submit *= 2; } submit_count++; } @@ -16461,7 +16480,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg ctx->fused_ops_write_mask = 0; } - ctx->last_total_mul_mat_bytes = total_mul_mat_bytes; + ctx->last_total_flops = total_flops; if (vk_perf_logger_enabled) { // End the command buffer and submit/wait From 6f4f53f2b7da54fcdbbecaaa734337c337ad6176 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adrien=20Gallou=C3=ABt?= Date: Mon, 29 Jun 2026 17:37:23 +0200 Subject: [PATCH 2/6] common : dedup preset and cached model entries in /v1/models (#25131) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Adrien Gallouët --- common/preset.cpp | 33 +++++++++++++++++++++++++++++---- 1 file changed, 29 insertions(+), 4 deletions(-) diff --git a/common/preset.cpp b/common/preset.cpp index f0cc1fa1a..4362c0621 100644 --- a/common/preset.cpp +++ b/common/preset.cpp @@ -7,6 +7,7 @@ #include #include #include +#include static std::string rm_leading_dashes(const std::string & str) { size_t pos = 0; @@ -16,6 +17,23 @@ static std::string rm_leading_dashes(const std::string & str) { return str.substr(pos); } +static std::string canonical_tag(const std::string & tag) { + static const std::regex re_tag("[-.]([A-Z0-9_]+)$", std::regex::icase); + std::smatch m; + if (std::regex_search(tag, m, re_tag)) { + std::string canon = m[1].str(); + for (char & c : canon) { + c = (char) std::toupper((unsigned char) c); + } + return canon; + } + std::string upper = tag; + for (char & c : upper) { + c = (char) std::toupper((unsigned char) c); + } + return upper; +} + std::vector common_preset::to_args(const std::string & bin_path) const { std::vector args; @@ -270,11 +288,18 @@ common_presets common_preset_context::load_from_ini(const std::string & path, co for (auto section : ini_data) { common_preset preset; - if (section.first.empty()) { - preset.name = COMMON_PRESET_DEFAULT_NAME; - } else { - preset.name = section.first; + std::string section_name = section.first.empty() ? std::string(COMMON_PRESET_DEFAULT_NAME) : section.first; + if (section_name != "*" && section_name != COMMON_PRESET_DEFAULT_NAME) { + auto colon_idx = section_name.rfind(':'); + if (colon_idx != std::string::npos) { + std::string tag = section_name.substr(colon_idx + 1); + std::string canon_tag = canonical_tag(tag); + if (canon_tag != tag) { + section_name = section_name.substr(0, colon_idx + 1) + canon_tag; + } + } } + preset.name = section_name; LOG_DBG("loading preset: %s\n", preset.name.c_str()); for (const auto & [key, value] : section.second) { if (key == "version") { From 86b94708f22478f900b76ca02e316f4f3418faff Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Tue, 30 Jun 2026 02:41:45 +0200 Subject: [PATCH 3/6] Revert "sched : reintroduce less synchronizations during split compute (#20793)" (#25138) --- ggml/src/ggml-backend.cpp | 10 +++------- ggml/src/ggml-cuda/ggml-cuda.cu | 24 ++++-------------------- 2 files changed, 7 insertions(+), 27 deletions(-) diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index 644130cd1..87615921c 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -1551,8 +1551,6 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s int split_backend_id = split->backend_id; ggml_backend_t split_backend = sched->backends[split_backend_id]; - ggml_backend_synchronize(split_backend); - // copy the input tensors to the split backend for (int input_id = 0; input_id < split->n_inputs; input_id++) { ggml_backend_t input_backend = ggml_backend_sched_get_tensor_backend(sched, split->inputs[input_id]); @@ -1563,15 +1561,15 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s // inputs from the user must be copied immediately to prevent the user overwriting the data before the copy is done if (sched->events[split_backend_id][sched->cur_copy] != NULL) { ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]); - } else if (!split_backend->iface.cpy_tensor_async) { + } else { ggml_backend_synchronize(split_backend); } - ggml_backend_tensor_copy_async(input_backend, split_backend, input, input_cpy); + ggml_backend_tensor_copy(input, input_cpy); } else { // wait for the split backend to finish using the input before overwriting it if (sched->events[split_backend_id][sched->cur_copy] != NULL) { ggml_backend_event_wait(split_backend, sched->events[split_backend_id][sched->cur_copy]); - } else if (!split_backend->iface.cpy_tensor_async) { + } else { ggml_backend_synchronize(split_backend); } @@ -1676,8 +1674,6 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } } - ggml_backend_synchronize(split_backend); - if (!sched->callback_eval) { enum ggml_status ec = ggml_backend_graph_compute_async(split_backend, &split->graph); if (ec != GGML_STATUS_SUCCESS) { diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f3fb32452..cca70592f 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3192,24 +3192,11 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ ggml_backend_buffer_t buf_src = src->view_src ? src->view_src->buffer : src->buffer; ggml_backend_buffer_t buf_dst = dst->view_src ? dst->view_src->buffer : dst->buffer; - // Enables async copies from CPU to CUDA, instead of only CUDA-to-CUDA - // Excluding this path for HIP and MUSA as a precaution. - // According to the summary in https://github.com/ggml-org/llama.cpp/pull/20793#issuecomment-4275794315, this change is not beneficial for hip anyways. - // Additionally, there is a lot of anectodal evidence that hip/musa stream behavior might not always 1:1 match CUDA behavior. - // e.g. https://github.com/ROCm/rocm-systems/issues/5109 - // It thus makes sense to exclude this path for HIP and MUSA. This PR was not aimed these backends, the majority of testing happened on CUDA. - // This can be revisited in the future if enabling copy_from_host benefits hip/MUSA, and if the PR author can extensively test on these backends. -#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) - const bool copy_from_host = false; -#else - const bool copy_from_host = ggml_backend_buffer_is_host(buf_src) && ggml_backend_dev_type(backend_src->device) == GGML_BACKEND_DEVICE_TYPE_CPU; -#endif - - if (!(copy_from_host || ggml_backend_is_cuda(backend_src)) || !ggml_backend_is_cuda(backend_dst)) { + if (!ggml_backend_is_cuda(backend_src) || !ggml_backend_is_cuda(backend_dst)) { return false; } - if (!(copy_from_host || ggml_backend_buffer_is_cuda(buf_src)) || !ggml_backend_buffer_is_cuda(buf_dst)) { + if (!ggml_backend_buffer_is_cuda(buf_src) || !ggml_backend_buffer_is_cuda(buf_dst)) { return false; } @@ -3220,17 +3207,14 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ ggml_backend_cuda_buffer_context * buf_ctx_src = (ggml_backend_cuda_buffer_context *) buf_src->context; ggml_backend_cuda_buffer_context * buf_ctx_dst = (ggml_backend_cuda_buffer_context *) buf_dst->context; - if ((copy_from_host && cuda_ctx_dst->device != buf_ctx_dst->device) || - !copy_from_host && (cuda_ctx_src->device != buf_ctx_src->device || cuda_ctx_dst->device != buf_ctx_dst->device)) { + if (cuda_ctx_src->device != buf_ctx_src->device || cuda_ctx_dst->device != buf_ctx_dst->device) { #ifndef NDEBUG GGML_LOG_DEBUG("%s: backend and buffer devices do not match\n", __func__); #endif // NDEBUG return false; } - if (copy_from_host) { - CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyHostToDevice, cuda_ctx_dst->stream())); - } else if (backend_src != backend_dst) { + if (backend_src != backend_dst) { // copy on src stream if (cuda_ctx_src->device == cuda_ctx_dst->device) { CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyDeviceToDevice, cuda_ctx_src->stream())); From 6c5de1cc83537bce5616ed08474f6fe119973a27 Mon Sep 17 00:00:00 2001 From: Masashi Yoshimura Date: Tue, 30 Jun 2026 17:20:04 +0900 Subject: [PATCH 4/6] ggml-webgpu: add support for NVFP4 (#25143) --- .../ggml-webgpu/ggml-webgpu-shader-lib.hpp | 7 +++ ggml/src/ggml-webgpu/ggml-webgpu.cpp | 3 ++ .../wgsl-shaders/common_decls.tmpl | 18 +++++++- .../ggml-webgpu/wgsl-shaders/get_rows.wgsl | 21 +++++++++ .../wgsl-shaders/mul_mat_decls.tmpl | 43 ++++++++++++++++- .../wgsl-shaders/mul_mat_vec_acc.tmpl | 46 +++++++++++++++++++ 6 files changed, 135 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp b/ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp index c00a2e9ee..80a16d16d 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp @@ -1563,6 +1563,7 @@ class ggml_webgpu_shader_lib { case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ4_NL: case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: { // Quantized types using u32 buffers for portability. defines.push_back("SRC_TYPE=u32"); @@ -1593,6 +1594,8 @@ class ggml_webgpu_shader_lib { } else if ((key.src_type >= GGML_TYPE_Q4_0 && key.src_type <= GGML_TYPE_Q8_1) || key.src_type == GGML_TYPE_IQ4_NL || key.src_type == GGML_TYPE_MXFP4) { defines.push_back("BLOCK_SIZE=32u"); + } else if (key.src_type == GGML_TYPE_NVFP4) { + defines.push_back("BLOCK_SIZE=64u"); } else if (key.src_type >= GGML_TYPE_Q2_K) { defines.push_back("BLOCK_SIZE=256u"); } else { @@ -1960,6 +1963,7 @@ class ggml_webgpu_shader_lib { defines.push_back(type_upper + "_TABLES"); break; case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: defines.push_back(type_upper + "_LUT"); break; default: @@ -2103,6 +2107,7 @@ class ggml_webgpu_shader_lib { defines.push_back(type_upper + "_TABLES"); break; case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: defines.push_back(type_upper + "_LUT"); break; default: @@ -2274,6 +2279,7 @@ class ggml_webgpu_shader_lib { defines.push_back(type_upper + "_TABLES"); break; case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: defines.push_back(type_upper + "_LUT"); break; default: @@ -2394,6 +2400,7 @@ class ggml_webgpu_shader_lib { defines.push_back(type_upper + "_TABLES"); break; case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: defines.push_back(type_upper + "_LUT"); break; default: diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index f0ec18abd..29025e9ba 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -4056,6 +4056,7 @@ static bool ggml_webgpu_supported_qtype(ggml_type type) { case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: return true; default: return false; @@ -4156,6 +4157,7 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: supports_op = true; break; default: @@ -4196,6 +4198,7 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_MXFP4: + case GGML_TYPE_NVFP4: supports_op = true; break; default: diff --git a/ggml/src/ggml-webgpu/wgsl-shaders/common_decls.tmpl b/ggml/src/ggml-webgpu/wgsl-shaders/common_decls.tmpl index 758efa17d..6634fbd65 100644 --- a/ggml/src/ggml-webgpu/wgsl-shaders/common_decls.tmpl +++ b/ggml/src/ggml-webgpu/wgsl-shaders/common_decls.tmpl @@ -896,9 +896,23 @@ const kvalues_iq4nl = array( #endif -#ifdef MXFP4_LUT +#if defined(MXFP4_LUT) || defined(NVFP4_LUT) const kvalues_mxfp4 = array( 0, 1, 2, 3, 4, 6, 8, 12, 0, -1, -2, -3, -4, -6, -8, -12 ); -#endif +#endif // MXFP4_LUT || NVFP4_LUT +#ifdef NVFP4_LUT +fn ue4m3_to_fp32(u: u32) -> f32 { + if (u == 0u || u == 127u) { + return 0.0; + } + let exp = (u >> 3u) & 15u; + let man = u & 7u; + if (exp == 0u) { + return f32(man) * (1.0 / 512.0); + } + let bits = ((exp + 120u) << 23u) | (man << 20u); + return bitcast(bits); +} +#endif // NVFP4_LUT diff --git a/ggml/src/ggml-webgpu/wgsl-shaders/get_rows.wgsl b/ggml/src/ggml-webgpu/wgsl-shaders/get_rows.wgsl index 78d61a93d..487edb327 100644 --- a/ggml/src/ggml-webgpu/wgsl-shaders/get_rows.wgsl +++ b/ggml/src/ggml-webgpu/wgsl-shaders/get_rows.wgsl @@ -672,6 +672,27 @@ fn copy_elements(src_base: u32, dst_base: u32, offset: u32) { } #endif +#ifdef NVFP4 +fn copy_elements(src_base: u32, dst_base: u32, offset: u32) { + let block_byte_base = (src_base + offset) * 36; + let d_word = load_u32_at_src(block_byte_base); + for (var sub: u32 = 0u; sub < 4; sub++) { + let d = ue4m3_to_fp32(get_byte(d_word, sub)) * 0.5; + for (var j: u32 = 0u; j < 2; j++) { + let q_packed = load_u32_at_src(block_byte_base + 4 + sub * 8 + j * 4); + for (var k: u32 = 0; k < 4; k++) { + let q_byte = get_byte(q_packed, k); + let q_lo = f32(kvalues_mxfp4[q_byte & 0xFu]) * d; + let q_hi = f32(kvalues_mxfp4[(q_byte >> 4) & 0xF]) * d; + let dst_offset = dst_base + offset * 64 + sub * 16 + j * 4 + k; + dst[dst_offset] = q_lo; + dst[dst_offset + 8u] = q_hi; + } + } + } +} +#endif + @group(0) @binding(0) var src: array; diff --git a/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_decls.tmpl b/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_decls.tmpl index 6a2eb8c82..13996ab51 100644 --- a/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_decls.tmpl +++ b/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_decls.tmpl @@ -241,7 +241,7 @@ fn init_shmem_src0(thread_id: u32, batch_offset: u32, offset_m: u32, k_outer: u3 #endif // INIT_SRC0_SHMEM_Q8_1 #if defined(INIT_SRC0_SHMEM_MXFP4) - let block_byte_base = src0_idx * 17u; + let block_byte_base = src0_idx * 17u; // BLOCK_SIZE_BYTES = 17u; let eu8 = get_byte(load_u32_at_src0_aligned(block_byte_base), block_byte_base & 3u); let e = ldexp(1.0, i32(eu8) - 128); @@ -263,6 +263,47 @@ fn init_shmem_src0(thread_id: u32, batch_offset: u32, offset_m: u32, k_outer: u3 } #endif // legacy-quants +#if defined(INIT_SRC0_SHMEM_NVFP4) +const BLOCK_SIZE = 64u; +const BLOCK_SIZE_BYTES = 36u; +const SUB_BLOCK_SIZE = 16u; // elements sharing one UE4M3 scale +const NQ = 16u; +const BYTES_PER_THREAD = 8u; +const BYTES_PER_INNER_LOOP = 4u; + +fn init_shmem_src0(thread_id: u32, batch_offset: u32, offset_m: u32, k_outer: u32) { + for (var i = thread_id * NQ; i < TILE_SRC0_SHMEM; i += TOTAL_WORKGROUP_SIZE * NQ) { + let tile_m = i / TILE_K; + let tile_k_start = i % TILE_K; + let global_m = offset_m + tile_m; + let global_k_start = k_outer + tile_k_start; + + if (global_m >= params.m) { + break; + } + + let block_k = global_k_start / BLOCK_SIZE; + let sub_block = (global_k_start % BLOCK_SIZE) / SUB_BLOCK_SIZE; + let src0_idx = batch_offset + global_m * params.stride_01 + block_k; + + let block_byte_base = src0_idx * BLOCK_SIZE_BYTES; + let d_byte_base = block_byte_base; + let qs_byte_base = block_byte_base + 4u; + + let d = ue4m3_to_fp32(get_byte(load_u32_at_src0_aligned(d_byte_base), sub_block)) * 0.5; + + for (var j = 0u; j < BYTES_PER_THREAD / BYTES_PER_INNER_LOOP; j++) { + let q_packed = load_u32_at_src0_aligned(qs_byte_base + sub_block * 8u + j * 4u); + for (var k = 0u; k < BYTES_PER_INNER_LOOP; k++) { + let q_byte = get_byte(q_packed, k); + shmem[i + j * BYTES_PER_INNER_LOOP + k] = f16(f32(kvalues_mxfp4[q_byte & 0xF]) * d); + shmem[i + j * BYTES_PER_INNER_LOOP + k + 8u] = f16(f32(kvalues_mxfp4[(q_byte >> 4) & 0xF]) * d); + } + } + } +} +#endif // INIT_SRC0_SHMEM_NVFP4 + // k-quants #if defined(INIT_SRC0_SHMEM_Q2_K) || defined(INIT_SRC0_SHMEM_Q3_K) || defined(INIT_SRC0_SHMEM_Q4_K) || defined(INIT_SRC0_SHMEM_Q5_K) || defined(INIT_SRC0_SHMEM_Q6_K) const BLOCK_SIZE = 256u; diff --git a/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_vec_acc.tmpl b/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_vec_acc.tmpl index b0703fe90..8fd0d1907 100644 --- a/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_vec_acc.tmpl +++ b/ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_vec_acc.tmpl @@ -1505,3 +1505,49 @@ fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src return acc; } #endif + +#ifdef MUL_ACC_NVFP4 +#define BLOCK_SIZE 64 +#define BLOCK_SIZE_BYTES 36 +#define THREADS_PER_BLOCK 4 +#define ELEMS_PER_THREAD (BLOCK_SIZE/THREADS_PER_BLOCK) +fn accumulate_vec_dot(thread_id: u32, row_base: u32, src0_batch_offset: u32, src1_idx_base: u32) -> array, NUM_COLS> { + var acc: array, NUM_COLS>; + + let num_blocks = params.k / BLOCK_SIZE; + let sub = thread_id % THREADS_PER_BLOCK; + for (var block = thread_id/THREADS_PER_BLOCK; block < num_blocks; block += WG_SIZE/THREADS_PER_BLOCK) { + let x_base = src1_idx_base + block * BLOCK_SIZE + sub * ELEMS_PER_THREAD; + var x_block: array, NUM_COLS>; + for (var col = 0u; col < NUM_COLS;col += 1) { + for (var i = 0u; i < ELEMS_PER_THREAD / 2; i++) { + x_block[col][i] = f32(src1[x_base + col * params.stride_11 + i]); + x_block[col][i + 8] = f32(src1[x_base + col * params.stride_11 + i + 8]); + } + } + for (var row = 0u; row < OUTPUTS_PER_WG; row++) { + let output_row = row_base + row; + if (output_row < params.m) { + let block_byte_base = (src0_batch_offset + output_row * params.stride_01 + block) * BLOCK_SIZE_BYTES; + let d = ue4m3_to_fp32(get_byte(load_u32_at_src0_aligned(block_byte_base), sub)) * 0.5; + let q_w0 = load_u32_at_src0_aligned(block_byte_base + 4u + 8u * sub); + let q_w1 = load_u32_at_src0_aligned(block_byte_base + 8u + 8u * sub); + for (var col = 0u;col < NUM_COLS;col += 1) { + var row_sum = 0.0; + for (var l = 0u; l < 8u; l++) { + let q_word = select(q_w0, q_w1, l >= 4u); + let q_byte = get_byte(q_word, l % 4u); + let q_lo = f32(kvalues_mxfp4[q_byte & 0xFu]) * d; + let q_hi = f32(kvalues_mxfp4[(q_byte >> 4u) & 0xFu]) * d; + row_sum += q_lo * x_block[col][l]; + row_sum += q_hi * x_block[col][l + 8u]; + } + acc[col][row] += row_sum; + } + } + } + } + + return acc; +} +#endif From d9df11006f6d3cc34772379bf6897f892874656f Mon Sep 17 00:00:00 2001 From: zduford Date: Tue, 30 Jun 2026 05:51:38 -0400 Subject: [PATCH 5/6] HIP: use hipBLAS for dense prefill on gfx900, keep MMQ for MoE (#24588) * HIP: keep MMQ for gfx900 MoE and Q8_0, use hipBLAS for dense K-quants Assisted-by: GitHub Copilot CLI * HIP: tighten conditional block to be explicitly for gfx900 * HIP: Further simplified gfx900 conditional block * removed unnecessary comment --- ggml/src/ggml-cuda/mmq.cu | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index e1add5e03..6b3b0d064 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -368,5 +368,12 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t return true; } + // gfx900 (Vega 10) lacks native dp4a, loses to dequant + hipBLAS + // for dense matrices; keep MMQ only for MoE, where the + // hipBLAS path is much slower. + if (cc == GGML_CUDA_CC_VEGA) { + return n_experts > 0; + } + return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; } From f708a5b2caaee0226c0af220e366785699ba41e2 Mon Sep 17 00:00:00 2001 From: Kevin Liu <4396kevinliu@gmail.com> Date: Tue, 30 Jun 2026 06:27:38 -0400 Subject: [PATCH 6/6] vulkan: roll bk loop in matmul for asahi linux (#24663) * vulkan: roll bk loop in matmul for asahi linux * vulkan: fix inline comment * vulkan: revert BK-loop unroll change * vulkan: edit spirv directly for asahi roll bk loop * vulkan: remove trailing whitespace at the end of comments --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 95 ++++++++++++++++++++++++++++ 1 file changed, 95 insertions(+) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index ac021c7a6..c0ab9f1c6 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -2475,6 +2475,85 @@ static bool ggml_vk_strip_decode_vector(const uint32_t * code, size_t word_count return true; } +// Remove the loop unrolling hint of the matmul shader's BK loop +// and replace it with the dont_unroll hint for better performance on +// hardware like Apple M1/M2. +// Assumes 1. code comes from mul_mm.comp 2. the K-tile loop has no loop +// control hint and 3. the BK loop is the last loop nested directly inside +// the K-tile loop. +// Returns true when the input was modified; returns false otherwise +// without touching `out`. +static bool ggml_vk_roll_bk_loop(const uint32_t * code, size_t word_count, std::vector & out) { + if (word_count < 5) { + return false; + } + + struct vk_spv_loop { + size_t header; + size_t end; + uint32_t control; + }; + + std::vector loops; + + // Collect a list of all loops in the module. + for (size_t pos = 5; pos < word_count; ) { + const uint32_t wc = code[pos] >> spv::WordCountShift; + const uint32_t op = code[pos] & spv::OpCodeMask; + if (wc == 0 || pos + wc > word_count) { + return false; + } + + if (op == spv::OpLoopMerge && wc >= 4) { loops.push_back({ pos, 0, code[pos + 3] }); } + + if (op == spv::OpLabel && wc >= 2) { + for (auto & l : loops) { + if (l.end == 0 && code[l.header + 1] == code[pos + 1]) { l.end = pos; } + } + } + + pos += wc; + } + + auto encloses = [](const vk_spv_loop & a, const vk_spv_loop & b) { + return a.header < b.header && b.header < a.end; + }; + + // Find the BK loop. + const vk_spv_loop * bk = nullptr; + for (const auto & h : loops) { + if (h.control != spv::LoopControlUnrollMask) { + continue; + } + const vk_spv_loop * parent = nullptr; + bool has_child = false; + for (const auto & g : loops) { + if (encloses(g, h) && (!parent || g.header > parent->header)) { + parent = &g; + } + if (encloses(h, g)) { + has_child = true; + } + } + // BK loop should be the last loop nested inside the loop with no hint + // and have at least one child loop. + if (parent && + parent->control == spv::LoopControlMaskNone && + has_child && + (!bk || h.header > bk->header)) { + bk = &h; + } + } + if (!bk) { + return false; + } + + // set DontUnroll instead of Unroll + out.assign(code, code + word_count); + out[bk->header + 3] = spv::LoopControlDontUnrollMask; + return true; +} + static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipeline, size_t spv_size, const void* spv_data, const std::string entrypoint, uint32_t parameter_count, std::array wg_denoms, std::vector specialization_constants, bool disable_robustness, bool require_full_subgroups, uint32_t required_subgroup_size) { @@ -2558,6 +2637,22 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin } #endif +#if VK_HEADER_VERSION >= 287 + // Roll the mul_mm BK loop on Asahi Linux. Skip bf16 and the mul_mmq pipelines. + if (device->driver_id == vk::DriverId::eMesaHoneykrisp && + pipeline->name.rfind("matmul", 0) == 0 && + pipeline->name.find("bf16") == std::string::npos && + pipeline->name.find("q8_1") == std::string::npos) { + const uint32_t * src = spirv.empty() ? reinterpret_cast(spv_data) : spirv.data(); + size_t src_n = spirv.empty() ? spv_size / sizeof(uint32_t) : spirv.size(); + std::vector rolled; + if (ggml_vk_roll_bk_loop(src, src_n, rolled)) { + spirv = std::move(rolled); + shader_module_create_info = vk::ShaderModuleCreateInfo({}, spirv.size() * sizeof(uint32_t), spirv.data()); + } + } +#endif + pipeline->shader_module = device->device.createShaderModule(shader_module_create_info); vk::PushConstantRange pcr(