diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 8c01e7685..2b252f541 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -172,13 +172,18 @@ jobs: -DDFLASH27B_ENABLE_BSA=OFF \ -DDFLASH27B_FA_ALL_QUANTS=OFF \ -DCMAKE_BUILD_TYPE=Release - cmake --build build --target test_flash_attn_sparse -j"$(nproc)" + cmake --build build \ + --target test_flash_attn_sparse test_deepseek4_mmid_grouped_cuda \ + -j"$(nproc)" - name: Run flash-attn sparse kernel test on the 3090 # Self-contained kernel correctness test (no model weights), so it is # safe and fast in CI. This is the execution the hosted jobs cannot do. run: ./server/build/test_flash_attn_sparse + - name: Run grouped MMID dispatch and parity test on the 3090 + run: ./server/build/test_deepseek4_mmid_grouped_cuda + # Optional model-backed end-to-end smoke (real spec-decode on the 3090), # disabled by default because it builds dflash_server and lazy-loads the # ~16 GB Qwen3.6-27B target + draft (~1-2 min). The weights are already @@ -285,7 +290,8 @@ jobs: -DCMAKE_HIP_FLAGS=-DDFLASH_WAVE_SIZE=32 cmake --build "$RUNNER_TEMP/rocmfp-build" \ --target test_rocmfp4 test_rocmfpx test_rocmfp4_hip_tail test_rocmfpx_mmq \ + test_deepseek4_mmid_grouped_cuda \ --parallel 8 ctest --test-dir "$RUNNER_TEMP/rocmfp-build" \ --output-on-failure \ - -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail|rocmfpx_mmq' + -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail|rocmfpx_mmq|deepseek4_mmid_grouped_cuda' diff --git a/server/CMakeLists.txt b/server/CMakeLists.txt index a07701382..243b5b0e3 100644 --- a/server/CMakeLists.txt +++ b/server/CMakeLists.txt @@ -749,6 +749,22 @@ if(DFLASH27B_TESTS) target_link_libraries(test_deepseek4_hc_cuda PRIVATE dflash_common ${DFLASH27B_GGML_BACKEND_TARGET}) add_test(NAME deepseek4_hc_cuda COMMAND test_deepseek4_hc_cuda) endif() + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_deepseek4_mmid_grouped_cuda.cpp") + add_executable(test_deepseek4_mmid_grouped_cuda test/test_deepseek4_mmid_grouped_cuda.cpp) + if(DFLASH27B_GPU_BACKEND STREQUAL "hip") + set_source_files_properties(test/test_deepseek4_mmid_grouped_cuda.cpp PROPERTIES LANGUAGE HIP) + set_target_properties(test_deepseek4_mmid_grouped_cuda PROPERTIES HIP_ARCHITECTURES "${_dflash_archs}") + target_include_directories(test_deepseek4_mmid_grouped_cuda PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/hip_compat) + target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE hip::host) + else() + target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE CUDA::cudart) + endif() + target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE + ggml ${DFLASH27B_GGML_BACKEND_TARGET}) + add_test(NAME deepseek4_mmid_grouped_cuda COMMAND test_deepseek4_mmid_grouped_cuda) + set_tests_properties(deepseek4_mmid_grouped_cuda PROPERTIES SKIP_RETURN_CODE 77) + endif() # GPU draft top-K kernel vs CPU reference (extract_draft_topk). Built on both # backends: geometric_draft_topk_cuda.cu is compiled into dflash_common on the # cuda backend directly and on hip via LANGUAGE HIP + the hip_compat shim. diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu index 14e464ea8..b2801befc 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2462,6 +2462,10 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) { // fusion is not universally faster on Pascal const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + if (is_mul_mat_id && ggml_cuda_mmvq_mmid_grouped_enabled( + src0->type, cc, dst->ne[2], dst->ne[1]*dst->ne[2])) { + return false; + } if (cc <= GGML_CUDA_CC_PASCAL) { return false; } @@ -2498,6 +2502,12 @@ static bool ggml_cuda_try_fuse_mul_mat_glu( const ggml_tensor * src1 = up->src[1]; const ggml_tensor * ids = up->src[2]; + const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + if (ids && ggml_cuda_mmvq_mmid_grouped_enabled( + src0->type, cc, up->ne[2], up->ne[1]*up->ne[2])) { + return false; + } + if (ggml_cuda_should_fuse_mul_mat_vec_f(up)) { ggml_cuda_mm_fusion_args_host fusion_data{}; fusion_data.gate = gate->src[0]; @@ -2531,7 +2541,6 @@ static bool ggml_cuda_try_fuse_mul_mat_glu( } } - const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; const int64_t ncols = ids ? src1->ne[2] : src1->ne[1]; if (ggml_cuda_should_use_mmq(src0->type, cc, ncols, src0->ne[2])) { ggml_cuda_mul_mat_q_pair( diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu index 4906f88fd..5ab168f3b 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu @@ -256,7 +256,8 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type // weight traffic toward the union of routed experts. Bit-exact per // (row, token) vs mul_mat_vec_q_moe (same vec_dot sequence and reduction). // Model-agnostic: applies to any MoE with n_expert_used*n_tokens <= 256. -// DFLASH_MMID_GROUPED=1 opt-in +// DFLASH_MMID_GROUPED 1 = enable, 0 = disable +// (CUDA default on; HIP default off) // DFLASH_MMID_GROUPED_TYPES bitmask, 1 = Q4_K (default), 2 = Q6_K, // 4 = Q4_0/Q8_0/Q5_K. Q6_K stays on its tuned // MMQ route above 5 tokens unless enabled. @@ -285,7 +286,9 @@ struct mmid_gate_extra { static bool mmid_grouped_env() { // Bit-exact and measured equal-or-faster on small MoE verify batches, so // enabled by default on CUDA; DFLASH_MMID_GROUPED=0 is the kill switch. - // HIP is unvalidated and stays opt-in (DFLASH_MMID_GROUPED=1). + // HIP (RDNA3/RDNA4) correctness is validated on gfx1151, but performance is + // still unmeasured, so it stays opt-in and default-off + // (DFLASH_MMID_GROUPED=1 to enable). static const bool on = []() { const char * e = std::getenv("DFLASH_MMID_GROUPED"); if (e != nullptr) { @@ -324,11 +327,26 @@ static bool mmid_grouped_type_ok(ggml_type type) { } } +static bool mmid_grouped_arch_ok(int cc) { + return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_TURING) || + GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc); +} + +bool ggml_cuda_mmvq_mmid_grouped_enabled( + ggml_type type, int cc, int64_t ncols_dst, int64_t routed_pairs) { + return ncols_dst >= 2 && ncols_dst <= MMVQ_MAX_MOE_BATCH_SIZE && + routed_pairs <= MMID_GROUPED_MAX_PAIRS && + mmid_grouped_env() && mmid_grouped_type_ok(type) && mmid_grouped_arch_ok(cc); +} + // Host function: returns the max batch size for the current arch+type at runtime. int get_mmvq_mmid_max_batch(ggml_type type, int cc) { // [TAG_MMID_GROUPED] the grouped kernel handles any supported type up to the // MoE batch ceiling; this also keeps CUDA graphs on for these batches. - if (mmid_grouped_env() && mmid_grouped_type_ok(type) && GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_TURING) { + // RDNA3/RDNA4 (wave32) share the non-grouped kernel's wave-width warp_reduce. + // IS_RDNA3/IS_RDNA4 are pure cc-range checks, safe above the IS_AMD guard below. + // The HIP path remains opt-in until on-hardware parity and performance validation. + if (mmid_grouped_env() && mmid_grouped_type_ok(type) && mmid_grouped_arch_ok(cc)) { return MMVQ_MAX_MOE_BATCH_SIZE; } // Dedicated multi-token MoE kernel: extend the MUL_MAT_ID ceiling to 16 @@ -1122,6 +1140,8 @@ static __global__ void mul_mat_vec_q_moe_grouped( const float * x_bias = nullptr; const float * gate_bias = nullptr; ggml_glu_op active_glu; + float glu_param0 = 0.0f; + float glu_param1 = 0.0f; if constexpr (has_fusion) { use_gate = fusion.gate != nullptr; @@ -1131,6 +1151,8 @@ static __global__ void mul_mat_vec_q_moe_grouped( x_bias = (const float *) fusion.x_bias; gate_bias = (const float *) fusion.gate_bias; active_glu = fusion.glu_op; + glu_param0 = fusion.glu_param0; + glu_param1 = fusion.glu_param1; } float tmp[c_rows_per_block] = {0.0f}; @@ -1180,7 +1202,11 @@ static __global__ void mul_mat_vec_q_moe_grouped( result *= ggml_cuda_op_gelu_single(gate_value); break; case GGML_GLU_OP_SWIGLU_OAI: - result = ggml_cuda_op_swiglu_oai_single(gate_value, result); + result = ggml_cuda_op_swiglu_oai_single( + gate_value, result, glu_param0, glu_param1); + break; + case GGML_GLU_OP_SWIGLU_DS4: + result = ggml_cuda_op_swiglu_ds4_single(gate_value, result, glu_param0); break; default: result = result * gate_value; @@ -1192,7 +1218,9 @@ static __global__ void mul_mat_vec_q_moe_grouped( } if constexpr (!has_fusion) { - GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, vgate, x_bias, gate_bias, active_glu, tmp_gate); + GGML_UNUSED_VARS( + use_gate, use_bias, use_gate_bias, vgate, x_bias, gate_bias, + active_glu, glu_param0, glu_param1, tmp_gate); } } @@ -2000,15 +2028,15 @@ void ggml_cuda_mul_mat_vec_q( const int64_t stride_channel_y = ids ? s11 : s12; const int64_t ids_stride = ids ? ids->nb[1] / ggml_type_size(ids->type) : 0; + const int cc = ggml_cuda_info().devices[ctx.device].cc; static const bool mmid_telemetry = []() { const char * value = std::getenv("DFLASH_MMID_TELEMETRY"); return value != nullptr && std::strcmp(value, "0") != 0; }(); // [TAG_MMID_GROUPED] grouped-expert path for small MUL_MAT_ID batches. - if (ids && ncols_dst >= 2 && ncols_dst <= MMVQ_MAX_MOE_BATCH_SIZE && - (int) (nchannels_dst*ncols_dst) <= MMID_GROUPED_MAX_PAIRS && - mmid_grouped_env() && mmid_grouped_type_ok(src0->type)) { + if (ids && ggml_cuda_mmvq_mmid_grouped_enabled( + src0->type, cc, ncols_dst, nchannels_dst*ncols_dst)) { // Batches above MMID_GROUPED_MAX_PAIRS fall through to the legacy // per-expert kernel instead of aborting the request. const int np = (int) (nchannels_dst*ncols_dst); @@ -2080,7 +2108,8 @@ void ggml_cuda_mul_mat_vec_q( ncols_dst > MMVQ_MAX_MOE_BATCH_SIZE ? "width_gt_16" : (int) (nchannels_dst*ncols_dst) > MMID_GROUPED_MAX_PAIRS ? "pairs_gt_256" : !mmid_grouped_env() ? "flag_off" : - !mmid_grouped_type_ok(src0->type) ? "unsupported_type" : "dispatch_rejected"; + !mmid_grouped_type_ok(src0->type) ? "unsupported_type" : + !mmid_grouped_arch_ok(cc) ? "unsupported_arch" : "dispatch_rejected"; std::fprintf(stderr, "[dflash-mmid] event=mmvq type=%s width=%lld pairs=%lld variant=%s reason=%s\n", ggml_type_name(src0->type), (long long) ncols_dst, diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cuh index 5fbb753f3..209557fd9 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cuh @@ -11,6 +11,9 @@ // based on the quantization type and GPU architecture (compute capability). int get_mmvq_mmid_max_batch(ggml_type type, int cc); +bool ggml_cuda_mmvq_mmid_grouped_enabled( + ggml_type type, int cc, int64_t ncols_dst, int64_t routed_pairs); + void ggml_cuda_mul_mat_vec_q(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst, const ggml_cuda_mm_fusion_args_host * fusion = nullptr); diff --git a/server/test/test_deepseek4_mmid_grouped_cuda.cpp b/server/test/test_deepseek4_mmid_grouped_cuda.cpp new file mode 100644 index 000000000..975b772fd --- /dev/null +++ b/server/test/test_deepseek4_mmid_grouped_cuda.cpp @@ -0,0 +1,403 @@ +#include "ggml-alloc.h" +#include "ggml-backend.h" +#include "ggml-cuda.h" +#include "ggml.h" + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#if defined(_WIN32) +#include +#else +#include +#endif + +static bool run_case( + ggml_backend_t backend, + ggml_type type, + int width, + bool fused_ds4, + bool write_output, + std::ofstream & output) { + constexpr int k_dim = 256; + constexpr int n_rows = 128; + constexpr int n_experts = 32; + constexpr int top_k = 8; + + ggml_init_params params = {16 * 1024 * 1024, nullptr, true}; + ggml_context * ctx = ggml_init(params); + if (ctx == nullptr) { + return false; + } + + ggml_tensor * weights = ggml_new_tensor_3d(ctx, type, k_dim, n_rows, n_experts); + ggml_tensor * gate_weights = + fused_ds4 ? ggml_new_tensor_3d(ctx, type, k_dim, n_rows, n_experts) : nullptr; + ggml_tensor * input = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, k_dim, 1, width); + ggml_tensor * ids = ggml_new_tensor_2d(ctx, GGML_TYPE_I32, top_k, width); + ggml_set_input(weights); + if (gate_weights != nullptr) { + ggml_set_input(gate_weights); + } + ggml_set_input(input); + ggml_set_input(ids); + + ggml_tensor * result = ggml_mul_mat_id(ctx, weights, input, ids); + if (gate_weights != nullptr) { + ggml_tensor * gate = ggml_mul_mat_id(ctx, gate_weights, input, ids); + result = ggml_swiglu_ds4_split(ctx, gate, result, 1.5f); + } + ggml_set_output(result); + ggml_cgraph * graph = ggml_new_graph(ctx); + ggml_build_forward_expand(graph, result); + + ggml_gallocr_t alloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend)); + if (!ggml_gallocr_alloc_graph(alloc, graph)) { + ggml_gallocr_free(alloc); + ggml_free(ctx); + return false; + } + + std::mt19937 rng(20260713u + (unsigned) type * 97u + (unsigned) width); + std::uniform_real_distribution dist(-1.0f, 1.0f); + std::vector weights_f((size_t) k_dim * n_rows * n_experts); + std::vector input_f((size_t) k_dim * width); + for (float & value : weights_f) { + value = dist(rng); + } + for (float & value : input_f) { + value = dist(rng); + } + + std::vector weights_q(ggml_nbytes(weights)); + const size_t quantized = ggml_quantize_chunk( + type, weights_f.data(), weights_q.data(), 0, n_rows * n_experts, k_dim, nullptr); + if (quantized != weights_q.size()) { + std::fprintf(stderr, "quantize size mismatch type=%s got=%zu expected=%zu\n", + ggml_type_name(type), quantized, weights_q.size()); + ggml_gallocr_free(alloc); + ggml_free(ctx); + return false; + } + + std::vector ids_h((size_t) top_k * width); + for (int token = 0; token < width; ++token) { + for (int slot = 0; slot < top_k; ++slot) { + ids_h[(size_t) token * top_k + slot] = (token * 3 + slot * 5) % n_experts; + } + } + + ggml_backend_tensor_set(weights, weights_q.data(), 0, weights_q.size()); + if (gate_weights != nullptr) { + ggml_backend_tensor_set(gate_weights, weights_q.data(), 0, weights_q.size()); + } + ggml_backend_tensor_set(input, input_f.data(), 0, input_f.size() * sizeof(float)); + ggml_backend_tensor_set(ids, ids_h.data(), 0, ids_h.size() * sizeof(int32_t)); + ggml_backend_synchronize(backend); + + const ggml_status status = ggml_backend_graph_compute(backend, graph); + std::vector result_h(ggml_nelements(result)); + if (status == GGML_STATUS_SUCCESS) { + ggml_backend_synchronize(backend); + ggml_backend_tensor_get(result, result_h.data(), 0, result_h.size() * sizeof(float)); + if (write_output) { + output.write(reinterpret_cast(result_h.data()), + (std::streamsize) (result_h.size() * sizeof(float))); + } + } + + std::printf("[mmid-grouped-test] type=%s width=%d fused_ds4=%d pairs=%d status=%d bytes=%zu\n", + ggml_type_name(type), width, fused_ds4 ? 1 : 0, width * top_k, (int) status, + result_h.size() * sizeof(float)); + ggml_gallocr_free(alloc); + ggml_free(ctx); + return status == GGML_STATUS_SUCCESS && output.good(); +} + +static int run_child(const char * mode, const char * output_path) { + const bool grouped = std::strcmp(mode, "grouped") == 0; + if (!grouped && std::strcmp(mode, "legacy") != 0) { + return 2; + } +#if defined(_WIN32) + if (!grouped) { + _putenv_s("GGML_CUDA_DISABLE_FUSION", "1"); + } else { + _putenv_s("GGML_CUDA_DISABLE_FUSION", ""); + } +#else + if (!grouped) { + setenv("GGML_CUDA_DISABLE_FUSION", "1", 1); + } else { + unsetenv("GGML_CUDA_DISABLE_FUSION"); + } +#endif + + ggml_backend_t backend = ggml_backend_cuda_init(0); + if (backend == nullptr) { + std::fprintf(stderr, "GPU backend unavailable\n"); + return 1; + } + + std::ofstream output(output_path, std::ios::binary | std::ios::trunc); + const ggml_type types[] = { + GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q5_K, + }; + const int widths[] = {2, 4, 8, 9, 16}; + bool ok = output.good(); + for (ggml_type type : types) { + for (int width : widths) { + ok = run_case(backend, type, width, false, true, output) && ok; + ok = run_case(backend, type, width, true, true, output) && ok; + } + } + output.close(); + ggml_backend_free(backend); + return ok ? 0 : 1; +} + +static std::vector read_file(const char * path) { + std::ifstream input(path, std::ios::binary | std::ios::ate); + if (!input) { + return {}; + } + const std::streamsize size = input.tellg(); + input.seekg(0); + std::vector data((size_t) size); + input.read(data.data(), size); + return data; +} + +static size_t count_records(const std::vector & data, const char * needle) { + const std::string text(data.begin(), data.end()); + size_t count = 0; + size_t offset = 0; + while ((offset = text.find(needle, offset)) != std::string::npos) { + ++count; + offset += std::strlen(needle); + } + return count; +} + +static bool has_mmvq_record( + const std::vector & log, + ggml_type type, + int width, + const char * variant = nullptr) { + const std::string type_field = "type=" + std::string(ggml_type_name(type)) + " "; + const std::string width_field = "width=" + std::to_string(width) + " "; + const std::string variant_field = variant ? "variant=" + std::string(variant) : std::string(); + std::istringstream lines(std::string(log.begin(), log.end())); + std::string line; + while (std::getline(lines, line)) { + if (line.find("[dflash-mmid] event=mmvq ") == std::string::npos || + line.find(type_field) == std::string::npos || + line.find(width_field) == std::string::npos) { + continue; + } + if (!variant || line.find(variant_field) != std::string::npos) { + return true; + } + } + return false; +} + +static bool compare_case_outputs( + const std::vector & legacy, + const std::vector & grouped, + size_t offset, + size_t case_bytes, + bool require_exact) { + if (offset + case_bytes > legacy.size() || offset + case_bytes > grouped.size()) { + return false; + } + if (require_exact) { + return std::equal( + legacy.begin() + (std::ptrdiff_t) offset, + legacy.begin() + (std::ptrdiff_t) (offset + case_bytes), + grouped.begin() + (std::ptrdiff_t) offset); + } + + double squared_error = 0.0; + double reference_power = 0.0; + bool finite = true; + for (size_t byte = 0; byte < case_bytes; byte += sizeof(float)) { + float expected; + float actual; + std::memcpy(&expected, legacy.data() + offset + byte, sizeof(float)); + std::memcpy(&actual, grouped.data() + offset + byte, sizeof(float)); + finite = finite && std::isfinite(expected) && std::isfinite(actual); + const double error = (double) actual - expected; + squared_error += error * error; + reference_power += (double) expected * expected; + } + + // The unfused legacy graph is an independent numerical oracle for fused + // DS4 cases, and legacy MMQ is the oracle for cases that did not previously + // fit the architecture's MMVQ ceiling. Match the repository's MMQ tolerance + // for the resulting reduction-order and fused-operation differences. + constexpr double max_nmse = 5e-4; + const double nmse = squared_error / std::max(reference_power, 1e-30); + return finite && nmse <= max_nmse; +} + +static bool grouped_supported_device() { + cudaDeviceProp prop{}; + if (cudaGetDeviceProperties(&prop, 0) != cudaSuccess) { + return false; + } +#if defined(__HIP_PLATFORM_AMD__) + const std::string arch = prop.gcnArchName; + return arch.rfind("gfx11", 0) == 0 || arch.rfind("gfx12", 0) == 0; +#else + return prop.major * 10 + prop.minor >= 75; +#endif +} + +static std::string shell_quote(const std::string & value) { +#if defined(_WIN32) + std::string quoted = "\""; + for (const char c : value) { + if (c == '"') { + quoted += '\\'; + } + quoted += c; + } + return quoted + "\""; +#else + std::string quoted = "'"; + for (const char c : value) { + if (c == '\'') { + quoted += "'\\''"; + } else { + quoted += c; + } + } + return quoted + "'"; +#endif +} + +static std::string child_command( + const std::string & executable, + const char * mode, + const std::string & output_path, + const std::string & log_path) { +#if defined(_WIN32) + return "set \"DFLASH_MMID_TELEMETRY=1\" && set \"DFLASH_MMID_GROUPED_TYPES=7\" && " + "set \"DFLASH_MMID_GROUPED=" + + std::string(std::strcmp(mode, "grouped") == 0 ? "1" : "0") + "\" && " + + shell_quote(executable) + " --child " + mode + " " + shell_quote(output_path) + + " 2>" + shell_quote(log_path); +#else + return "DFLASH_MMID_TELEMETRY=1 DFLASH_MMID_GROUPED_TYPES=7 DFLASH_MMID_GROUPED=" + + std::string(std::strcmp(mode, "grouped") == 0 ? "1" : "0") + " " + + shell_quote(executable) + " --child " + mode + " " + shell_quote(output_path) + + " 2>" + shell_quote(log_path); +#endif +} + +int main(int argc, char ** argv) { + if (argc == 4 && std::strcmp(argv[1], "--child") == 0) { + return run_child(argv[2], argv[3]); + } + if (argc != 1) { + std::fprintf(stderr, "usage: %s [--child legacy|grouped OUTPUT]\n", argv[0]); + return 2; + } + + if (!grouped_supported_device()) { + std::printf("[mmid-grouped-test] SKIP: grouped MMID requires NVIDIA Turing+ or AMD RDNA3/RDNA4\n"); + return 77; + } + +#if defined(_WIN32) + const long long pid = (long long) _getpid(); +#else + const long long pid = (long long) getpid(); +#endif + const std::string prefix = "mmid_grouped_" + std::to_string(pid); + const std::string legacy_path = prefix + "_legacy.bin"; + const std::string grouped_path = prefix + "_enabled.bin"; + const std::string legacy_log_path = prefix + "_legacy.log"; + const std::string grouped_log_path = prefix + "_enabled.log"; + const std::string executable = argv[0]; + const std::string legacy_cmd = + child_command(executable, "legacy", legacy_path, legacy_log_path); + const std::string grouped_cmd = + child_command(executable, "grouped", grouped_path, grouped_log_path); + + const int legacy_status = std::system(legacy_cmd.c_str()); + const int grouped_status = std::system(grouped_cmd.c_str()); + const std::vector legacy = read_file(legacy_path.c_str()); + const std::vector grouped = read_file(grouped_path.c_str()); + const std::vector legacy_log = read_file(legacy_log_path.c_str()); + const std::vector grouped_log = read_file(grouped_log_path.c_str()); + const size_t legacy_grouped = count_records(legacy_log, "variant=grouped"); + const size_t grouped_grouped = count_records(grouped_log, "variant=grouped"); + + const ggml_type types[] = { + GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q5_K, + }; + const int widths[] = {2, 4, 8, 9, 16}; + size_t offset = 0; + size_t compared_bytes = 0; + int compared_cases = 0; + int exact_cases = 0; + int tolerant_cases = 0; + bool output_parity = legacy.size() == grouped.size() && !legacy.empty(); + bool grouped_dispatch = true; + for (ggml_type type : types) { + for (int width : widths) { + const bool legacy_mmvq = has_mmvq_record(legacy_log, type, width); + for (bool fused_ds4 : {false, true}) { + const size_t case_bytes = (size_t) 128 * 8 * width * sizeof(float); + const bool require_exact = legacy_mmvq && !fused_ds4; + grouped_dispatch = + has_mmvq_record(grouped_log, type, width, "grouped") && grouped_dispatch; + if (output_parity) { + output_parity = compare_case_outputs( + legacy, grouped, offset, case_bytes, require_exact); + if (!output_parity) { + std::fprintf(stderr, + "output mismatch type=%s width=%d fused_ds4=%d exact=%d\n", + ggml_type_name(type), width, fused_ds4 ? 1 : 0, + require_exact ? 1 : 0); + } + } + compared_bytes += case_bytes; + ++compared_cases; + exact_cases += require_exact ? 1 : 0; + tolerant_cases += require_exact ? 0 : 1; + offset += case_bytes; + } + } + } + output_parity = output_parity && offset == legacy.size() && compared_cases == 50; + const bool pass = legacy_status == 0 && grouped_status == 0 && + output_parity && grouped_dispatch && legacy_grouped == 0 && grouped_grouped == 75; + if (pass) { + std::remove(legacy_path.c_str()); + std::remove(grouped_path.c_str()); + std::remove(legacy_log_path.c_str()); + std::remove(grouped_log_path.c_str()); + } + std::printf("[mmid-grouped-test] legacy_status=%d grouped_status=%d bytes=%zu " + "compared_cases=%d exact_cases=%d tolerant_cases=%d compared_bytes=%zu " + "legacy_grouped=%zu grouped_grouped=%zu " + "parity=%s\n", + legacy_status, grouped_status, legacy.size(), compared_cases, exact_cases, + tolerant_cases, compared_bytes, legacy_grouped, grouped_grouped, + pass ? "PASS" : "FAIL"); + return pass ? 0 : 1; +}