diff --git a/ggml/src/ggml-cuda/mmvdq.cu b/ggml/src/ggml-cuda/mmvdq.cu index 5a27f325a4b2..62011e78927e 100644 --- a/ggml/src/ggml-cuda/mmvdq.cu +++ b/ggml/src/ggml-cuda/mmvdq.cu @@ -186,6 +186,31 @@ static void launch_dq(const void * vx, const float * y, float * d, int ncols_x, } } +template <> struct dq_traits { + using block_t = block_q4_0; + static constexpr int qk = QK4_0; + static constexpr int lanes_per_block = 2; + struct geom { int ix, q_offset; }; + struct act { float4 lo0, lo4, hi0, hi4; }; + + static __device__ __forceinline__ geom setup(int tid) { + return { tid / lanes_per_block, 8 * (tid % lanes_per_block) }; + } + + static __device__ __forceinline__ act load(const float * yb, const geom & g) { + const float * y = yb + g.q_offset; + return { *(const float4 *) y, *(const float4 *) (y + 4), + *(const float4 *) (y + 16), *(const float4 *) (y + 20) }; + } + + static __device__ __forceinline__ float dot(const block_t * b, const geom & g, const act & a) { + const uint16_t * qs = (const uint16_t *) (b->qs + g.q_offset); + const uint32_t q0 = (uint32_t) qs[0] | ((uint32_t) qs[1] << 16); + const uint32_t q4 = (uint32_t) qs[2] | ((uint32_t) qs[3] << 16); + return __half2float(b->d) * (dq_dot_q4_0_packed(q0, a.lo0, a.hi0) + dq_dot_q4_0_packed(q4, a.lo4, a.hi4)); + } +}; + // ---- Q4_K geometry ---- struct dq_geom_q4_K { int itid, ix, v_im, q_offset, y_offset; @@ -452,6 +477,7 @@ static void dq_switch_ncols_dst(const dq_launch_args & a, int ncols_dst) { // numerics live in dq_traits; adding a type is a new case here plus a Traits. static void dq_switch_type(ggml_type type, const dq_launch_args & a, int ncols_dst) { switch (type) { + case GGML_TYPE_Q4_0: dq_switch_ncols_dst(a, ncols_dst); break; case GGML_TYPE_Q8_0: dq_switch_ncols_dst(a, ncols_dst); break; case GGML_TYPE_Q4_K: dq_switch_ncols_dst(a, ncols_dst); break; case GGML_TYPE_Q5_K: dq_switch_ncols_dst(a, ncols_dst); break; @@ -488,7 +514,7 @@ bool ggml_cuda_should_use_mmv_dq( } const ggml_type type = src0->type; - if (!(type == GGML_TYPE_Q8_0 || type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K || type == GGML_TYPE_Q6_K)) { + if (!(type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q8_0 || type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K || type == GGML_TYPE_Q6_K)) { return false; } @@ -509,7 +535,7 @@ bool ggml_cuda_should_use_mmv_dq( } } - const int qk = type == GGML_TYPE_Q8_0 ? QK8_0 : QK_K; + const int qk = ggml_blck_size(type); if (src0->ne[0] % qk != 0) { return false; } diff --git a/ggml/src/ggml-cuda/vecdotdq.cuh b/ggml/src/ggml-cuda/vecdotdq.cuh index a53fcf9cb3a7..e666080a1b7f 100644 --- a/ggml/src/ggml-cuda/vecdotdq.cuh +++ b/ggml/src/ggml-cuda/vecdotdq.cuh @@ -160,6 +160,17 @@ static __device__ __forceinline__ float dq_dot_q6_K( return d * (sum0*sc0 + sum1*sc2 + sum2*sc4 + sum3*sc6); } +static __device__ __forceinline__ float dq_dot_q4_0_packed(uint32_t q, const float4 & lo, const float4 & hi) { + return lo.x * (float) ((int) ((q >> 0) & 0xF) - 8) + + lo.y * (float) ((int) ((q >> 8) & 0xF) - 8) + + lo.z * (float) ((int) ((q >> 16) & 0xF) - 8) + + lo.w * (float) ((int) ((q >> 24) & 0xF) - 8) + + hi.x * (float) ((int) ((q >> 4) & 0xF) - 8) + + hi.y * (float) ((int) ((q >> 12) & 0xF) - 8) + + hi.z * (float) ((int) ((q >> 20) & 0xF) - 8) + + hi.w * (float) ((int) ((q >> 28) & 0xF) - 8); +} + struct dq_weight_q8_0 { uint32_t qs0, qs4; half d; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index b25accc18a07..494924804daf 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -1592,11 +1592,12 @@ struct test_case { ggml_cgraph * gf = ggml_new_graph_custom(ctx.get(), graph_nodes, false); ggml_build_forward_expand(gf, out); - // warmup run - ggml_status status = ggml_backend_graph_compute(backend, gf); - if (status != GGML_STATUS_SUCCESS) { - fprintf(stderr, "%s: ggml_backend_graph_compute failed. status=%s \n", __func__, ggml_status_to_string(status)); - return false; + { + ggml_status status = ggml_backend_graph_compute(backend, gf); + if (status != GGML_STATUS_SUCCESS) { + fprintf(stderr, "%s: ggml_backend_graph_compute failed. status=%s \n", __func__, ggml_status_to_string(status)); + return false; + } } // optional cache-flush scratch (GGML_TEST_FLUSH_CACHE): evict the op's working set @@ -1706,7 +1707,7 @@ struct test_case { if (flush_cache && total_runs >= flush_iters) { break; } - } while (total_time_us < 1000*1000); // run for at least 1 second + } while (total_time_us < 1000*1000); // Create test result double avg_time_us = (double) total_time_us / total_runs; @@ -8994,6 +8995,15 @@ static constexpr mmq_test_shape rdna35_mmq_shapes[] = { }; // RDNA3.5 MMQ production shapes (gfx1151). Shared by eval (correctness) and perf suites. +static void add_q8_0_decode_cases(std::vector> & test_cases) { + for (const auto & shape : {std::pair{960, 1280}, {12288, 4096}, {6144, 8192}}) { + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, GGML_GLU_OP_SWIGLU, 1, shape.first, shape.second, false, 1, 1, false, false, true, false, {1, 1})); + } + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 1024, 1, 4096, {1, 1}, {1, 1})); + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 4096, 1, 12288, {1, 1}, {1, 1})); + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 151936, 1, 4096, {1, 1}, {1, 1})); +} + static void add_rdna35_mmq_cases(std::vector> & test_cases) { for (const ggml_type type : rdna35_mmq_types) { for (const mmq_test_shape & shape : rdna35_mmq_shapes) { @@ -9006,7 +9016,21 @@ static void add_rdna35_mmq_cases(std::vector> & test_ test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_K, GGML_TYPE_F32, 4096, 16, 4096, {1, 1}, {1, 1})); test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_K, GGML_TYPE_F32, 4096, 1024, 12288, {1, 1}, {1, 1})); - // corpus (tools/model-parse/kernel-shape-weights/unique_kernel_shapes.tsv). n=1 exercises + static constexpr mmq_test_shape q4_0_decode_model_shapes[] = { + {17408, 1, 5120}, + { 5120, 1, 17408}, + {12288, 1, 4096}, + }; + for (const mmq_test_shape & shape : q4_0_decode_model_shapes) { + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, shape.m, shape.n, shape.k, {1, 1}, {1, 1})); + } + for (ggml_glu_op glu_op : {GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU}) { + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_0, glu_op, 1, 17408, 5120, + false, 1, 1, false, false, true, false, {1, 1})); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_0, glu_op, 1, 12288, 4096, + false, 1, 1, false, false, true, false, {1, 1})); + } + // Array is {m=N (out rows), n=1, k=K (contraction)}. static constexpr mmq_test_shape q4k_decode_model_shapes[] = { { 21504, 1, 5376}, @@ -10045,6 +10069,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1})); test_cases.emplace_back(new test_mul_mat(GGML_TYPE_MXFP4, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1})); + add_q8_0_decode_cases(test_cases); add_rdna35_mmq_cases(test_cases); // m == 1, with n on both sides of MMVF_MAX_BATCH_SIZE (8): mmvf below, operand swap above @@ -11366,25 +11391,7 @@ static std::vector> make_test_cases_perf() { } } - for (int k : {768, 960, 1024, 1280, 2048, 4096, 8192}) { - for (int n : {320, 960, 2560, 4096, 12288}) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, GGML_GLU_OP_SWIGLU, 1, n, k, false, 1, 1, false, false, true, false, {1, 1})); - } - } - for (int k : {4096, 8192}) { - for (int n : {3072, 3584, 4608, 5120, 6144}) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, GGML_GLU_OP_SWIGLU, 1, n, k, false, 1, 1, false, false, true, false, {1, 1})); - } - } - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 4096, 1, 12288, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 1024, 1, 4096, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 151936, 1, 896, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 262144, 1, 1536, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 262144, 1, 2560, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 262144, 1, 2816, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 49280, 1, 960, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 248320, 1, 2048, {1, 1}, {1, 1})); - test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 151936, 1, 4096, {1, 1}, {1, 1})); + add_q8_0_decode_cases(test_cases); add_rdna35_mmq_cases(test_cases); // Q4_K multi-column mat-vec