From 5a65131b5ccf6b61267b97ae24bcc8aa5cdad157 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Sun, 25 Oct 2026 09:46:59 +0800 Subject: [PATCH] cuda: extend the GDN conv-chain fusion and alpha/beta mat-vec pair to speculative verify widths (GGML_CUDA_GDN_FUSION=1) The fused GDN conv chain (#383) now matches batches of 1 to 16 tokens of one sequence, with one conv-state write (n_rs_seq = 0) or one per rollback snapshot slot (1 + n_rs_seq rows, up to 16, as build_conv_state writes them for MTP / DFlash verification). One launch replaces the conv-state GET_ROWS, the CONCAT, the snapshot CPYs (one batched launch since #379), SSM_CONV + SILU and the q/k L2_NORM. Same arithmetic per element as the unfused kernels: the conv sum as ssm_conv_f32, the L2 norm as l2_norm_f32 with WARP_SIZE threads (one warp per token). The ssm_alpha/ssm_beta pair now also runs at 2..8 columns through the same mmvf template instance (ncols_dst = the column count) the #380 thin-F32 path uses for each product alone; above 8 columns (mul_mat_f) it is left unpaired. The check on nodes evaluated between the CONCAT and the L2_NORM is narrowed to what moving the writes earlier requires (the SILU/L2 outputs written early, the cache, weights and ids), so a node that reuses the dead projection output (the DFlash hidden-state taps make the allocator do this in 6 of 24 layers on 9B) no longer blocks the fusion. Still opt-in under GGML_CUDA_GDN_FUSION=1. --- ggml/src/ggml-cuda/ggml-cuda.cu | 207 ++++++++++++++++++++------------ ggml/src/ggml-cuda/mmvf.cu | 24 ++-- ggml/src/ggml-cuda/ssm-conv.cu | 107 +++++++++++------ ggml/src/ggml-cuda/ssm-conv.cuh | 20 ++- tests/test-backend-ops.cpp | 76 ++++++++---- 5 files changed, 280 insertions(+), 154 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 534f6b4b8947..5e641b238942 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3066,13 +3066,15 @@ static int ggml_cuda_try_fwht_q8(ggml_backend_cuda_context & ctx, const ggml_cgr return consumed; } -// GDN decode conv chain (GGML_CUDA_GDN_FUSION=1, off by default). For one token of one sequence, a hybrid-model -// gated delta net layer runs five launch-bound kernels between the qkv projection and the GDN op: -// GET_ROWS(conv cache, s_copy) -> RESHAPE -> CONCAT(state, qkv^T) -> CPY(last window -> conv cache) +// GDN conv chain (GGML_CUDA_GDN_FUSION=1, off by default). For n_t tokens of one sequence (a decode step, n_t = 1, or a +// speculative verify step, n_t = 2..GGML_CUDA_GDN_CONV_MAX_TOKENS), a hybrid-model gated delta net layer runs these +// launch-bound kernels between the qkv projection and the GDN op: +// GET_ROWS(conv cache, s_copy) -> RESHAPE -> CONCAT(state, qkv^T) +// -> CPY(window -> conv cache row), once, or once per rollback snapshot (1 + n_rs_seq rows, #379 batches them) // -> SSM_CONV + SILU -> L2_NORM(q/k heads of the SILU output) // ggml_cuda_ssm_conv_gdn_fused computes all of them in one launch with the same arithmetic per element (bit-identical // outputs). The GET_ROWS is skipped where it stands and registered under its CONCAT (the conv cache is read there -// instead: nothing in between may write the cache or the ids); the CONCAT launches the fused kernel and marks the CPY, +// instead: nothing in between may write the cache or the ids); the CONCAT launches the fused kernel and marks the CPYs, // SSM_CONV, SILU and L2_NORM done. The CONCAT and SSM_CONV outputs are never materialised, so they must have no other // readers, and every node evaluated between the CONCAT and the L2_NORM (after the fused launch instead of before parts // of it) must not touch any memory the fused kernel reads or writes. Megakernel-style launch removal in the spirit of @@ -3085,7 +3087,8 @@ static bool ggml_cuda_gdn_fusion_enabled() { struct ggml_cuda_gdn_conv_plan { int i_gather = -1; - int i_cpy = -1; + int n_cpy = 0; + int i_cpy[GGML_CUDA_GDN_CONV_MAX_SNAP] = {}; int i_conv = -1; // the SILU is i_conv + 1 int i_l2 = -1; const ggml_tensor * gather = nullptr; @@ -3125,18 +3128,22 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml (cc->flags & GGML_TENSOR_FLAG_OUTPUT) || !ggml_is_contiguous(cc) || cc->ne[2] != 1 || cc->ne[3] != 1) { return false; } - const int64_t d_conv = cc->ne[0]; - const int64_t C = cc->ne[1]; - if (d_conv < 3 || d_conv > 5 || C > INT_MAX / 8) { - return false; - } // conv state: RESHAPE(GET_ROWS(cache, ids)) of one cache row, [d_conv - 1, C] const ggml_tensor * s0 = cc->src[0]; const ggml_tensor * gr = s0 ? s0->src[0] : nullptr; - if (s0->op != GGML_OP_RESHAPE || s0->type != GGML_TYPE_F32 || s0->ne[0] != d_conv - 1 || s0->ne[1] != C || - !gr || gr->op != GGML_OP_GET_ROWS || gr->type != GGML_TYPE_F32 || (gr->flags & GGML_TENSOR_FLAG_OUTPUT) || - !ggml_is_contiguous(gr) || gr->ne[0] != (d_conv - 1) * C || ggml_nrows(gr) != 1) { + if (!s0 || s0->op != GGML_OP_RESHAPE || s0->type != GGML_TYPE_F32) { + return false; + } + const int64_t d_conv = s0->ne[0] + 1; + const int64_t C = cc->ne[1]; + const int64_t n_t = cc->ne[0] - (d_conv - 1); + if (d_conv < 3 || d_conv > 5 || C > INT_MAX / 8 || n_t < 1 || n_t > GGML_CUDA_GDN_CONV_MAX_TOKENS) { + return false; + } + if (s0->ne[1] != C || !gr || gr->op != GGML_OP_GET_ROWS || gr->type != GGML_TYPE_F32 || + (gr->flags & GGML_TENSOR_FLAG_OUTPUT) || !ggml_is_contiguous(gr) || gr->ne[0] != (d_conv - 1) * C || + ggml_nrows(gr) != 1) { return false; } const ggml_tensor * cache = gr->src[0]; @@ -3152,47 +3159,68 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml return false; } - // new token: [1, C] (the transposed projection output) + // new tokens: [n_t, C] (the transposed projection output), token t of channel c at x + t*nb[0] + c*nb[1] const ggml_tensor * x = cc->src[1]; - if (x->type != GGML_TYPE_F32 || x->data == nullptr || x->ne[0] != 1 || x->ne[1] != C || x->ne[2] != 1 || - x->ne[3] != 1 || x->nb[1] % sizeof(float) != 0) { + if (x->type != GGML_TYPE_F32 || x->data == nullptr || x->ne[0] != n_t || x->ne[1] != C || x->ne[2] != 1 || + x->ne[3] != 1 || x->nb[1] % sizeof(float) != 0 || (n_t > 1 && x->nb[0] % sizeof(float) != 0)) { return false; } + const int64_t x_stride = (int64_t) (x->nb[1] / sizeof(float)); + const int64_t x_tok_stride = n_t > 1 ? (int64_t) (x->nb[0] / sizeof(float)) : 0; - // the next real node: CPY of the last d_conv - 1 inputs back into one cache row - int i_cpy = ic + 1; - while (i_cpy < n_nodes && ggml_cuda_is_view_or_noop(cgraph->nodes[i_cpy])) { - ++i_cpy; - } - if (i_cpy >= n_nodes) { - return false; - } - const ggml_tensor * cpy = cgraph->nodes[i_cpy]; - const ggml_tensor * win = cpy->src[0]; - const ggml_tensor * dst = cpy->src[1]; - if (cpy->op != GGML_OP_CPY || (cpy->flags & GGML_TENSOR_FLAG_OUTPUT) || win->op != GGML_OP_VIEW || - win->view_src != cc || win->view_offs != sizeof(float) || win->type != GGML_TYPE_F32 || - win->ne[0] != d_conv - 1 || win->ne[1] != C || win->ne[2] != 1 || win->ne[3] != 1 || - win->nb[0] != sizeof(float) || win->nb[1] != cc->nb[1] || - dst->type != GGML_TYPE_F32 || dst->data == nullptr || !ggml_is_contiguous(dst) || - ggml_nelements(dst) != (d_conv - 1) * C) { - return false; - } - const int i_win = ggml_cuda_find_node(cgraph, win, ic + 1, i_cpy); - if (i_win < 0 || ggml_node_get_use_count(cgraph, i_win) != 1 || ggml_node_get_use_count(cgraph, ic) != 2) { - return false; - } - // the destination is one whole cache row: the same row the state is read from, or a disjoint one + // the next real nodes: CPYs of d_conv - 1 consecutive conv inputs back into whole cache rows, each row written once + // (one CPY at a plain step, one per rollback snapshot when the memory keeps n_rs_seq > 0 snapshots) const ggml_cuda_byte_range r_cache = ggml_cuda_tensor_range(cache); - const ggml_cuda_byte_range r_dst = ggml_cuda_tensor_range(dst); - if (r_dst.beg < r_cache.beg || r_dst.end > r_cache.end || (r_dst.beg - r_cache.beg) % cache->nb[1] != 0 || - r_dst.end - r_dst.beg != (uintptr_t) cache->ne[0] * sizeof(float)) { + ggml_cuda_byte_range r_dsts[GGML_CUDA_GDN_CONV_MAX_SNAP]; + int n_cpy = 0; + int i_last_cpy = ic; + for (int j = ic + 1; j < n_nodes; ++j) { + const ggml_tensor * cpy = cgraph->nodes[j]; + if (ggml_cuda_is_view_or_noop(cpy)) { + continue; + } + if (cpy->op != GGML_OP_CPY || cpy->src[0]->op != GGML_OP_VIEW || cpy->src[0]->view_src != cc) { + break; + } + const ggml_tensor * win = cpy->src[0]; + const ggml_tensor * dst = cpy->src[1]; + if (n_cpy == GGML_CUDA_GDN_CONV_MAX_SNAP || (cpy->flags & GGML_TENSOR_FLAG_OUTPUT) || + win->view_offs % sizeof(float) != 0 || win->view_offs / sizeof(float) > (size_t) n_t || + win->type != GGML_TYPE_F32 || win->ne[0] != d_conv - 1 || win->ne[1] != C || win->ne[2] != 1 || + win->ne[3] != 1 || win->nb[0] != sizeof(float) || win->nb[1] != cc->nb[1] || + dst->type != GGML_TYPE_F32 || dst->data == nullptr || !ggml_is_contiguous(dst) || + ggml_nelements(dst) != (d_conv - 1) * C) { + return false; + } + const int i_win = ggml_cuda_find_node(cgraph, win, ic + 1, j); + if (i_win < 0 || ggml_node_get_use_count(cgraph, i_win) != 1) { + return false; + } + // one whole cache row (the row the state is read from, or another one), not written by an earlier CPY + const ggml_cuda_byte_range r_dst = ggml_cuda_tensor_range(dst); + if (r_dst.beg < r_cache.beg || r_dst.end > r_cache.end || (r_dst.beg - r_cache.beg) % cache->nb[1] != 0 || + r_dst.end - r_dst.beg != (uintptr_t) cache->ne[0] * sizeof(float)) { + return false; + } + for (int k = 0; k < n_cpy; ++k) { + if (ggml_cuda_ranges_overlap(r_dst, r_dsts[k])) { + return false; + } + } + r_dsts[n_cpy] = r_dst; + p.i_cpy[n_cpy] = j; + p.args.snap_offs[n_cpy] = (int64_t) ((r_dst.beg - r_cache.beg) / sizeof(float)); + p.args.snap_pos [n_cpy] = (int) (win->view_offs / sizeof(float)); + ++n_cpy; + i_last_cpy = j; + } + if (n_cpy == 0 || ggml_node_get_use_count(cgraph, ic) != n_cpy + 1) { return false; } // SSM_CONV(concat) followed by its SILU int i_conv = -1; - for (int j = i_cpy + 1; j < n_nodes - 1 && j < i_cpy + 64; ++j) { + for (int j = i_last_cpy + 1; j < n_nodes - 1 && j < i_last_cpy + 64; ++j) { if (cgraph->nodes[j]->op == GGML_OP_SSM_CONV && cgraph->nodes[j]->src[0] == cc) { i_conv = j; break; @@ -3204,8 +3232,8 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml const ggml_tensor * conv = cgraph->nodes[i_conv]; const ggml_tensor * silu = cgraph->nodes[i_conv + 1]; const ggml_tensor * w = conv->src[1]; - if (conv->type != GGML_TYPE_F32 || (conv->flags & GGML_TENSOR_FLAG_OUTPUT) || conv->ne[0] != C || conv->ne[1] != 1 || - ggml_node_get_use_count(cgraph, i_conv) != 1 || + if (conv->type != GGML_TYPE_F32 || (conv->flags & GGML_TENSOR_FLAG_OUTPUT) || conv->ne[0] != C || + conv->ne[1] != n_t || conv->ne[2] != 1 || ggml_node_get_use_count(cgraph, i_conv) != 1 || w->type != GGML_TYPE_F32 || w->data == nullptr || w->ne[0] != d_conv || w->ne[1] != C || w->nb[0] != sizeof(float) || w->nb[1] % sizeof(float) != 0 || silu->op != GGML_OP_UNARY || ggml_get_unary_op(silu) != GGML_UNARY_OP_SILU || silu->src[0] != conv || @@ -3214,7 +3242,7 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml return false; } - // L2_NORM over rows of a leading view of the SILU output (the q and k heads) + // L2_NORM over rows of a leading view of each token's SILU output (the q and k heads) int i_l2 = i_conv + 2; while (i_l2 < n_nodes && ggml_cuda_is_view_or_noop(cgraph->nodes[i_l2])) { ++i_l2; @@ -3227,34 +3255,39 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml if (l2->op != GGML_OP_L2_NORM || l2->type != GGML_TYPE_F32 || (l2->flags & GGML_TENSOR_FLAG_OUTPUT) || !ggml_is_contiguous(l2) || qkv->op != GGML_OP_VIEW || qkv->view_src != silu || qkv->view_offs != 0 || !ggml_are_same_shape(qkv, l2) || qkv->nb[0] != sizeof(float) || - qkv->nb[1] != qkv->ne[0] * sizeof(float) || qkv->ne[2] != 1 || qkv->ne[3] != 1) { + qkv->nb[1] != qkv->ne[0] * sizeof(float) || qkv->ne[2] != n_t || qkv->ne[3] != 1 || + (n_t > 1 && qkv->nb[2] != silu->nb[1])) { return false; } const int64_t row_len = qkv->ne[0]; const int64_t n_l2 = qkv->ne[0] * qkv->ne[1]; float eps; memcpy(&eps, l2->op_params, sizeof(float)); - if (row_len % WARP_SIZE != 0 || row_len >= 1024 || C % row_len != 0 || n_l2 > C || eps < 0.0f) { + if (row_len % WARP_SIZE != 0 || row_len >= 1024 || C % row_len != 0 || n_l2 > C || eps < 0.0f || + ((d_conv - 1 + 2*n_t)*row_len + n_t)*sizeof(float) > 48*1024) { return false; } - // memory: outputs against inputs (the SILU may overwrite the new token in place, element for element) - const ggml_cuda_byte_range r_x = { (uintptr_t) x->data, (uintptr_t) x->data + (uintptr_t) ((C - 1) * x->nb[1] + sizeof(float)) }; + // memory: outputs against inputs (the SILU may overwrite the new tokens in place, element for element) + const uintptr_t x_span = (uintptr_t) ((n_t - 1) * x_tok_stride * sizeof(float) + (C - 1) * x->nb[1] + sizeof(float)); + const ggml_cuda_byte_range r_x = { (uintptr_t) x->data, (uintptr_t) x->data + x_span }; const ggml_cuda_byte_range r_w = ggml_cuda_tensor_range(w); const ggml_cuda_byte_range r_ids = ggml_cuda_tensor_range(ids); const ggml_cuda_byte_range r_y = ggml_cuda_tensor_range(silu); const ggml_cuda_byte_range r_l2 = ggml_cuda_tensor_range(l2); - const ggml_cuda_byte_range r_cc = ggml_cuda_tensor_range(cc); - const ggml_cuda_byte_range r_conv = ggml_cuda_tensor_range(conv); - const bool y_in_place = silu->data == x->data && x->nb[1] == sizeof(float); + const bool y_in_place = silu->data == x->data && x_stride == 1 && (n_t == 1 || x_tok_stride == C); if ((!y_in_place && ggml_cuda_ranges_overlap(r_y, r_x)) || ggml_cuda_ranges_overlap(r_y, r_cache) || ggml_cuda_ranges_overlap(r_y, r_w) || ggml_cuda_ranges_overlap(r_y, r_ids) || ggml_cuda_ranges_overlap(r_y, r_l2) || ggml_cuda_ranges_overlap(r_l2, r_x) || ggml_cuda_ranges_overlap(r_l2, r_cache) || - ggml_cuda_ranges_overlap(r_l2, r_w) || ggml_cuda_ranges_overlap(r_l2, r_ids) || - ggml_cuda_ranges_overlap(r_dst, r_x) || ggml_cuda_ranges_overlap(r_dst, r_w) || - ggml_cuda_ranges_overlap(r_dst, r_ids)) { + ggml_cuda_ranges_overlap(r_l2, r_w) || ggml_cuda_ranges_overlap(r_l2, r_ids)) { return false; } + for (int k = 0; k < n_cpy; ++k) { + if (ggml_cuda_ranges_overlap(r_dsts[k], r_x) || ggml_cuda_ranges_overlap(r_dsts[k], r_w) || + ggml_cuda_ranges_overlap(r_dsts[k], r_ids)) { + return false; + } + } // between the GET_ROWS and the CONCAT nothing may write the cache or the ids (the read moves to the CONCAT) for (int j = i_gr + 1; j < ic; ++j) { @@ -3267,14 +3300,18 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml return false; } } - // nodes evaluated after the fused launch instead of before parts of it: no shared memory with the chain + // nodes evaluated after the fused launch instead of before parts of it. The SILU and L2_NORM outputs are written + // early (at the CONCAT instead of at their own nodes), so no node in between may write or read their memory (the + // allocator may lend it out until those nodes run); the cache rows, weights and ids must not be written either, + // and the cache must not be read. The new tokens are read at the CONCAT position as before, and the CONCAT and + // SSM_CONV buffers are never written by the fused launch, so nodes in between may reuse that memory freely. for (int j = ic + 1; j < i_l2; ++j) { const ggml_tensor * t = cgraph->nodes[j]; - if (j == i_cpy || j == i_conv || j == i_conv + 1 || ggml_cuda_is_view_or_noop(t)) { - continue; + if (j <= i_last_cpy || j == i_conv || j == i_conv + 1 || ggml_cuda_is_view_or_noop(t)) { + continue; // (ic, i_last_cpy] holds only the CPYs and their views } const ggml_cuda_byte_range r = ggml_cuda_tensor_range(t); - for (const ggml_cuda_byte_range & o : { r_x, r_cache, r_w, r_ids, r_y, r_l2, r_cc, r_conv }) { + for (const ggml_cuda_byte_range & o : { r_y, r_l2, r_cache, r_w, r_ids }) { if (ggml_cuda_ranges_overlap(r, o)) { return false; } @@ -3284,7 +3321,7 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml continue; } const ggml_cuda_byte_range rs = ggml_cuda_tensor_range(t->src[s]); - for (const ggml_cuda_byte_range & o : { r_cache, r_y, r_l2, r_cc, r_conv }) { + for (const ggml_cuda_byte_range & o : { r_y, r_l2, r_cache }) { if (ggml_cuda_ranges_overlap(rs, o)) { return false; } @@ -3294,35 +3331,43 @@ static bool ggml_cuda_gdn_conv_match(ggml_backend_cuda_context & ctx, const ggml // concurrent streams would run other nodes alongside for (const auto & [fork_node, event] : ctx.stream_context().concurrent_events) { GGML_UNUSED(fork_node); - for (const ggml_tensor * t : { gr, cc, cpy, conv, silu, l2 }) { + for (const ggml_tensor * t : { gr, cc, conv, silu, l2 }) { if (event.stream_mapping.find(t) != event.stream_mapping.end()) { return false; } } + for (int k = 0; k < n_cpy; ++k) { + if (event.stream_mapping.find(cgraph->nodes[p.i_cpy[k]]) != event.stream_mapping.end()) { + return false; + } + } } p.i_gather = i_gr; - p.i_cpy = i_cpy; + p.n_cpy = n_cpy; p.i_conv = i_conv; p.i_l2 = i_l2; p.gather = gr; ggml_cuda_ssm_conv_gdn_fused_args & a = p.args; - a.cache = (const float *) cache->data; - a.ids = (const int32_t *) ids->data; - a.row_stride = (int64_t) (cache->nb[1] / sizeof(float)); - a.x = (const float *) x->data; - a.x_stride = (int64_t) (x->nb[1] / sizeof(float)); - a.w = (const float *) w->data; - a.w_stride = (int) (w->nb[1] / sizeof(float)); - a.cache_dst = (float *) dst->data; - a.y = (float *) silu->data; - a.l2 = (float *) l2->data; - a.n_l2 = (int) n_l2; - a.n_channels = (int) C; - a.row_len = (int) row_len; - a.d_conv = (int) d_conv; - a.eps = eps; + a.cache = (const float *) cache->data; + a.ids = (const int32_t *) ids->data; + a.row_stride = (int64_t) (cache->nb[1] / sizeof(float)); + a.x = (const float *) x->data; + a.x_stride = x_stride; + a.x_tok_stride = x_tok_stride; + a.w = (const float *) w->data; + a.w_stride = (int) (w->nb[1] / sizeof(float)); + a.cache_dst = (float *) cache->data; + a.n_snap = n_cpy; + a.y = (float *) silu->data; + a.l2 = (float *) l2->data; + a.n_l2 = (int) n_l2; + a.n_channels = (int) C; + a.n_t = (int) n_t; + a.row_len = (int) row_len; + a.d_conv = (int) d_conv; + a.eps = eps; return true; } @@ -3362,7 +3407,9 @@ static bool ggml_cuda_try_gdn_conv_fused(ggml_backend_cuda_context & ctx, const return false; } ggml_cuda_ssm_conv_gdn_fused(ctx, p.args); - reg.done.insert(cgraph->nodes[p.i_cpy]); + for (int k = 0; k < p.n_cpy; ++k) { + reg.done.insert(cgraph->nodes[p.i_cpy[k]]); + } reg.done.insert(cgraph->nodes[p.i_conv]); reg.done.insert(cgraph->nodes[p.i_conv + 1]); reg.done.insert(cgraph->nodes[p.i_l2]); diff --git a/ggml/src/ggml-cuda/mmvf.cu b/ggml/src/ggml-cuda/mmvf.cu index 1c505fa39344..2486ff28e579 100644 --- a/ggml/src/ggml-cuda/mmvf.cu +++ b/ggml/src/ggml-cuda/mmvf.cu @@ -883,8 +883,10 @@ bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, const int64_t * src0 // Two independent mat-vecs that share src1 and the weight type, shape and strides, as one launch of the same kernel: // the second weight becomes channel 1 (its offset from the first is the channel stride, src1 is broadcast with channel -// stride 0). Every row runs the same template instance with the same block size as a single launch would, so the -// results are bit-identical. Returns false (nothing launched) when the pair does not qualify. +// stride 0). With 1..MMVF_MAX_BATCH_SIZE src1 columns (decode and speculative verify widths) every row and column runs +// the same template instance (ncols_dst = the column count) with the same block size as ggml_cuda_mul_mat_vec_f would +// for either product alone, so the results are bit-identical. Returns false (nothing launched) when the pair does not +// qualify. bool ggml_cuda_mul_mat_vec_f_pair(ggml_backend_cuda_context & ctx, ggml_tensor * dst0, ggml_tensor * dst1) { const ggml_tensor * a = dst0->src[0]; const ggml_tensor * b = dst1->src[0]; @@ -892,9 +894,11 @@ bool ggml_cuda_mul_mat_vec_f_pair(ggml_backend_cuda_context & ctx, ggml_tensor * if (dst1->src[1] != src1 || a->type != b->type || !ggml_are_same_shape(a, b) || !ggml_are_same_stride(a, b) || (a->type != GGML_TYPE_F32 && a->type != GGML_TYPE_F16 && a->type != GGML_TYPE_BF16) || src1->type != GGML_TYPE_F32 || dst0->type != GGML_TYPE_F32 || dst1->type != GGML_TYPE_F32 || - a->ne[2] != 1 || a->ne[3] != 1 || src1->ne[1] != 1 || src1->ne[2] != 1 || src1->ne[3] != 1 || - src1->nb[0] != sizeof(float) || !ggml_are_same_shape(dst0, dst1) || !ggml_is_contiguous(dst0) || - !ggml_is_contiguous(dst1) || dst0->ne[1] != 1 || dst0->ne[2] != 1 || dst0->ne[3] != 1 || + a->ne[2] != 1 || a->ne[3] != 1 || src1->ne[1] < 1 || src1->ne[1] > MMVF_MAX_BATCH_SIZE || + src1->ne[2] != 1 || src1->ne[3] != 1 || src1->nb[0] != sizeof(float) || + src1->nb[1] % (2*sizeof(float)) != 0 || + !ggml_are_same_shape(dst0, dst1) || !ggml_is_contiguous(dst0) || !ggml_is_contiguous(dst1) || + dst0->ne[1] != src1->ne[1] || dst0->ne[2] != 1 || dst0->ne[3] != 1 || a->nb[0] != ggml_type_size(a->type)) { return false; } @@ -914,6 +918,7 @@ bool ggml_cuda_mul_mat_vec_f_pair(ggml_backend_cuda_context & ctx, ggml_tensor * const int64_t ne00 = a->ne[0]; const int64_t ne01 = a->ne[1]; + const int64_t ne11 = src1->ne[1]; const int64_t s01 = a->nb[1] / ts; const int64_t s11 = src1->nb[1] / sizeof(float); const int64_t s1 = dst0->nb[1] / sizeof(float); @@ -927,16 +932,17 @@ bool ggml_cuda_mul_mat_vec_f_pair(ggml_backend_cuda_context & ctx, ggml_tensor * switch (a->type) { case GGML_TYPE_F32: - mul_mat_vec_f_cuda((const float *) a->data, src1_d, nullptr, fusion, dst_d, ne00, ne01, 1, s01, s11, s1, + mul_mat_vec_f_cuda((const float *) a->data, src1_d, nullptr, fusion, dst_d, ne00, ne01, ne11, s01, s11, s1, 2, 2, 2, dx / (int64_t) ts, 0, dd / (int64_t) sizeof(float), 1, 1, s03, s13, s3, 0, prec, ctx.stream()); break; case GGML_TYPE_F16: - mul_mat_vec_f_cuda((const half *) a->data, src1_d, nullptr, fusion, dst_d, ne00, ne01, 1, s01, s11, s1, + mul_mat_vec_f_cuda((const half *) a->data, src1_d, nullptr, fusion, dst_d, ne00, ne01, ne11, s01, s11, s1, 2, 2, 2, dx / (int64_t) ts, 0, dd / (int64_t) sizeof(float), 1, 1, s03, s13, s3, 0, prec, ctx.stream()); break; case GGML_TYPE_BF16: - mul_mat_vec_f_cuda((const nv_bfloat16 *) a->data, src1_d, nullptr, fusion, dst_d, ne00, ne01, 1, s01, s11, s1, - 2, 2, 2, dx / (int64_t) ts, 0, dd / (int64_t) sizeof(float), 1, 1, s03, s13, s3, 0, prec, ctx.stream()); + mul_mat_vec_f_cuda((const nv_bfloat16 *) a->data, src1_d, nullptr, fusion, dst_d, ne00, ne01, ne11, s01, + s11, s1, 2, 2, 2, dx / (int64_t) ts, 0, dd / (int64_t) sizeof(float), 1, 1, s03, s13, s3, 0, prec, + ctx.stream()); break; default: return false; diff --git a/ggml/src/ggml-cuda/ssm-conv.cu b/ggml/src/ggml-cuda/ssm-conv.cu index d86fe14ee1eb..3a1f3d68a740 100644 --- a/ggml/src/ggml-cuda/ssm-conv.cu +++ b/ggml/src/ggml-cuda/ssm-conv.cu @@ -205,90 +205,123 @@ void ggml_cuda_op_ssm_conv(ggml_backend_cuda_context & ctx, ggml_tensor * dst, g } } -// GDN decode conv chain in one launch (see ggml_cuda_gdn_conv_fusion in ggml-cuda.cu): the conv-state row gather, the -// concat with the new token, the conv-state write-back, the depthwise conv + SiLU and the L2 norm of the q/k heads. -// One thread per channel, one block per L2 row (head). Every value is computed with the same expression, in the same -// order, as the unfused kernels: the conv sum as ssm_conv_f32 (token 0, no bias), the L2 norm as l2_norm_f32 with -// WARP_SIZE threads (per-lane strided partial sums, then the warp butterfly), so the outputs are bit-identical. +// GDN conv chain in one launch (see ggml_cuda_gdn_conv_match in ggml-cuda.cu): the conv-state row gather, the concat +// with the n_t new tokens, the conv-state write-backs (one per rollback snapshot), the depthwise conv + SiLU and the L2 +// norm of the q/k heads of every token. One thread per channel, one block per L2 row (head). Every value is computed +// with the same expression, in the same order, as the unfused kernels: the conv sum as ssm_conv_f32 (no bias), the L2 +// norm as l2_norm_f32 with WARP_SIZE threads (per-lane strided partial sums, then the warp butterfly; one warp per +// token here), so the outputs are bit-identical. +struct ggml_cuda_gdn_conv_snaps { + int64_t offs[GGML_CUDA_GDN_CONV_MAX_SNAP]; + int pos [GGML_CUDA_GDN_CONV_MAX_SNAP]; +}; + template static __global__ void ssm_conv_gdn_fused_f32( - const float * __restrict__ cache, const int32_t * __restrict__ ids, const int64_t row_stride, - const float * x, const int64_t x_stride, + const float * cache, const int32_t * __restrict__ ids, const int64_t row_stride, + const float * x, const int64_t x_stride, const int64_t x_tok_stride, const float * __restrict__ w, const int w_stride, - float * __restrict__ cache_dst, - float * y, - float * __restrict__ l2, const int n_l2, const float eps) { + float * cache_dst, const ggml_cuda_gdn_conv_snaps snaps, const int n_snap, + float * y, float * __restrict__ l2, const int n_l2, const int n_channels, const int n_t, const float eps) { ggml_cuda_pdl_lc(); - extern __shared__ float smem[]; // [blockDim.x] silu outputs of this row, then the row scale + // per thread (strided by blockDim.x): the d_conv - 1 + n_t conv inputs of its channel, then its n_t silu outputs; + // then the n_t row scales of the block + extern __shared__ float smem[]; const int tid = threadIdx.x; - const int c = blockIdx.x*blockDim.x + tid; + const int bd = blockDim.x; + const int c = blockIdx.x*bd + tid; + + float * win = smem; + float * yv = win + (d_conv - 1 + n_t)*bd; + float * sc = yv + n_t*bd; ggml_cuda_pdl_sync(); + // every read of this channel's state and tokens comes first: a conv-state row written below may be the row read + // here, and y may be x itself (element for element); both only at this thread's own channel const float * s = cache + ids[0]*row_stride + (int64_t) c*(d_conv - 1); - - float xw[d_conv]; - float ww[d_conv]; #pragma unroll for (int j = 0; j < d_conv - 1; ++j) { - xw[j] = s[j]; + win[j*bd + tid] = s[j]; + } + for (int t = 0; t < n_t; ++t) { + win[(d_conv - 1 + t)*bd + tid] = x[t*x_tok_stride + c*x_stride]; } - xw[d_conv - 1] = x[c*x_stride]; + float ww[d_conv]; #pragma unroll for (int j = 0; j < d_conv; ++j) { ww[j] = w[c*w_stride + j]; } - // conv-state write-back: the last d_conv - 1 inputs (the same row when ids[0] is the destination row: this thread - // read its values above) + // conv-state write-backs: d_conv - 1 consecutive conv inputs per row + for (int k = 0; k < n_snap; ++k) { + float * dk = cache_dst + snaps.offs[k] + (int64_t) c*(d_conv - 1); + const int p = snaps.pos[k]; #pragma unroll - for (int j = 0; j < d_conv - 1; ++j) { - cache_dst[(int64_t) c*(d_conv - 1) + j] = xw[j + 1]; + for (int j = 0; j < d_conv - 1; ++j) { + dk[j] = win[(p + j)*bd + tid]; + } } const float b = 0.0f; // ssm_conv_f32 adds its (absent) bias as 0.0f - float sumf = 0.0f; + for (int t = 0; t < n_t; ++t) { + float sumf = 0.0f; #pragma unroll - for (int j = 0; j < d_conv; ++j) { - sumf += xw[j] * ww[j]; + for (int j = 0; j < d_conv; ++j) { + sumf += win[(t + j)*bd + tid] * ww[j]; + } + sumf += b; + const float v = ggml_cuda_op_silu_single(sumf); + y[(int64_t) t*n_channels + c] = v; + yv[t*bd + tid] = v; } - sumf += b; - const float v = ggml_cuda_op_silu_single(sumf); - y[c] = v; - if (blockIdx.x*blockDim.x >= n_l2) { + if (blockIdx.x*bd >= n_l2) { return; // a v channel block: no L2 norm (block-uniform) } - smem[tid] = v; __syncthreads(); - if (tid < WARP_SIZE) { + const int lane = tid % WARP_SIZE; + const int n_warps = bd / WARP_SIZE; + for (int t = tid / WARP_SIZE; t < n_t; t += n_warps) { float tmp = 0.0f; - for (int col = tid; col < (int) blockDim.x; col += WARP_SIZE) { - const float xi = smem[col]; + for (int col = lane; col < bd; col += WARP_SIZE) { + const float xi = yv[t*bd + col]; tmp += xi * xi; } tmp = warp_reduce_sum(tmp); - if (tid == 0) { - smem[blockDim.x] = rsqrtf(fmaxf(tmp, eps * eps)); + if (lane == 0) { + sc[t] = rsqrtf(fmaxf(tmp, eps * eps)); } } __syncthreads(); - l2[c] = smem[blockDim.x] * v; + for (int t = 0; t < n_t; ++t) { + l2[(int64_t) t*n_l2 + c] = sc[t] * yv[t*bd + tid]; + } } void ggml_cuda_ssm_conv_gdn_fused(ggml_backend_cuda_context & ctx, const ggml_cuda_ssm_conv_gdn_fused_args & a) { GGML_ASSERT(a.row_len % WARP_SIZE == 0 && a.row_len < 1024 && a.n_channels % a.row_len == 0 && a.n_l2 % a.row_len == 0); + GGML_ASSERT(a.n_t >= 1 && a.n_t <= GGML_CUDA_GDN_CONV_MAX_TOKENS); + GGML_ASSERT(a.n_snap >= 1 && a.n_snap <= GGML_CUDA_GDN_CONV_MAX_SNAP); const dim3 blocks(a.n_channels / a.row_len, 1, 1); const dim3 threads(a.row_len, 1, 1); - const size_t smem = (a.row_len + 1) * sizeof(float); + const size_t smem = ((size_t) (a.d_conv - 1 + 2*a.n_t)*a.row_len + a.n_t) * sizeof(float); + GGML_ASSERT(smem <= 48*1024); const ggml_cuda_kernel_launch_params launch_params(blocks, threads, smem, ctx.stream()); + ggml_cuda_gdn_conv_snaps snaps; + for (int k = 0; k < GGML_CUDA_GDN_CONV_MAX_SNAP; ++k) { + snaps.offs[k] = k < a.n_snap ? a.snap_offs[k] : 0; + snaps.pos[k] = k < a.n_snap ? a.snap_pos[k] : 0; + } + auto launch = [&](auto NC) { constexpr int kNC = decltype(NC)::value; ggml_cuda_kernel_launch(ssm_conv_gdn_fused_f32, launch_params, a.cache, a.ids, a.row_stride, a.x, a.x_stride, - a.w, a.w_stride, a.cache_dst, a.y, a.l2, a.n_l2, a.eps); + a.x_tok_stride, a.w, a.w_stride, a.cache_dst, snaps, a.n_snap, a.y, a.l2, a.n_l2, + a.n_channels, a.n_t, a.eps); }; switch (a.d_conv) { case 3: launch(std::integral_constant{}); break; diff --git a/ggml/src/ggml-cuda/ssm-conv.cuh b/ggml/src/ggml-cuda/ssm-conv.cuh index d9da8e3af153..6c0f8567a545 100644 --- a/ggml/src/ggml-cuda/ssm-conv.cuh +++ b/ggml/src/ggml-cuda/ssm-conv.cuh @@ -2,20 +2,30 @@ void ggml_cuda_op_ssm_conv(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * bias_add_node = nullptr, ggml_tensor * silu_dst = nullptr); -// fused GDN decode conv chain (one token, one sequence); see ggml_cuda_gdn_conv_fusion in ggml-cuda.cu +// fused GDN conv chain for one sequence and 1..GGML_CUDA_GDN_CONV_MAX_TOKENS tokens, with up to +// GGML_CUDA_GDN_CONV_MAX_SNAP conv-state writes (one at a plain step, 1 + n_rs_seq rollback snapshots at a verify +// step); see ggml_cuda_gdn_conv_match in ggml-cuda.cu +#define GGML_CUDA_GDN_CONV_MAX_TOKENS 16 +#define GGML_CUDA_GDN_CONV_MAX_SNAP 16 + struct ggml_cuda_ssm_conv_gdn_fused_args { const float * cache = nullptr; // conv-state cache, rows of row_stride floats const int32_t * ids = nullptr; // ids[0]: the row to read int64_t row_stride = 0; - const float * x = nullptr; // new token, channel c at x[c*x_stride] + const float * x = nullptr; // new tokens, channel c of token t at x[t*x_tok_stride + c*x_stride] int64_t x_stride = 0; + int64_t x_tok_stride = 0; const float * w = nullptr; // conv weights, channel c at w[c*w_stride + j] int w_stride = 0; - float * cache_dst = nullptr; // conv-state row written back, [n_channels*(d_conv - 1)] - float * y = nullptr; // silu(conv), [n_channels] - float * l2 = nullptr; // l2-normalised first n_l2 channels, rows of row_len + float * cache_dst = nullptr; // base of the conv-state rows written back + int n_snap = 0; // conv-state rows written, each [n_channels*(d_conv - 1)] + int64_t snap_offs[GGML_CUDA_GDN_CONV_MAX_SNAP] = {}; // row k starts at cache_dst + snap_offs[k] (floats) + int snap_pos [GGML_CUDA_GDN_CONV_MAX_SNAP] = {}; // holding conv-input columns snap_pos[k] + [0, d-1) + float * y = nullptr; // silu(conv), [n_channels, n_t] + float * l2 = nullptr; // l2-normalised first n_l2 channels per token, rows of row_len, [n_l2, n_t] int n_l2 = 0; int n_channels = 0; + int n_t = 1; int row_len = 0; int d_conv = 0; float eps = 0.0f; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 048c12f941d7..385689dad32f 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -4136,17 +4136,20 @@ struct test_ssm_conv_bias_silu : public test_case { } }; -// The decode conv chain of a gated delta net layer (one token, one sequence), as the hybrid models build it: -// GET_ROWS (conv state row) -> CONCAT with the new token -> CPY of the last window back into a cache row, -// SSM_CONV -> SILU -> L2_NORM of the leading q/k heads. The CUDA backend can run it as one launch -// (GGML_CUDA_GDN_FUSION=1); the L2_NORM, the SILU output and the written cache row are all checked. +// The conv chain of a gated delta net layer (n_t tokens of one sequence), as the hybrid models build it: +// GET_ROWS (conv state row) -> CONCAT with the new tokens -> CPY of the last window back into a cache row (n_rs == 0), +// or one CPY per rollback snapshot slot (n_rs > 0: 1 + n_rs slots, as llm_build_delta_net_base::build_conv_state writes +// them), SSM_CONV -> SILU -> L2_NORM of the leading q/k heads of every token. The CUDA backend can run it as one +// launch (GGML_CUDA_GDN_FUSION=1); the L2_NORM, the SILU output and the written cache rows are all checked. struct test_gdn_conv_chain : public test_case { const int64_t d_conv; const int64_t head_dim; // L2 row length const int64_t n_qk_rows; // q + k heads const int64_t n_ch; // all channels (q, k and v) const int src_row; // cache row read - const int dst_row; // cache row written + const int dst_row; // cache row written (n_rs == 0), or the slot-0 row (kv head) of the snapshot planes + const int64_t n_t; // tokens + const int64_t n_rs; // rollback snapshots (n_rs_seq) std::vector outs; std::string op_desc(ggml_tensor * t) override { @@ -4155,7 +4158,7 @@ struct test_gdn_conv_chain : public test_case { } std::string vars() override { - return VARS_TO_STR6(d_conv, head_dim, n_qk_rows, n_ch, src_row, dst_row); + return VARS_TO_STR8(d_conv, head_dim, n_qk_rows, n_ch, src_row, dst_row, n_t, n_rs); } bool run_whole_graph() override { return true; } @@ -4163,16 +4166,18 @@ struct test_gdn_conv_chain : public test_case { std::vector fusion_test_nodes() override { return outs; } test_gdn_conv_chain(int64_t d_conv = 4, int64_t head_dim = 128, int64_t n_qk_rows = 32, int64_t n_ch = 5120, - int src_row = 1, int dst_row = 1) - : d_conv(d_conv), head_dim(head_dim), n_qk_rows(n_qk_rows), n_ch(n_ch), src_row(src_row), dst_row(dst_row) {} + int src_row = 1, int dst_row = 1, int64_t n_t = 1, int64_t n_rs = 0) + : d_conv(d_conv), head_dim(head_dim), n_qk_rows(n_qk_rows), n_ch(n_ch), src_row(src_row), dst_row(dst_row), + n_t(n_t), n_rs(n_rs) {} ggml_tensor * build_graph(ggml_context * ctx) override { - const int64_t n_rows = 3; + const int64_t mem_size = 3; // rows per snapshot plane + const int64_t n_rows = n_rs == 0 ? 3 : mem_size * (1 + n_rs); ggml_tensor * cache = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, (d_conv - 1) * n_ch, n_rows); ggml_set_name(cache, "cache"); ggml_tensor * ids = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, 1); ggml_set_name(ids, "ids"); - ggml_tensor * x = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_ch); + ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, n_ch, n_t); ggml_set_name(x, "x"); ggml_tensor * w = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, d_conv, n_ch); ggml_set_name(w, "w"); @@ -4180,24 +4185,32 @@ struct test_gdn_conv_chain : public test_case { ggml_tensor * state = ggml_get_rows(ctx, cache, ids); state = ggml_reshape_3d(ctx, state, d_conv - 1, n_ch, 1); x = ggml_scale(ctx, x, 1.0f); // a computed tensor, as the projection is in the model - ggml_tensor * xt = ggml_transpose(ctx, ggml_reshape_3d(ctx, x, n_ch, 1, 1)); + ggml_tensor * xt = ggml_transpose(ctx, ggml_reshape_3d(ctx, x, n_ch, n_t, 1)); ggml_tensor * conv_input = ggml_concat(ctx, state, xt, 0); - ggml_tensor * last = ggml_view_3d(ctx, conv_input, d_conv - 1, n_ch, 1, conv_input->nb[1], conv_input->nb[2], - ggml_row_size(conv_input->type, 1)); - ggml_tensor * upd = ggml_view_2d(ctx, cache, (d_conv - 1) * n_ch, 1, cache->nb[1], dst_row * cache->nb[1]); - ggml_tensor * cpy = ggml_cpy(ctx, last, upd); - if (gf) { - ggml_build_forward_expand(gf, cpy); + const int64_t K = n_rs + 1; + std::vector cpys; + for (int64_t t = n_rs == 0 ? K : 1; t <= K; ++t) { + const int64_t s_idx = n_rs == 0 ? n_t : std::max(0, n_t - K + t); + const int64_t row = n_rs == 0 ? dst_row : (K - t) * mem_size + dst_row; + ggml_tensor * last = ggml_view_3d(ctx, conv_input, d_conv - 1, n_ch, 1, conv_input->nb[1], + conv_input->nb[2], ggml_row_size(conv_input->type, s_idx)); + ggml_tensor * upd = ggml_view_2d(ctx, cache, (d_conv - 1) * n_ch, 1, cache->nb[1], row * cache->nb[1]); + ggml_tensor * cpy = ggml_cpy(ctx, last, upd); + if (gf) { + ggml_build_forward_expand(gf, cpy); + } + cpys.push_back(cpy); } ggml_tensor * y = ggml_silu(ctx, ggml_ssm_conv(ctx, conv_input, w)); - ggml_tensor * qk = ggml_view_4d(ctx, y, head_dim, n_qk_rows, 1, 1, ggml_row_size(y->type, head_dim), y->nb[1], - y->nb[2], 0); + ggml_tensor * qk = ggml_view_4d(ctx, y, head_dim, n_qk_rows, n_t, 1, ggml_row_size(y->type, head_dim), y->nb[1], + y->nb[1] * n_t, 0); ggml_tensor * out = ggml_l2_norm(ctx, qk, 1e-6f); ggml_set_name(out, "out"); - outs = { out, y, cpy }; + outs = { out, y }; + outs.insert(outs.end(), cpys.begin(), cpys.end()); return out; } @@ -4222,6 +4235,7 @@ struct test_mul_mat_vec_pair : public test_case { const ggml_type type; const int64_t k; const int64_t m; + const int64_t n; // columns (tokens) std::vector outs; std::string op_desc(ggml_tensor * t) override { @@ -4230,7 +4244,7 @@ struct test_mul_mat_vec_pair : public test_case { } std::string vars() override { - return VARS_TO_STR3(type, k, m); + return VARS_TO_STR4(type, k, m, n); } bool run_whole_graph() override { return true; } @@ -4239,10 +4253,11 @@ struct test_mul_mat_vec_pair : public test_case { double max_nmse_err() override { return 5e-4; } - test_mul_mat_vec_pair(ggml_type type = GGML_TYPE_F32, int64_t k = 1024, int64_t m = 8) : type(type), k(k), m(m) {} + test_mul_mat_vec_pair(ggml_type type = GGML_TYPE_F32, int64_t k = 1024, int64_t m = 8, int64_t n = 1) + : type(type), k(k), m(m), n(n) {} ggml_tensor * build_graph(ggml_context * ctx) override { - ggml_tensor * x = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, k); + ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, k, n); ggml_set_name(x, "x"); ggml_tensor * wa = ggml_new_tensor_2d(ctx, type, k, m); ggml_set_name(wa, "wa"); @@ -10220,10 +10235,25 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_gdn_conv_chain(4, 256, 8, 4096 + 2048, 1, 0)); test_cases.emplace_back(new test_gdn_conv_chain(3, 64, 4, 512, 1, 1)); test_cases.emplace_back(new test_gdn_conv_chain(5, 32, 6, 256, 0, 1)); + // speculative verify widths: n_t tokens, plain (n_rs = 0) and with 1 + n_rs rollback snapshot rows, including + // n_t < 1 + n_rs (short drafts: several slots get the pre-batch window), the read row among the written ones + for (int64_t n_t : { 1, 2, 3, 4, 5, 8, 9, 16 }) { + test_cases.emplace_back(new test_gdn_conv_chain(4, 128, 32, 4096 + 1024, 1, 1, n_t, 0)); + test_cases.emplace_back(new test_gdn_conv_chain(4, 128, 32, 4096 + 1024, 1, 1, n_t, 3)); + test_cases.emplace_back(new test_gdn_conv_chain(4, 128, 32, 4096 + 1024, 4, 1, n_t, 7)); + } + test_cases.emplace_back(new test_gdn_conv_chain(4, 128, 16, 4096, 0, 2, 16, 15)); + test_cases.emplace_back(new test_gdn_conv_chain(4, 256, 8, 4096 + 2048, 2, 0, 4, 3)); + test_cases.emplace_back(new test_gdn_conv_chain(3, 64, 4, 512, 1, 1, 3, 2)); + test_cases.emplace_back(new test_gdn_conv_chain(5, 32, 6, 256, 0, 1, 7, 7)); for (ggml_type t : { GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16 }) { test_cases.emplace_back(new test_mul_mat_vec_pair(t, 1024, 8)); test_cases.emplace_back(new test_mul_mat_vec_pair(t, 3072, 24)); test_cases.emplace_back(new test_mul_mat_vec_pair(t, 130, 3)); + for (int64_t n : { 2, 3, 4, 5, 8, 9 }) { + test_cases.emplace_back(new test_mul_mat_vec_pair(t, 1024, 8, n)); + test_cases.emplace_back(new test_mul_mat_vec_pair(t, 4096, 32, n)); + } } test_cases.emplace_back(new test_ssm_scan(GGML_TYPE_F32, 16, 1, 1024, 1, 32, 4)); // Mamba-1