From b8fed9fae4b1af1479646f6af183628442441d2b Mon Sep 17 00:00:00 2001 From: Robert Esclapez Garcia Date: Wed, 29 Jul 2026 02:52:17 -0700 Subject: [PATCH] tools: add mmq-tune, an MMQ tile-width autotuning harness Measures the MMQ kernel at a forced mmq_x over a catalog of shapes taken from real models, so tile-width decisions can be made from measurement rather than from an occupancy model alone. The core mirrors both branches of ggml_cuda_mul_mat_q and then calls launch_mul_mat_q directly. For MoE it builds the ids tensor itself and drives ggml_cuda_launch_mm_ids_helper, so expert_bounds and the compacted grid are produced exactly as in production. How the measurement is set up: - Routing is drawn from a Zipf distribution, which is what measured MoE routing follows; a uniform distribution would be the one case the tile heuristic already assumes. - Expert weights are generated distinctly per expert rather than cycling a few slices, which would shrink the weight working set far below the real one. - Correctness is decided by majority vote over a position-weighted checksum of the output. Diffing against one chosen reference cannot work, because a wrong reference makes every other config look broken; the position weight is what makes a misplaced row visible. - Timing interleaves one iteration per config per pass, so clock drift spreads across all configs instead of penalising whichever runs late. The catalog holds the 99 distinct mul_mat_q shapes that 20 models execute (Qwen2.5/3/3.5/3.6, Llama-3.1, gemma-2/3/4, SmolLM2, GLM-4.7; 0.5B to 35B, dense and MoE), each tagged with its invocation count so aggregates reflect real time. The token count is a sweep parameter rather than part of the catalog, and it matters: the dense q8_0 mmq_x=64 cliff only appears around 64 tokens while the MoE tile-width gap needs 2048. mmq_y and the warp count are device-side constexpr, so sweeping them needs a rebuild. They become GGML_CUDA_MMQ_Y_RDNA3_5 / GGML_CUDA_MMQ_NWARPS_RDNA3_5, wired to the GGML_HIP_MMQ_Y / GGML_HIP_MMQ_NWARPS cache variables, with sweep.py driving the rebuild loop. Defaults are unchanged, and only three pairs compile at all since the kernel asserts nwarps*tile_C::I == mmq_y; 64/4 measured best of the three. All of it is behind GGML_MMQ_TUNE (default OFF), which keeps it out of the build entirely. --- ggml/CMakeLists.txt | 3 + ggml/include/ggml-mmq-tune.h | 55 +++++ ggml/src/ggml-cuda/CMakeLists.txt | 5 + ggml/src/ggml-cuda/mmq-tune.cu | 386 ++++++++++++++++++++++++++++++ ggml/src/ggml-cuda/mmq.cuh | 17 +- ggml/src/ggml-hip/CMakeLists.txt | 17 ++ tools/CMakeLists.txt | 3 + tools/mmq-tune/CMakeLists.txt | 9 + tools/mmq-tune/mmq-tune.cpp | 303 +++++++++++++++++++++++ tools/mmq-tune/sweep.py | 162 +++++++++++++ 10 files changed, 956 insertions(+), 4 deletions(-) create mode 100644 ggml/include/ggml-mmq-tune.h create mode 100644 ggml/src/ggml-cuda/mmq-tune.cu create mode 100644 tools/mmq-tune/CMakeLists.txt create mode 100644 tools/mmq-tune/mmq-tune.cpp create mode 100644 tools/mmq-tune/sweep.py diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt index c4e99c9de654..33401b55a7b2 100644 --- a/ggml/CMakeLists.txt +++ b/ggml/CMakeLists.txt @@ -220,6 +220,9 @@ option(GGML_HIP_ROCWMMA_FATTN "ggml: enable rocWMMA for FlashAtten option(GGML_HIP_MMQ_MFMA "ggml: enable MFMA MMA for CDNA in MMQ" ON) option(GGML_HIP_EXPORT_METRICS "ggml: enable kernel perf metrics output" OFF) option(GGML_HIP_ROOFLINE "ggml: enable per-op roofline profiling" OFF) +option(GGML_MMQ_TUNE "ggml: build the MMQ autotuning core" OFF) +set (GGML_HIP_MMQ_Y "" CACHE STRING "ggml: override MMQ mmq_y on RDNA3.5 (default 64)") +set (GGML_HIP_MMQ_NWARPS "" CACHE STRING "ggml: override MMQ nwarps on RDNA3.5 (default 4)") option(GGML_MUSA_GRAPHS "ggml: use MUSA graph, experimental, unstable" OFF) option(GGML_MUSA_MUDNN_COPY "ggml: enable muDNN for accelerated copy" OFF) option(GGML_VULKAN "ggml: use Vulkan" OFF) diff --git a/ggml/include/ggml-mmq-tune.h b/ggml/include/ggml-mmq-tune.h new file mode 100644 index 000000000000..1ccf3b26ba21 --- /dev/null +++ b/ggml/include/ggml-mmq-tune.h @@ -0,0 +1,55 @@ +#pragma once + +// Autotuning entry point for the MMQ (mul_mat_q) kernel, built only with -DGGML_MMQ_TUNE=ON. +// +// Launches mul_mat_q with a caller-chosen tile width instead of the one mul_mat_q_case would +// pick, for dense and MoE shapes alike. mmq_y and the warp count are compile-time constants +// (GGML_CUDA_MMQ_Y_RDNA3_5 / GGML_CUDA_MMQ_NWARPS_RDNA3_5 in mmq.cuh); they are reported back so +// results from builds with different values can be merged. + +#include "ggml.h" +#include "ggml-backend.h" + +#include +#include + +#ifdef __cplusplus +extern "C" { +#endif + +// mmq_x value meaning "let the production heuristic choose" - used as the baseline measurement. +#define GGML_MMQ_TUNE_MMQ_X_AUTO (-1) + +struct ggml_mmq_tune_case { + enum ggml_type type; // weight type of src0 + int64_t n; // src0 rows (output columns) + int64_t k; // reduction dimension + int64_t m; // tokens + int n_experts; // 0 = dense MUL_MAT, >0 = MoE MUL_MAT_ID + int top_k; // experts per token (MoE only) + float zipf_s; // routing skew (MoE only); 0 = uniform + int nwarmup; + int niter; +}; + +struct ggml_mmq_tune_point { + int mmq_x; + bool valid; // false when the config is rejected (granularity or shared memory) + double us_median; + double us_min; + double checksum; // position-weighted sum of |dst|; configs that disagree computed different results +}; + +// Runs one case over the given mmq_x list, writing n_mmq_x results into out. +// Buffers are allocated once for the whole list. Returns false if the case itself is unusable. +GGML_BACKEND_API bool ggml_mmq_tune_sweep( + const struct ggml_mmq_tune_case * tcase, + const int * mmq_x_list, int n_mmq_x, + struct ggml_mmq_tune_point * out); + +GGML_BACKEND_API int ggml_mmq_tune_mmq_y(void); +GGML_BACKEND_API int ggml_mmq_tune_nwarps(void); + +#ifdef __cplusplus +} +#endif diff --git a/ggml/src/ggml-cuda/CMakeLists.txt b/ggml/src/ggml-cuda/CMakeLists.txt index d3953eee962e..4294699b9239 100644 --- a/ggml/src/ggml-cuda/CMakeLists.txt +++ b/ggml/src/ggml-cuda/CMakeLists.txt @@ -103,6 +103,11 @@ if (CUDAToolkit_FOUND) list(APPEND GGML_HEADERS_CUDA "../../include/ggml-cuda.h") file(GLOB GGML_SOURCES_CUDA "*.cu") + # The autotuning core instantiates launch_mul_mat_q for every mmq_x; keep it out of normal builds. + list(FILTER GGML_SOURCES_CUDA EXCLUDE REGEX "mmq-tune\\.cu$") + if (GGML_MMQ_TUNE) + list(APPEND GGML_SOURCES_CUDA mmq-tune.cu) + endif() file(GLOB SRCS "template-instances/fattn-tile*.cu") list(APPEND GGML_SOURCES_CUDA ${SRCS}) file(GLOB SRCS "template-instances/fattn-mma*.cu") diff --git a/ggml/src/ggml-cuda/mmq-tune.cu b/ggml/src/ggml-cuda/mmq-tune.cu new file mode 100644 index 000000000000..280547a224a8 --- /dev/null +++ b/ggml/src/ggml-cuda/mmq-tune.cu @@ -0,0 +1,386 @@ +// MMQ autotuning core - compiled only with -DGGML_MMQ_TUNE=ON. +// +// Mirrors both branches of ggml_cuda_mul_mat_q in mmq.cu (dense and MoE) but dispatches to +// launch_mul_mat_q with an explicit mmq_x, so every tile width can be measured for +// a given shape instead of only the one mul_mat_q_case would pick. + +#include "ggml-mmq-tune.h" + +#include "common.cuh" +#include "mmid.cuh" +#include "mmq.cuh" +#include "quantize.cuh" + +#include +#include +#include +#include +#include + +#if defined(GGML_USE_HIP) +// vendors/hip.h aliases most of the event API but not these two. +#define cudaEventCreate hipEventCreate +#define cudaEventElapsedTime hipEventElapsedTime +#endif + +int ggml_mmq_tune_mmq_y(void) { + return get_mmq_y_host(ggml_cuda_info().devices[ggml_cuda_get_device()].cc); +} + +int ggml_mmq_tune_nwarps(void) { + const int id = ggml_cuda_get_device(); + return mmq_get_nwarps_host(ggml_cuda_info().devices[id].cc, ggml_cuda_info().devices[id].warp_size); +} + +#define MMQ_TUNE_RED_BLOCKS 64 + +// Position-weighted checksum of dst, reduced per block and summed on the host. Configs that +// disagree on it computed something different; a majority vote over the widths then says which +// ones are wrong. The index weight is what makes a misplaced row visible - a plain sum of |dst| +// is permutation invariant. +static __global__ void mmq_tune_checksum(const float * dst, const int64_t n, double * out) { + __shared__ double s[256]; + double local = 0.0; + for (int64_t i = blockIdx.x*blockDim.x + threadIdx.x; i < n; i += (int64_t) gridDim.x*blockDim.x) { + local += fabs((double) dst[i]) * (double) (i % 251 + 1); + } + s[threadIdx.x] = local; + __syncthreads(); + for (int stride = blockDim.x/2; stride > 0; stride /= 2) { + if (threadIdx.x < stride) { + s[threadIdx.x] += s[threadIdx.x + stride]; + } + __syncthreads(); + } + if (threadIdx.x == 0) { + out[blockIdx.x] = s[0]; + } +} + +// Draws top_k distinct experts per token from a Zipf distribution over a shuffled expert order, +// so the hot experts are not always the low indices. Measured routing in MoE models follows +// Zipf with s around 2. +static std::vector mmq_tune_build_ids(int n_tokens, int top_k, int n_experts, float zipf_s) { + std::mt19937 rng(5678); + + std::vector order(n_experts); + for (int e = 0; e < n_experts; ++e) { + order[e] = e; + } + std::shuffle(order.begin(), order.end(), rng); + + std::vector cdf(n_experts); + double acc = 0.0; + for (int r = 0; r < n_experts; ++r) { + acc += zipf_s > 0.0f ? 1.0/pow(r + 1.0, (double) zipf_s) : 1.0; + cdf[r] = acc; + } + for (double & c : cdf) { + c /= acc; + } + + std::uniform_real_distribution uni(0.0, 1.0); + std::vector ids((size_t) top_k*n_tokens); + + for (int t = 0; t < n_tokens; ++t) { + for (int j = 0; j < top_k; ++j) { + int e; + bool dup; + int guard = 0; + do { + const int rank = std::lower_bound(cdf.begin(), cdf.end(), uni(rng)) - cdf.begin(); + e = order[std::min(rank, n_experts - 1)]; + dup = false; + for (int p = 0; p < j; ++p) { + dup |= ids[(size_t) t*top_k + p] == e; + } + // A heavily skewed draw collides often; fall back to a linear probe rather than + // spinning, which would flatten the distribution we are trying to reproduce. + if (dup && ++guard > 64) { + for (int cand = 0; cand < n_experts; ++cand) { + bool used = false; + for (int p = 0; p < j; ++p) { + used |= ids[(size_t) t*top_k + p] == cand; + } + if (!used) { + e = cand; + dup = false; + break; + } + } + } + } while (dup); + ids[(size_t) t*top_k + j] = e; + } + } + return ids; +} + +template +static bool mmq_tune_launch(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream, const int mmq_x) { + switch (mmq_x) { + case 8: launch_mul_mat_q(ctx, args, stream); return true; + case 16: launch_mul_mat_q(ctx, args, stream); return true; + case 24: launch_mul_mat_q(ctx, args, stream); return true; + case 32: launch_mul_mat_q(ctx, args, stream); return true; + case 40: launch_mul_mat_q(ctx, args, stream); return true; + case 48: launch_mul_mat_q(ctx, args, stream); return true; + case 56: launch_mul_mat_q(ctx, args, stream); return true; + case 64: launch_mul_mat_q(ctx, args, stream); return true; + case 72: launch_mul_mat_q(ctx, args, stream); return true; + case 80: launch_mul_mat_q(ctx, args, stream); return true; + case 88: launch_mul_mat_q(ctx, args, stream); return true; + case 96: launch_mul_mat_q(ctx, args, stream); return true; + case 104: launch_mul_mat_q(ctx, args, stream); return true; + case 112: launch_mul_mat_q(ctx, args, stream); return true; + case 120: launch_mul_mat_q(ctx, args, stream); return true; + case 128: launch_mul_mat_q(ctx, args, stream); return true; + case GGML_MMQ_TUNE_MMQ_X_AUTO: mul_mat_q_case(ctx, args, stream); return true; + default: return false; + } +} + +// Same three rejection criteria the heuristic loop in mul_mat_q_case applies. +template +static bool mmq_tune_valid(const int mmq_x, const int cc, const int warp_size, const int nwarps, + const int mmq_y, const size_t smpbo) { + if (mmq_x == GGML_MMQ_TUNE_MMQ_X_AUTO) { + return true; + } + if (mmq_x < 8 || mmq_x % 8 != 0 || mmq_x > get_mmq_x_max_host(cc)) { + return false; + } + if (mmq_x % mmq_get_granularity_host(mmq_x, cc) != 0) { + return false; + } + return mmq_get_nbytes_shared(mmq_x, mmq_y, cc, warp_size, nwarps) <= smpbo; +} + +// Distinct weights for every expert, generated and uploaded one [K,N] slice at a time: building +// the whole f32 tensor at once would need gigabytes of host memory for the larger MoE shapes, +// and reusing a few slices would shrink the weight working set enough to make the measurement +// unrealistically cache-friendly. +template +static void mmq_tune_fill_weights(char * dst_d, int64_t k, int64_t n, int n_experts, cudaStream_t stream) { + const size_t slice_bytes = ggml_row_size(type, k)*n; + + std::mt19937 rng(1234); + std::uniform_real_distribution dist(-1.0f, 1.0f); + + std::vector f32(k*n); + std::vector q(slice_bytes); + + for (int e = 0; e < n_experts; ++e) { + for (float & v : f32) { + v = dist(rng); + } + ggml_quantize_chunk(type, f32.data(), q.data(), 0, n, k, nullptr); + CUDA_CHECK(cudaMemcpyAsync(dst_d + slice_bytes*e, q.data(), slice_bytes, + cudaMemcpyHostToDevice, stream)); + CUDA_CHECK(cudaStreamSynchronize(stream)); // q is reused for the next expert + } +} + +template +static bool mmq_tune_sweep_impl(const ggml_mmq_tune_case & tc, const int * mmq_x_list, const int n_mmq_x, + ggml_mmq_tune_point * out) { + const bool moe = tc.n_experts > 0; + const int64_t ne00 = tc.k; + const int64_t ne01 = tc.n; + const int64_t ne02 = moe ? tc.n_experts : 1; + const int64_t top_k = moe ? tc.top_k : 1; + // Dense: src1 is [K, m]. MoE: src1 is [K, top_k, m] and dst is [N, top_k, m]. + const int64_t ne11 = moe ? top_k : tc.m; + const int64_t ne12 = moe ? tc.m : 1; + + if (ne00 % ggml_blck_size(type) != 0 || (moe && (tc.top_k < 1 || tc.top_k > tc.n_experts))) { + return false; + } + + const int id = ggml_cuda_get_device(); + const int cc = ggml_cuda_info().devices[id].cc; + const int warp_size = ggml_cuda_info().devices[id].warp_size; + const size_t smpbo = ggml_cuda_info().devices[id].smpbo; + const int nwarps = mmq_get_nwarps_host(cc, warp_size); + const int mmq_y = get_mmq_y_host(cc); + + ggml_backend_cuda_context ctx(id); + cudaStream_t stream = ctx.stream(); + + const size_t x_row_size = ggml_row_size(type, ne00); + const int64_t n_dst = ne01*ne11*ne12; + + ggml_cuda_pool_alloc src0_d(ctx.pool(), x_row_size*ne01*ne02); + ggml_cuda_pool_alloc src1_d(ctx.pool(), ne00*ne11*ne12); + ggml_cuda_pool_alloc dst_d(ctx.pool(), n_dst); + ggml_cuda_pool_alloc checksum_d(ctx.pool(), MMQ_TUNE_RED_BLOCKS); + + mmq_tune_fill_weights(src0_d.get(), ne00, ne01, ne02, stream); + + { + std::mt19937 rng(4321); + std::uniform_real_distribution dist(-1.0f, 1.0f); + std::vector y(ne00*ne11*ne12); + for (float & v : y) { + v = dist(rng); + } + CUDA_CHECK(cudaMemcpyAsync(src1_d.get(), y.data(), y.size()*sizeof(float), cudaMemcpyHostToDevice, stream)); + CUDA_CHECK(cudaStreamSynchronize(stream)); + } + + const int64_t ne10_padded = GGML_PAD(ne00, MATRIX_ROW_PADDING); + const int64_t s01 = x_row_size / ggml_type_size(type); + const int64_t s02 = ne01*s01; + const int64_t s1 = ne01; + + const bool use_stream_k = (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA) + || GGML_CUDA_CC_IS_CDNA(cc); + + // Kept alive for the whole sweep; mmq_args holds raw pointers into them. + ggml_cuda_pool_alloc ids_dst(ctx.pool()); + ggml_cuda_pool_alloc ids_src1(ctx.pool()); + ggml_cuda_pool_alloc expert_bounds(ctx.pool()); + ggml_cuda_pool_alloc ids_d(ctx.pool()); + ggml_cuda_pool_alloc src1_q8_1(ctx.pool()); + + mmq_args args = {}; + + if (!moe) { + const size_t nbytes = ne11*ne10_padded*sizeof(block_q8_1)/QK8_1 + + get_mmq_x_max_host(cc)*sizeof(block_q8_1_mmq); + src1_q8_1.alloc(nbytes); + + quantize_mmq_q8_1_cuda(src1_d.get(), nullptr, src1_q8_1.get(), type, + ne00, ne00, ne00*ne11, ne00*ne11, ne10_padded, ne11, 1, 1, stream); + CUDA_CHECK(cudaGetLastError()); + + const int64_t s12 = ne11*ne10_padded*sizeof(block_q8_1) / (QK8_1*sizeof(int)); + + args = { + src0_d.get(), type, (const int *) src1_q8_1.get(), nullptr, nullptr, dst_d.get(), + ne00, ne01, ne11, s01, ne11, s1, + 1, 1, s02, s12, ne01*ne11, + 1, 1, s02, s12, ne01*ne11, + use_stream_k, ne11}; + } else { + const int64_t n_expert_used = top_k; + const int64_t ne_get_rows = ne12*n_expert_used; + + const std::vector ids_h = mmq_tune_build_ids(ne12, n_expert_used, ne02, tc.zipf_s); + ids_d.alloc(ids_h.size()); + CUDA_CHECK(cudaMemcpyAsync(ids_d.get(), ids_h.data(), ids_h.size()*sizeof(int32_t), + cudaMemcpyHostToDevice, stream)); + + ids_src1.alloc(ne_get_rows); + ids_dst.alloc(ne_get_rows); + expert_bounds.alloc(ne02 + 1); + + // ids is [top_k, m] contiguous, src1 is [K, top_k, m]: si1 = sis1 = top_k. + ggml_cuda_launch_mm_ids_helper(ids_d.get(), ids_src1.get(), ids_dst.get(), expert_bounds.get(), + ne02, ne12, n_expert_used, ne11, n_expert_used, n_expert_used, stream); + CUDA_CHECK(cudaGetLastError()); + + const size_t nbytes = ne12*n_expert_used*ne10_padded*sizeof(block_q8_1)/QK8_1 + + get_mmq_x_max_host(cc)*sizeof(block_q8_1_mmq); + src1_q8_1.alloc(nbytes); + + quantize_mmq_q8_1_cuda(src1_d.get(), ids_src1.get(), src1_q8_1.get(), type, + ne00, ne00, ne00*ne11, ne00*ne11*ne12, ne10_padded, + ne12*n_expert_used, 1, 1, stream); + CUDA_CHECK(cudaGetLastError()); + + const int64_t s12 = ne11*ne10_padded*sizeof(block_q8_1) / (QK8_1*sizeof(int)); + + args = { + src0_d.get(), type, (const int *) src1_q8_1.get(), ids_dst.get(), expert_bounds.get(), dst_d.get(), + ne00, ne01, ne_get_rows, s01, ne_get_rows, s1, + ne02, ne02, s02, s12, ne01*ne11, + 1, 1, ne02*s02, ne12*s12, ne01*ne11*ne12, + use_stream_k, ne12}; + } + + // ---- sweep ------------------------------------------------------------ + cudaEvent_t ev_start, ev_stop; + CUDA_CHECK(cudaEventCreate(&ev_start)); + CUDA_CHECK(cudaEventCreate(&ev_stop)); + + std::vector> samples(n_mmq_x); + + for (int i = 0; i < n_mmq_x; ++i) { + const int mmq_x = mmq_x_list[i]; + out[i] = { mmq_x, mmq_tune_valid(mmq_x, cc, warp_size, nwarps, mmq_y, smpbo), 0.0, 0.0, 0.0 }; + + if (!out[i].valid) { + continue; + } + samples[i].reserve(tc.niter); + for (int w = 0; w < tc.nwarmup; ++w) { + mmq_tune_launch(ctx, args, stream, mmq_x); + } + } + CUDA_CHECK(cudaStreamSynchronize(stream)); + CUDA_CHECK(cudaGetLastError()); + + // One timed iteration per config per pass, so clock drift and background load spread across + // all configs instead of penalising whichever ones happen to run late. + for (int pass = 0; pass < tc.niter; ++pass) { + for (int i = 0; i < n_mmq_x; ++i) { + if (!out[i].valid) { + continue; + } + CUDA_CHECK(cudaEventRecord(ev_start, stream)); + mmq_tune_launch(ctx, args, stream, mmq_x_list[i]); + CUDA_CHECK(cudaEventRecord(ev_stop, stream)); + CUDA_CHECK(cudaEventSynchronize(ev_stop)); + float ms = 0.0f; + CUDA_CHECK(cudaEventElapsedTime(&ms, ev_start, ev_stop)); + samples[i].push_back(ms*1000.0); + } + } + CUDA_CHECK(cudaGetLastError()); + + for (int i = 0; i < n_mmq_x; ++i) { + if (!out[i].valid) { + continue; + } + + // Re-run alone so dst holds only this config's output. MoE writes just the routed rows, + // so clear it first or rows left by the previous config would hide a mismatch. + CUDA_CHECK(cudaMemsetAsync(dst_d.get(), 0, n_dst*sizeof(float), stream)); + mmq_tune_launch(ctx, args, stream, mmq_x_list[i]); + mmq_tune_checksum<<>>(dst_d.get(), n_dst, checksum_d.get()); + + double partial[MMQ_TUNE_RED_BLOCKS]; + CUDA_CHECK(cudaMemcpyAsync(partial, checksum_d.get(), sizeof(partial), cudaMemcpyDeviceToHost, stream)); + CUDA_CHECK(cudaStreamSynchronize(stream)); + + std::sort(samples[i].begin(), samples[i].end()); + out[i].us_median = samples[i][tc.niter/2]; + out[i].us_min = samples[i][0]; + out[i].checksum = std::accumulate(partial, partial + MMQ_TUNE_RED_BLOCKS, 0.0); + } + + CUDA_CHECK(cudaEventDestroy(ev_start)); + CUDA_CHECK(cudaEventDestroy(ev_stop)); + + return true; +} + +bool ggml_mmq_tune_sweep(const ggml_mmq_tune_case * tcase, const int * mmq_x_list, const int n_mmq_x, + ggml_mmq_tune_point * out) { + if (tcase->niter < 1 || n_mmq_x < 1) { + return false; + } + + switch (tcase->type) { + case GGML_TYPE_Q8_0: return mmq_tune_sweep_impl(*tcase, mmq_x_list, n_mmq_x, out); + case GGML_TYPE_Q5_0: return mmq_tune_sweep_impl(*tcase, mmq_x_list, n_mmq_x, out); + case GGML_TYPE_Q5_1: return mmq_tune_sweep_impl(*tcase, mmq_x_list, n_mmq_x, out); + case GGML_TYPE_Q4_K: return mmq_tune_sweep_impl(*tcase, mmq_x_list, n_mmq_x, out); + case GGML_TYPE_Q5_K: return mmq_tune_sweep_impl(*tcase, mmq_x_list, n_mmq_x, out); + case GGML_TYPE_Q6_K: return mmq_tune_sweep_impl(*tcase, mmq_x_list, n_mmq_x, out); + default: + fprintf(stderr, "%s: unsupported type %s\n", __func__, ggml_type_name(tcase->type)); + return false; + } +} diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 924d8195204d..a0812f492233 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -15,6 +15,15 @@ using namespace ggml_cuda_mma; #define MMQ_ITER_K_FP4 512 #define MMQ_NWARPS 8 +// Tile height and warp count for RDNA3.5. get_mmq_y_* and mmq_get_nwarps_* must return the same +// values on the host and the device or the shared memory layout breaks, so both read these. +#ifndef GGML_CUDA_MMQ_Y_RDNA3_5 +#define GGML_CUDA_MMQ_Y_RDNA3_5 64 +#endif +#ifndef GGML_CUDA_MMQ_NWARPS_RDNA3_5 +#define GGML_CUDA_MMQ_NWARPS_RDNA3_5 4 +#endif + typedef void (*load_tiles_mmq_t)(const char * __restrict__ x, int * x_tile, const int kbx0, const int i_max, const int stride); typedef void (*vec_dot_mmq_t)(const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00); typedef void (*mmq_write_back_t)(const float * __restrict__ sum, const int32_t * __restrict__ get_rows_to_sorted, @@ -142,7 +151,7 @@ static constexpr __device__ int get_mmq_x_max_device() { static int get_mmq_y_host(const int cc) { if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { - return 64; + return GGML_CUDA_MMQ_Y_RDNA3_5; } return GGML_CUDA_CC_IS_AMD(cc) ? (GGML_CUDA_CC_IS_RDNA1(cc) ? 64 : 128) : ((GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA) ? 128 : 64); @@ -160,7 +169,7 @@ if (type == GGML_TYPE_NVFP4 || type == GGML_TYPE_MXFP4) { static constexpr __device__ int get_mmq_y_device() { #if defined(GGML_USE_HIP) #if defined(RDNA3_5) - return 64; + return GGML_CUDA_MMQ_Y_RDNA3_5; #elif defined(RDNA1) return 64; #else @@ -325,7 +334,7 @@ static constexpr __device__ int mmq_get_granularity_device(const int /*mmq_x*/) #if defined(GGML_USE_HIP) static int mmq_get_nwarps_host(const int cc, const int warp_size) { if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { - return 4; + return GGML_CUDA_MMQ_NWARPS_RDNA3_5; } return amd_mfma_available(cc) ? 8 : 256/warp_size; } @@ -337,7 +346,7 @@ static int mmq_get_nwarps_host(const int /*cc*/, const int warp_size) { static constexpr __device__ int mmq_get_nwarps_device() { #if defined(RDNA3_5) - return 4; + return GGML_CUDA_MMQ_NWARPS_RDNA3_5; #elif defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) return 8; #else diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt index 5eb581b66b25..9bbaedb95f04 100644 --- a/ggml/src/ggml-hip/CMakeLists.txt +++ b/ggml/src/ggml-hip/CMakeLists.txt @@ -61,6 +61,11 @@ file(GLOB GGML_HEADERS_ROCM "../ggml-cuda/*.cuh") list(APPEND GGML_HEADERS_ROCM "../../include/ggml-cuda.h") file(GLOB GGML_SOURCES_ROCM "../ggml-cuda/*.cu") +# The autotuning core instantiates launch_mul_mat_q for every mmq_x; keep it out of normal builds. +list(FILTER GGML_SOURCES_ROCM EXCLUDE REGEX "mmq-tune\\.cu$") +if (GGML_MMQ_TUNE) + list(APPEND GGML_SOURCES_ROCM ../ggml-cuda/mmq-tune.cu) +endif() file(GLOB SRCS "../ggml-cuda/template-instances/fattn-tile*.cu") list(APPEND GGML_SOURCES_ROCM ${SRCS}) file(GLOB SRCS "../ggml-cuda/template-instances/fattn-mma*.cu") @@ -175,4 +180,16 @@ if (GGML_HIP_ROOFLINE) target_link_libraries(ggml-hip PRIVATE ${CMAKE_DL_LIBS}) endif() +if (GGML_MMQ_TUNE) + target_compile_definitions(ggml-hip PRIVATE GGML_MMQ_TUNE) +endif() + +# MMQ tile hyperparameters for RDNA3.5, swept by tools/mmq-tune/sweep.py. +if (GGML_HIP_MMQ_Y) + target_compile_definitions(ggml-hip PRIVATE GGML_CUDA_MMQ_Y_RDNA3_5=${GGML_HIP_MMQ_Y}) +endif() +if (GGML_HIP_MMQ_NWARPS) + target_compile_definitions(ggml-hip PRIVATE GGML_CUDA_MMQ_NWARPS_RDNA3_5=${GGML_HIP_MMQ_NWARPS}) +endif() + target_compile_options(ggml-hip PRIVATE "$<$:-ffast-math;-fno-finite-math-only>") diff --git a/tools/CMakeLists.txt b/tools/CMakeLists.txt index 780df3266132..cb9dffce55f7 100644 --- a/tools/CMakeLists.txt +++ b/tools/CMakeLists.txt @@ -40,4 +40,7 @@ else() endif() add_subdirectory(fit-params) add_subdirectory(results) + if (GGML_MMQ_TUNE) + add_subdirectory(mmq-tune) + endif() endif() diff --git a/tools/mmq-tune/CMakeLists.txt b/tools/mmq-tune/CMakeLists.txt new file mode 100644 index 000000000000..fe2b300e5653 --- /dev/null +++ b/tools/mmq-tune/CMakeLists.txt @@ -0,0 +1,9 @@ +set(TARGET mmq-tune) + +add_executable(${TARGET} mmq-tune.cpp) +target_link_libraries(${TARGET} PRIVATE ggml ${CMAKE_THREAD_LIBS_INIT}) +target_compile_features(${TARGET} PRIVATE cxx_std_17) + +if(LLAMA_TOOLS_INSTALL) + install(TARGETS ${TARGET} RUNTIME) +endif() diff --git a/tools/mmq-tune/mmq-tune.cpp b/tools/mmq-tune/mmq-tune.cpp new file mode 100644 index 000000000000..7f9b882b474f --- /dev/null +++ b/tools/mmq-tune/mmq-tune.cpp @@ -0,0 +1,303 @@ +// Autotuner for the MMQ (mul_mat_q) tile width. The catalog holds every distinct shape that a +// set of 20 models dispatches to mul_mat_q, as captured from GGML_ROOFLINE_OUT traces. +// +// `weight` is the invocation count per forward pass, so aggregates reflect real GPU time rather +// than treating a shape run 10 times as equal to one run 900 times. Counts come from the +// -ub 2048 capture and are reused as a proxy at the other token counts the sweep covers. +// +// The token count is a sweep parameter, not part of the catalog, and it matters: the dense +// q8_0 mmq_x=64 cliff only shows up around 64 tokens, while the MoE tile-width gap needs 2048. +// +// mmq_x is swept at runtime. mmq_y and the warp count are compile-time constants, so sweeping +// those means rebuilding - see sweep.py. + +#include "ggml.h" +#include "ggml-mmq-tune.h" + +#include +#include +#include +#include +#include +#include + +struct shape { + const char * model; + ggml_type type; + int64_t n; + int64_t k; + int n_experts; // 0 = dense + int top_k; + int weight; // invocations per pp128 run +}; + +static const shape g_shapes[] = { + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q4_K, 17408, 5120, 0, 0, 128 }, // 1853.7 ms + { "gemma-3-12b-it_Q4_", GGML_TYPE_Q4_K, 15360, 3840, 0, 0, 188 }, // 1736.0 ms + { "Qwen2.5-7B-Instruc", GGML_TYPE_Q4_K, 18944, 3584, 0, 0, 54 }, // 603.9 ms + { "Llama-3.1-8B-Instr", GGML_TYPE_Q4_K, 14336, 4096, 0, 0, 62 }, // 555.2 ms + { "Qwen3-8B_Q4_K_M", GGML_TYPE_Q4_K, 12288, 4096, 0, 0, 70 }, // 542.7 ms + { "Qwen2.5-3B-Instruc", GGML_TYPE_Q4_K, 11008, 2048, 0, 0, 140 }, // 479.7 ms + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q4_K, 5120, 17408, 0, 0, 32 }, // 459.2 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q4_K, 512, 2048, 256, 8, 160 }, // 452.0 ms + { "gemma-3-12b-it_Q4_", GGML_TYPE_Q4_K, 3840, 15360, 0, 0, 48 }, // 424.5 ms + { "Llama-3.1-8B-Instr", GGML_TYPE_Q4_K, 4096, 4096, 0, 0, 136 }, // 337.8 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q4_K, 1536, 2048, 64, 4, 90 }, // 277.2 ms + { "gemma-3-4b-it_Q4_K", GGML_TYPE_Q4_K, 10240, 2560, 0, 0, 66 }, // 262.8 ms + { "Qwen3-4B_Q4_K_M", GGML_TYPE_Q4_K, 9728, 2560, 0, 0, 70 }, // 258.6 ms + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q5_K, 5120, 6144, 0, 0, 48 }, // 254.3 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q5_K, 2048, 512, 256, 8, 77 }, // 250.0 ms + { "gemma-3-12b-it_Q4_", GGML_TYPE_Q4_K, 4096, 3840, 0, 0, 96 }, // 242.1 ms + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q4_K, 6144, 5120, 0, 0, 48 }, // 241.7 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q4_K, 1408, 2816, 128, 8, 30 }, // 232.6 ms + { "gemma-3-12b-it_Q4_", GGML_TYPE_Q4_K, 3840, 4096, 0, 0, 96 }, // 221.1 ms + { "Qwen2.5-3B-Instruc", GGML_TYPE_Q4_K, 2048, 2048, 0, 0, 320 }, // 207.2 ms + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q4_K, 10240, 5120, 0, 0, 24 }, // 206.1 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q5_1, 2816, 704, 128, 8, 29 }, // 195.4 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q8_0, 8192, 2048, 0, 0, 80 }, // 186.7 ms + { "gemma-3-12b-it_Q4_", GGML_TYPE_Q4_K, 2048, 3840, 0, 0, 144 }, // 176.3 ms + { "gemma-2b-it_Q4_K_M", GGML_TYPE_Q4_K, 16384, 2048, 0, 0, 34 }, // 172.6 ms + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q4_K, 12288, 5120, 0, 0, 16 }, // 165.1 ms + { "Qwen2.5-7B-Instruc", GGML_TYPE_Q4_K, 3584, 18944, 0, 0, 14 }, // 155.8 ms + { "Llama-3.1-8B-Instr", GGML_TYPE_Q4_K, 4096, 14336, 0, 0, 16 }, // 142.6 ms + { "Qwen3-8B_Q4_K_M", GGML_TYPE_Q4_K, 4096, 12288, 0, 0, 18 }, // 140.2 ms + { "Qwen2.5-7B-Instruc", GGML_TYPE_Q4_K, 3584, 3584, 0, 0, 56 }, // 129.8 ms + { "Qwen2.5-3B-Instruc", GGML_TYPE_Q4_K, 2048, 11008, 0, 0, 36 }, // 125.2 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 12288, 1536, 0, 0, 40 }, // 115.6 ms + { "SmolLM2-1.7B-Instr", GGML_TYPE_Q4_K, 8192, 2048, 0, 0, 46 }, // 111.2 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q6_K, 2048, 1536, 64, 4, 21 }, // 109.6 ms + { "Qwen3-1.7B_Q4_K_M", GGML_TYPE_Q4_K, 6144, 2048, 0, 0, 54 }, // 97.8 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q8_0, 2048, 4096, 0, 0, 80 }, // 94.0 ms + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q4_K, 5120, 6144, 0, 0, 16 }, // 77.3 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q4_K, 2048, 1536, 64, 4, 24 }, // 75.5 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q8_0, 4096, 2048, 0, 0, 60 }, // 75.0 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q4_K, 2048, 5120, 0, 0, 47 }, // 71.3 ms + { "Qwen3-4B_Q4_K_M", GGML_TYPE_Q4_K, 2560, 9728, 0, 0, 18 }, // 69.4 ms + { "gemma-3-4b-it_Q4_K", GGML_TYPE_Q4_K, 2560, 10240, 0, 0, 17 }, // 68.3 ms + { "Llama-3.1-8B-Instr", GGML_TYPE_Q4_K, 1024, 4096, 0, 0, 102 }, // 64.9 ms + { "Qwen3-4B_Q4_K_M", GGML_TYPE_Q4_K, 4096, 2560, 0, 0, 36 }, // 60.6 ms + { "Qwen3-4B_Q4_K_M", GGML_TYPE_Q4_K, 2560, 4096, 0, 0, 36 }, // 55.0 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 2112, 2816, 0, 0, 60 }, // 48.3 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q5_K, 1536, 2048, 0, 0, 90 }, // 48.3 ms + { "Qwen2.5-0.5B-Instr", GGML_TYPE_Q5_0, 4864, 896, 0, 0, 46 }, // 45.2 ms + { "gemma-2b-it_Q4_K_M", GGML_TYPE_Q4_K, 2048, 16384, 0, 0, 9 }, // 44.4 ms + { "Qwen3-4B_Q4_K_M", GGML_TYPE_Q4_K, 1024, 2560, 0, 0, 105 }, // 43.6 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 6144, 1536, 0, 0, 30 }, // 43.2 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 2048, 2816, 0, 0, 50 }, // 41.9 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 4096, 2816, 0, 0, 25 }, // 41.2 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 2816, 4096, 0, 0, 25 }, // 39.7 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q8_0, 512, 2048, 0, 0, 200 }, // 32.7 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q4_K, 5120, 768, 0, 0, 47 }, // 30.6 ms + { "SmolLM2-1.7B-Instr", GGML_TYPE_Q4_K, 2048, 8192, 0, 0, 12 }, // 30.0 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q4_K, 3584, 1024, 0, 0, 48 }, // 28.8 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 1536, 12288, 0, 0, 10 }, // 28.5 ms + { "gemma-3-4b-it_Q4_K", GGML_TYPE_Q4_K, 2048, 2560, 0, 0, 34 }, // 28.3 ms + { "gemma-3-4b-it_Q4_K", GGML_TYPE_Q4_K, 2560, 2048, 0, 0, 34 }, // 27.8 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 2816, 2112, 0, 0, 30 }, // 26.9 ms + { "Qwen3-1.7B_Q4_K_M", GGML_TYPE_Q4_K, 2048, 6144, 0, 0, 14 }, // 26.3 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q8_0, 512, 192, 0, 0, 47 }, // 23.9 ms + { "Qwen3.6-27B_Q4_K_M", GGML_TYPE_Q4_K, 1024, 5120, 0, 0, 24 }, // 20.6 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q5_K, 6144, 1024, 0, 0, 18 }, // 18.7 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q8_0, 256, 512, 0, 0, 47 }, // 18.5 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 8192, 2816, 0, 0, 5 }, // 16.4 ms + { "Qwen3-1.7B_Q4_K_M", GGML_TYPE_Q4_K, 1024, 2048, 0, 0, 48 }, // 15.7 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 2816, 8192, 0, 0, 5 }, // 15.1 ms + { "Qwen2.5-3B-Instruc", GGML_TYPE_Q4_K, 256, 2048, 0, 0, 135 }, // 15.0 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 1536, 2048, 0, 0, 28 }, // 14.0 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 2048, 1536, 0, 0, 28 }, // 13.9 ms + { "Qwen2.5-7B-Instruc", GGML_TYPE_Q4_K, 512, 3584, 0, 0, 42 }, // 13.6 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q8_0, 2048, 512, 0, 0, 80 }, // 13.0 ms + { "Qwen3.6-35B-A3B_UD", GGML_TYPE_Q6_K, 2048, 512, 256, 8, 3 }, // 12.8 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q4_K, 768, 2048, 0, 0, 47 }, // 12.0 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 1536, 6144, 0, 0, 8 }, // 11.8 ms + { "Qwen2.5-0.5B-Instr", GGML_TYPE_Q5_0, 896, 896, 0, 0, 48 }, // 9.4 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q8_0, 576, 2048, 0, 0, 47 }, // 9.2 ms + { "Qwen2.5-0.5B-Instr", GGML_TYPE_Q4_K, 896, 4864, 0, 0, 11 }, // 8.7 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 4096, 1536, 0, 0, 7 }, // 7.0 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q4_K, 1024, 3584, 0, 0, 12 }, // 6.9 ms + { "GLM-4.7-Flash_Q4_K", GGML_TYPE_Q4_K, 10240, 2048, 0, 0, 2 }, // 6.7 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q5_K, 1024, 2048, 0, 0, 18 }, // 6.7 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 1536, 4096, 0, 0, 7 }, // 6.5 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q4_K, 2048, 1024, 0, 0, 18 }, // 6.2 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 2816, 704, 128, 8, 1 }, // 4.3 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q4_K, 4096, 1024, 0, 0, 6 }, // 4.1 ms + { "Qwen3.5-35B-A3B_Q4", GGML_TYPE_Q8_0, 32, 2048, 0, 0, 60 }, // 2.5 ms + { "gemma-4-26B-A4B-it", GGML_TYPE_Q8_0, 1024, 2816, 0, 0, 5 }, // 2.2 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 256, 1536, 0, 0, 18 }, // 1.5 ms + { "Qwen2.5-0.5B-Instr", GGML_TYPE_Q5_0, 128, 896, 0, 0, 36 }, // 1.5 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q8_0, 16, 1024, 0, 0, 36 }, // 0.9 ms + { "Qwen3.5-0.8B_Q4_K_", GGML_TYPE_Q4_K, 512, 1024, 0, 0, 8 }, // 0.7 ms + { "gemma-4-E2B-it_Q4_", GGML_TYPE_Q4_K, 512, 1536, 0, 0, 5 }, // 0.7 ms + { "Qwen2.5-0.5B-Instr", GGML_TYPE_Q8_0, 128, 896, 0, 0, 12 }, // 0.4 ms + // MMQ only at smaller batches - at -ub 2048 these two are dispatched to hipBLASLt instead, + // so the capture above does not see them. Kept because the sweep also covers small M. + { "GLM-4.7-Flash", GGML_TYPE_Q6_K, 2048, 1536, 0, 0, 450 }, + { "GLM-4.7-Flash", GGML_TYPE_Q6_K, 2048, 10240, 0, 0, 10 }, +}; + + +static const int g_n_shapes = sizeof(g_shapes)/sizeof(g_shapes[0]); + +static std::vector parse_i64_list(const char * s) { + std::vector out; + for (const char * p = s; *p; ) { + char * end = nullptr; + const int64_t v = strtoll(p, &end, 10); + if (end == p) { + break; + } + out.push_back(v); + p = (*end == ',') ? end + 1 : end; + } + return out; +} + +static void usage(const char * argv0) { + printf("usage: %s [options]\n", argv0); + printf(" -s shape indices to run (default: all)\n"); + printf(" -m token counts to sweep (default: 64,128,512,2048)\n"); + printf(" -x mmq_x values to sweep (default: 8,16,...,128)\n"); + printf(" -r timed iterations per config (default: 11)\n"); + printf(" -w warmup iterations per config (default: 3)\n"); + printf(" --zipf MoE routing skew (default: 2.0, fitted from real traces; 0 = uniform)\n"); + printf(" --dense-only skip MoE shapes\n"); + printf(" --moe-only skip dense shapes\n"); + printf(" --output \n"); + printf(" --list print the shape catalog and exit\n"); +} + +int main(int argc, char ** argv) { + std::vector shape_ids; + std::vector ms = { 64, 128, 512, 2048 }; + std::vector mmq_xs; + int niter = 11; + int nwarmup = 3; + float zipf = 2.0f; + bool csv = false; + bool dense_only = false; + bool moe_only = false; + + for (int i = 1; i < argc; i++) { + const std::string a = argv[i]; + if (a == "-s" && i + 1 < argc) { + shape_ids = parse_i64_list(argv[++i]); + } else if (a == "-m" && i + 1 < argc) { + ms = parse_i64_list(argv[++i]); + } else if (a == "-x" && i + 1 < argc) { + for (int64_t v : parse_i64_list(argv[++i])) { + mmq_xs.push_back((int) v); + } + } else if (a == "-r" && i + 1 < argc) { + niter = atoi(argv[++i]); + } else if (a == "-w" && i + 1 < argc) { + nwarmup = atoi(argv[++i]); + } else if (a == "--zipf" && i + 1 < argc) { + zipf = atof(argv[++i]); + } else if (a == "--dense-only") { + dense_only = true; + } else if (a == "--moe-only") { + moe_only = true; + } else if (a == "--output" && i + 1 < argc) { + csv = std::string(argv[++i]) == "csv"; + } else if (a == "--list") { + for (int s = 0; s < g_n_shapes; s++) { + const shape & sh = g_shapes[s]; + printf("%2d %-12s %-5s %-5s N=%-6" PRId64 " K=%-6" PRId64 " E=%-4d topk=%-3d w=%d\n", + s, sh.model, ggml_type_name(sh.type), sh.n_experts ? "moe" : "dense", + sh.n, sh.k, sh.n_experts, sh.top_k, sh.weight); + } + return 0; + } else { + usage(argv[0]); + return a == "-h" || a == "--help" ? 0 : 1; + } + } + + if (shape_ids.empty()) { + for (int s = 0; s < g_n_shapes; s++) { + const bool is_moe = g_shapes[s].n_experts > 0; + if ((dense_only && is_moe) || (moe_only && !is_moe)) { + continue; + } + shape_ids.push_back(s); + } + } + if (mmq_xs.empty()) { + for (int x = 8; x <= 128; x += 8) { + mmq_xs.push_back(x); + } + } + // The heuristic's choice is the baseline every other config is compared against. + mmq_xs.insert(mmq_xs.begin(), GGML_MMQ_TUNE_MMQ_X_AUTO); + + const int mmq_y = ggml_mmq_tune_mmq_y(); + const int nwarps = ggml_mmq_tune_nwarps(); + + if (csv) { + printf("model,kind,type,n,k,m,n_experts,top_k,zipf_s,weight,mmq_y,nwarps,mmq_x," + "us_median,us_min,tflops,checksum\n"); + } else { + printf("mmq_y=%d nwarps=%d zipf_s=%.2f\n\n", mmq_y, nwarps, zipf); + printf("%-12s %-5s %-5s %7s %7s %6s | %10s | %5s %10s %8s\n", + "model", "kind", "type", "N", "K", "M", "auto us", "best", "us", "speedup"); + } + + std::vector pts(mmq_xs.size()); + + for (int64_t sid : shape_ids) { + if (sid < 0 || sid >= g_n_shapes) { + fprintf(stderr, "bad shape index %" PRId64 "\n", sid); + return 1; + } + const shape & sh = g_shapes[sid]; + const bool is_moe = sh.n_experts > 0; + + for (int64_t m : ms) { + const ggml_mmq_tune_case tc = { + sh.type, sh.n, sh.k, m, sh.n_experts, sh.top_k, is_moe ? zipf : 0.0f, nwarmup, niter + }; + + if (!ggml_mmq_tune_sweep(&tc, mmq_xs.data(), (int) mmq_xs.size(), pts.data())) { + fprintf(stderr, "sweep failed for shape %" PRId64 " m=%" PRId64 "\n", sid, m); + return 1; + } + + // MoE does top_k times the work of a dense GEMM of the same N,K per token. + const double flop = 2.0*m*sh.n*sh.k*(is_moe ? sh.top_k : 1); + + if (csv) { + for (const auto & p : pts) { + if (!p.valid) { + continue; + } + printf("%s,%s,%s,%" PRId64 ",%" PRId64 ",%" PRId64 ",%d,%d,%.2f,%d,%d,%d,%d," + "%.3f,%.3f,%.4f,%.9e\n", + sh.model, is_moe ? "moe" : "dense", ggml_type_name(sh.type), + sh.n, sh.k, m, sh.n_experts, sh.top_k, is_moe ? zipf : 0.0f, sh.weight, + mmq_y, nwarps, p.mmq_x, p.us_median, p.us_min, + flop/(p.us_median*1e6), p.checksum); + } + fflush(stdout); + continue; + } + + const double us_auto = pts[0].us_median; + + int best_x = 0; + double best_us = 0.0; + for (size_t i = 1; i < pts.size(); i++) { + if (pts[i].valid && (best_x == 0 || pts[i].us_median < best_us)) { + best_x = pts[i].mmq_x; + best_us = pts[i].us_median; + } + } + + printf("%-12s %-5s %-5s %7" PRId64 " %7" PRId64 " %6" PRId64 " | %10.2f | %5d %10.2f %7.2fx\n", + sh.model, is_moe ? "moe" : "dense", ggml_type_name(sh.type), sh.n, sh.k, m, + us_auto, best_x, best_us, best_us > 0.0 ? us_auto/best_us : 0.0); + fflush(stdout); + } + } + + return 0; +} diff --git a/tools/mmq-tune/sweep.py b/tools/mmq-tune/sweep.py new file mode 100644 index 000000000000..d9efff918766 --- /dev/null +++ b/tools/mmq-tune/sweep.py @@ -0,0 +1,162 @@ +#!/usr/bin/env python3 +"""Sweep the MMQ compile-time tile hyperparameters. + +mmq_x is swept inside mmq-tune at runtime, but mmq_y and the warp count are device-side +constexpr, so each combination needs its own build. This drives that outer loop: configure, +build only the mmq-tune target, run it, append the CSV. + + python tools/mmq-tune/sweep.py --mmq-y 32,64,128 --mmq-nwarps 2,4,8 + +Combinations that fail to build or run are recorded rather than aborting the sweep: the kernel +asserts nwarps*tile_C::I == mmq_y, so most pairs do not compile. +""" +from __future__ import annotations + +import argparse +import csv +import io +import os +import subprocess +import sys +from collections import defaultdict +from pathlib import Path + +REPO = Path(__file__).resolve().parents[2] + + +def run(cmd, **kw): + print("+", " ".join(str(c) for c in cmd), flush=True) + return subprocess.run(cmd, **kw) + + +def build(build_dir, mmq_y, nwarps, jobs, cmake_args): + cfg = [ + "cmake", "-S", str(REPO), "-B", str(build_dir), + "-DGGML_HIP=ON", "-DGGML_MMQ_TUNE=ON", + f"-DGGML_HIP_MMQ_Y={mmq_y}", f"-DGGML_HIP_MMQ_NWARPS={nwarps}", + *cmake_args, + ] + if run(cfg, capture_output=True, text=True).returncode != 0: + return "configure_failed" + bld = ["cmake", "--build", str(build_dir), "--target", "mmq-tune", "-j", str(jobs)] + if run(bld, capture_output=True, text=True).returncode != 0: + return "build_failed" + return None + + +def measure(build_dir, extra_args, gpu_lock, retries=6): + cmd = ([gpu_lock] if gpu_lock else []) + [ + str(build_dir / "bin" / "mmq-tune"), "--output", "csv", *extra_args + ] + # amd-gpu-lock gives up after 10 minutes; retry rather than lose a build's measurements. + for attempt in range(retries): + p = run(cmd, capture_output=True, text=True) + if p.returncode == 0: + # The launcher and the ggml backend print to stdout before the CSV header. + lines = p.stdout.splitlines() + for i, line in enumerate(lines): + if line.startswith("model,kind,"): + return "\n".join(lines[i:]) + sys.stderr.write(p.stderr[-2000:]) + print(f" measure attempt {attempt + 1}/{retries} failed", flush=True) + return None + + +def report(rows): + """Best width per shape, then the (mmq_y, nwarps) ranking those widths imply.""" + rows = [r for r in rows if r.get("mmq_x")] + if not rows: + return + + def key(r): + return (r["kind"], r["type"], r["n"], r["k"], r["m"]) + + def usable(r): + # mmq_x < 0 is the baseline measurement, not a candidate width. + return int(r["mmq_x"]) > 0 + + print("\n=== best config per shape ===") + best = {} + for r in rows: + if not usable(r): + continue + cur = best.get(key(r)) + if cur is None or float(r["us_median"]) < float(cur["us_median"]): + best[key(r)] = r + for k in sorted(best, key=lambda t: (t[0], t[1], int(t[2]), int(t[3]), int(t[4]))): + r = best[k] + print(f" {r['kind']:<5} {r['type']:>5} N={r['n']:>6} K={r['k']:>6} M={r['m']:>5} -> " + f"mmq_y={r['mmq_y']:>3} nwarps={r['nwarps']} mmq_x={r['mmq_x']:>3} " + f"{float(r['us_median']):9.2f} us") + + # mmq_y and nwarps are fixed per build, so the ranking is over their totals, weighted by how + # often each shape runs. + print("\n=== global (mmq_y, nwarps), weighted by invocations per forward pass ===") + per_build = defaultdict(dict) + weight = {} + for r in rows: + if not usable(r): + continue + b = (int(r["mmq_y"]), int(r["nwarps"])) + us = float(r["us_median"]) + if key(r) not in per_build[b] or us < per_build[b][key(r)]: + per_build[b][key(r)] = us + weight[key(r)] = int(r["weight"]) + + print(f" {'build':>16}{'dense ms':>12}{'MoE ms':>12}{'total ms':>12}") + ranked = [] + for b, v in per_build.items(): + dense = sum(us*weight[k] for k, us in v.items() if k[0] == "dense")/1000 + moe = sum(us*weight[k] for k, us in v.items() if k[0] == "moe")/1000 + ranked.append((dense + moe, dense, moe, b)) + for total, dense, moe, b in sorted(ranked): + print(f" mmq_y={b[0]:>3} nwarps={b[1]}{dense:12.1f}{moe:12.1f}{total:12.1f}") + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--mmq-y", default="32,64,128") + ap.add_argument("--mmq-nwarps", default="2,4,8") + ap.add_argument("--build-dir", default=str(REPO / "build-tune")) + ap.add_argument("--out", default="sweep.csv") + ap.add_argument("--jobs", type=int, default=os.cpu_count()) + ap.add_argument("--gpu-lock", default="amd-gpu-lock", help="launcher prefix, empty to disable") + ap.add_argument("--cmake-arg", action="append", default=[], help="extra -D flag forwarded to cmake configure") + ap.add_argument("mmq_tune_args", nargs="*", help="extra args forwarded to mmq-tune") + args = ap.parse_args() + + build_dir = Path(args.build_dir).resolve() + rows = [] + + with open(args.out, "w", newline="") as fout: + writer = None + for mmq_y in [int(v) for v in args.mmq_y.split(",")]: + for nwarps in [int(v) for v in args.mmq_nwarps.split(",")]: + print(f"\n### mmq_y={mmq_y} nwarps={nwarps}", flush=True) + + status = build(build_dir, mmq_y, nwarps, args.jobs, args.cmake_arg) + csv_text = None if status else measure(build_dir, args.mmq_tune_args, args.gpu_lock) + if csv_text is None: + status = status or "run_failed" + print(f" {status}", flush=True) + if writer: + writer.writerow({**{k: "" for k in writer.fieldnames}, + "mmq_y": mmq_y, "nwarps": nwarps, "status": status}) + continue + + for row in csv.DictReader(io.StringIO(csv_text)): + row["status"] = "ok" + if writer is None: + writer = csv.DictWriter(fout, fieldnames=list(row.keys())) + writer.writeheader() + writer.writerow(row) + rows.append(row) + fout.flush() + print(f" ok ({len(rows)} rows total)", flush=True) + + report(rows) + print(f"\nwrote {args.out}") + + +if __name__ == "__main__": + main()