diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index fa9d27046b5e..e027325cd0a6 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -10477,34 +10477,40 @@ static void ggml_compute_forward_gated_delta_net_one_chunk( const float beta_val = *(const float *)((const char *)src_beta->data + iv3 * nbb3 + t * nbb2 + iv1 * nbb1); const float * g_d = (const float *)((const char *)src_g->data + iv3 * nbg3 + t * nbg2 + iv1 * nbg1); + // state is stored transposed: s_out[j*S_v + i] = S[i][j] + // so row j of s_out = column j of S (contiguous access) + if (kda) { + // precompute exp(g) into delta scratch (reused below) for (int64_t i = 0; i < S_v; ++i) { - ggml_vec_scale_f32(S_v, &s_out[i * S_v], expf(g_d[i])); + delta[i] = expf(g_d[i]); + } + // S[i][:] *= exp(g[i]) => for each row j of M: M[j][i] *= exp(g[i]) + for (int64_t j = 0; j < S_v; ++j) { + ggml_vec_mul_f32(S_v, &s_out[j * S_v], &s_out[j * S_v], delta); } } else { ggml_vec_scale_f32(S_v * S_v, s_out, expf(g_d[0])); } - // delta[j] = sum_i S[j][i] * k[i] - memset(delta, 0, S_v * sizeof(float)); - for (int64_t i = 0; i < S_v; ++i) { - ggml_vec_mad_f32(S_v, delta, &s_out[i * S_v], k_d[i]); - } + // delta[j] = sum_i S[i][j] * k[i] = dot(row j of M, k) for (int64_t j = 0; j < S_v; ++j) { - delta[j] = (v_d[j] - delta[j]) * beta_val; + float sum = 0.0f; + ggml_vec_dot_f32(S_v, &sum, 0, &s_out[j * S_v], 0, k_d, 0, 1); + delta[j] = (v_d[j] - sum) * beta_val; } - // outer product: S[j][i] += k[i] * delta[j] - for (int64_t i = 0; i < S_v; ++i) { - ggml_vec_mad_f32(S_v, &s_out[i * S_v], delta, k_d[i]); + // outer product: S[i][j] += k[i] * delta[j] => M[j][i] += delta[j] * k[i] + for (int64_t j = 0; j < S_v; ++j) { + ggml_vec_mad_f32(S_v, &s_out[j * S_v], k_d, delta[j]); } - // attn_out[j] = sum_i S[j][i] * q[i] - memset(attn_data, 0, S_v * sizeof(float)); - for (int64_t i = 0; i < S_v; ++i) { - ggml_vec_mad_f32(S_v, attn_data, &s_out[i * S_v], q_d[i]); + // attn_out[j] = sum_i S[i][j] * q[i] = dot(row j of M, q) + for (int64_t j = 0; j < S_v; ++j) { + float sum = 0.0f; + ggml_vec_dot_f32(S_v, &sum, 0, &s_out[j * S_v], 0, q_d, 0, 1); + attn_data[j] = sum * scale; } - ggml_vec_scale_f32(S_v, attn_data, scale); attn_data += S_v * H; // advance to next token } diff --git a/ggml/src/ggml-cuda/gated_delta_net.cu b/ggml/src/ggml-cuda/gated_delta_net.cu index 5f0fa8e58dfe..706f1097fe16 100644 --- a/ggml/src/ggml-cuda/gated_delta_net.cu +++ b/ggml/src/ggml-cuda/gated_delta_net.cu @@ -1,28 +1,28 @@ #include "gated_delta_net.cuh" -template -__global__ void gated_delta_net_cuda(const float * q, - const float * k, - const float * v, - const float * g, - const float * beta, - const float * curr_state, - float * dst, - int64_t H, - int64_t n_tokens, - int64_t n_seqs, - int64_t sq1, - int64_t sq2, - int64_t sq3, - int64_t sv1, - int64_t sv2, - int64_t sv3, - int64_t sb1, - int64_t sb2, - int64_t sb3, - const uint3 neqk1_magic, - const uint3 rq3_magic, - float scale) { +template +__global__ void gated_delta_net_cuda_no_tail(const float * q, + const float * k, + const float * v, + const float * g, + const float * beta, + const float * curr_state, + float * dst, + int64_t H, + int64_t n_tokens, + int64_t n_seqs, + int64_t sq1, + int64_t sq2, + int64_t sq3, + int64_t sv1, + int64_t sv2, + int64_t sv3, + int64_t sb1, + int64_t sb2, + int64_t sb3, + const uint3 neqk1_magic, + const uint3 rq3_magic, + float scale) { const uint32_t h_idx = blockIdx.x; const uint32_t sequence = blockIdx.y; // each warp owns one column, using warp-level primitives to reduce across rows @@ -45,102 +45,377 @@ __global__ void gated_delta_net_cuda(const float * q, static_assert(S_v % warp_size == 0, "S_v must be a multiple of warp_size"); constexpr int rows_per_lane = (S_v + warp_size - 1) / warp_size; float s_shard[rows_per_lane]; + float qs[n_tokens_per_loop][rows_per_lane]; + float ks[n_tokens_per_loop][rows_per_lane]; + float vs[n_tokens_per_loop]; + float beta_ts[n_tokens_per_loop]; + // non-KDA has one g per column, KDA has one g per row + float g_vals[n_tokens_per_loop][rows_per_lane]; + int idx[rows_per_lane]; + float attn_col_values[n_tokens_per_loop]; + + // state is stored transposed: M[col][i] = S[i][col], row col is contiguous #pragma unroll for (int r = 0; r < rows_per_lane; r++) { const int i = r * warp_size + lane; - s_shard[r] = curr_state[i * S_v + col]; + s_shard[r] = curr_state[col * S_v + i]; + idx[r] = i; } - for (int t = 0; t < n_tokens; t++) { - const float * q_t = q + iq3 * sq3 + t * sq2 + iq1 * sq1; - const float * k_t = k + iq3 * sq3 + t * sq2 + iq1 * sq1; - const float * v_t = v + sequence * sv3 + t * sv2 + h_idx * sv1; + for (int ts = 0; ts < n_tokens; ts += n_tokens_per_loop) { +#pragma unroll + for (int i = 0; i < n_tokens_per_loop; i++) { + const int t = ts + i; + const float * q_t = q + iq3 * sq3 + t * sq2 + iq1 * sq1; + const float * k_t = k + iq3 * sq3 + t * sq2 + iq1 * sq1; + const float * v_t = v + sequence * sv3 + t * sv2 + h_idx * sv1; + const int64_t gb_offset = sequence * sb3 + t * sb2 + h_idx * sb1; + const float * beta_t = beta + gb_offset; + const float * g_t = g + gb_offset * (KDA ? S_v : 1); + + beta_ts[i] = *beta_t; + vs[i] = v_t[col]; + if constexpr (!KDA) { + g_vals[i][0] = expf(*g_t); + } +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + qs[i][r] = q_t[idx[r]]; + ks[i][r] = k_t[idx[r]]; + if constexpr (KDA) { + g_vals[i][r] = expf(g_t[idx[r]]); + } + } + } +#pragma unroll + for (int i = 0; i < n_tokens_per_loop; i++) { + if constexpr (!KDA) { + const float g_val = g_vals[i][0]; - const int64_t gb_offset = sequence * sb3 + t * sb2 + h_idx * sb1; - const float * beta_t = beta + gb_offset; - const float * g_t = g + gb_offset * (KDA ? S_v : 1); + // kv[col] = (S^T @ k)[col] = sum_i S[i][col] * k[i] + float kv_shard = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + kv_shard += s_shard[r] * ks[i][r]; + } + float kv_col = warp_reduce_sum(kv_shard); + + // delta[col] = (v[col] - g * kv[col]) * beta + float delta_col = (vs[i] - g_val * kv_col) * beta_ts[i]; - const float beta_val = *beta_t; + // fused: S[i][col] = g * S[i][col] + k[i] * delta[col] + // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] + float attn_partial = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + s_shard[r] = g_val * s_shard[r] + ks[i][r] * delta_col; + attn_partial += s_shard[r] * qs[i][r]; + } - if constexpr (!KDA) { - const float g_val = expf(*g_t); + float attn_col = warp_reduce_sum(attn_partial); - // kv[col] = (S^T @ k)[col] = sum_i S[i][col] * k[i] - float kv_shard = 0.0f; + attn_col_values[i] = attn_col * scale; + } else { + // kv[col] = sum_i g[i] * S[i][col] * k[i] + float kv_shard = 0.0f; #pragma unroll - for (int r = 0; r < rows_per_lane; r++) { - const int i = r * warp_size + lane; - kv_shard += s_shard[r] * k_t[i]; - } - float kv_col = warp_reduce_sum(kv_shard); + for (int r = 0; r < rows_per_lane; r++) { + kv_shard += g_vals[i][r] * s_shard[r] * ks[i][r]; + } + + float kv_col = warp_reduce_sum(kv_shard); - // delta[col] = (v[col] - g * kv[col]) * beta - float delta_col = (v_t[col] - g_val * kv_col) * beta_val; + // delta[col] = (v[col] - kv[col]) * beta + float delta_col = (vs[i] - kv_col) * beta_ts[i]; - // fused: S[i][col] = g * S[i][col] + k[i] * delta[col] - // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] - float attn_partial = 0.0f; + // fused: S[i][col] = g[i] * S[i][col] + k[i] * delta[col] + // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] + float attn_partial = 0.0f; #pragma unroll - for (int r = 0; r < rows_per_lane; r++) { - const int i = r * warp_size + lane; - s_shard[r] = g_val * s_shard[r] + k_t[i] * delta_col; - attn_partial += s_shard[r] * q_t[i]; - } + for (int r = 0; r < rows_per_lane; r++) { + s_shard[r] = g_vals[i][r] * s_shard[r] + ks[i][r] * delta_col; + attn_partial += s_shard[r] * qs[i][r]; + } - float attn_col = warp_reduce_sum(attn_partial); + float attn_col = warp_reduce_sum(attn_partial); - if (lane == 0) { - attn_data[col] = attn_col * scale; + attn_col_values[i] = attn_col * scale; + } + } + + // Pulling the writes out of the loop is compiler-friendly + if constexpr (warp_size < n_tokens_per_loop) { + for (int i = lane; i < n_tokens_per_loop; i += warp_size) { + if (i < n_tokens_per_loop) { + attn_data[i * S_v * H + col] = attn_col_values[i]; + } } } else { - // kv[col] = sum_i g[i] * S[i][col] * k[i] - float kv_shard = 0.0f; + if (lane < n_tokens_per_loop) { + attn_data[lane * S_v * H + col] = attn_col_values[lane]; + } + } + + attn_data += n_tokens_per_loop * S_v * H; + } + // Write state back to global memory (transposed layout) +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + state[col * S_v + idx[r]] = s_shard[r]; + } +} + +template +__global__ void gated_delta_net_cuda_tail(const float * q, + const float * k, + const float * v, + const float * g, + const float * beta, + const float * curr_state, + float * dst, + int64_t H, + int64_t n_tokens, + int64_t n_seqs, + int64_t sq1, + int64_t sq2, + int64_t sq3, + int64_t sv1, + int64_t sv2, + int64_t sv3, + int64_t sb1, + int64_t sb2, + int64_t sb3, + const uint3 neqk1_magic, + const uint3 rq3_magic, + float scale) { + const uint32_t h_idx = blockIdx.x; + const uint32_t sequence = blockIdx.y; + // each warp owns one column, using warp-level primitives to reduce across rows + const int lane = threadIdx.x; + const int col = blockIdx.z * blockDim.y + threadIdx.y; + + const uint32_t iq1 = fastmodulo(h_idx, neqk1_magic); + const uint32_t iq3 = fastdiv(sequence, rq3_magic); + + const int64_t attn_score_elems = S_v * H * n_tokens * n_seqs; + float * attn_data = dst; + float * state = dst + attn_score_elems; + + const int64_t state_offset = (sequence * H + h_idx) * S_v * S_v; + state += state_offset; + curr_state += state_offset; + attn_data += (sequence * n_tokens * H + h_idx) * S_v; + + constexpr int warp_size = ggml_cuda_get_physical_warp_size() < S_v ? ggml_cuda_get_physical_warp_size() : S_v; + static_assert(S_v % warp_size == 0, "S_v must be a multiple of warp_size"); + constexpr int rows_per_lane = (S_v + warp_size - 1) / warp_size; + float s_shard[rows_per_lane]; + float qs[n_tokens_per_loop][rows_per_lane]; + float ks[n_tokens_per_loop][rows_per_lane]; + float vs[n_tokens_per_loop]; + float beta_ts[n_tokens_per_loop]; + // non-KDA has one g per column, KDA has one g per row + float g_vals[n_tokens_per_loop][rows_per_lane]; + int idx[rows_per_lane]; + // state is stored transposed: M[col][i] = S[i][col], row col is contiguous +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + const int i = r * warp_size + lane; + s_shard[r] = curr_state[col * S_v + i]; + idx[r] = i; + } + + // Iterate over all valid chunks first, then handle the tail case if n_tokens is not divisible by n_tokens_per_loop + const int tokens_to_process = n_tokens - (n_tokens % n_tokens_per_loop); + for (int ts = 0; ts < tokens_to_process; ts += n_tokens_per_loop) { +#pragma unroll + for (int i = 0; i < n_tokens_per_loop; i++) { + const int t = ts + i; + const float * q_t = q + iq3 * sq3 + t * sq2 + iq1 * sq1; + const float * k_t = k + iq3 * sq3 + t * sq2 + iq1 * sq1; + const float * v_t = v + sequence * sv3 + t * sv2 + h_idx * sv1; + const int64_t gb_offset = sequence * sb3 + t * sb2 + h_idx * sb1; + const float * beta_t = beta + gb_offset; + const float * g_t = g + gb_offset * (KDA ? S_v : 1); + + beta_ts[i] = *beta_t; + vs[i] = v_t[col]; + if constexpr (!KDA) { + g_vals[i][0] = expf(*g_t); + } #pragma unroll for (int r = 0; r < rows_per_lane; r++) { - const int i = r * warp_size + lane; - kv_shard += expf(g_t[i]) * s_shard[r] * k_t[i]; + qs[i][r] = q_t[idx[r]]; + ks[i][r] = k_t[idx[r]]; + if constexpr (KDA) { + g_vals[i][r] = expf(g_t[idx[r]]); + } } + } +#pragma unroll + for (int i = 0; i < n_tokens_per_loop; i++) { + if constexpr (!KDA) { + const float g_val = g_vals[i][0]; - float kv_col = warp_reduce_sum(kv_shard); + // kv[col] = (S^T @ k)[col] = sum_i S[i][col] * k[i] + float kv_shard = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + kv_shard += s_shard[r] * ks[i][r]; + } + float kv_col = warp_reduce_sum(kv_shard); - // delta[col] = (v[col] - kv[col]) * beta - float delta_col = (v_t[col] - kv_col) * beta_val; + // delta[col] = (v[col] - g * kv[col]) * beta + float delta_col = (vs[i] - g_val * kv_col) * beta_ts[i]; - // fused: S[i][col] = g[i] * S[i][col] + k[i] * delta[col] - // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] - float attn_partial = 0.0f; + // fused: S[i][col] = g * S[i][col] + k[i] * delta[col] + // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] + float attn_partial = 0.0f; #pragma unroll - for (int r = 0; r < rows_per_lane; r++) { - const int i = r * warp_size + lane; - s_shard[r] = expf(g_t[i]) * s_shard[r] + k_t[i] * delta_col; - attn_partial += s_shard[r] * q_t[i]; + for (int r = 0; r < rows_per_lane; r++) { + s_shard[r] = g_val * s_shard[r] + ks[i][r] * delta_col; + attn_partial += s_shard[r] * qs[i][r]; + } + + float attn_col = warp_reduce_sum(attn_partial); + + if (lane == 0) { + attn_data[col] = attn_col * scale; + } + } else { + // kv[col] = sum_i g[i] * S[i][col] * k[i] + float kv_shard = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + kv_shard += g_vals[i][r] * s_shard[r] * ks[i][r]; + } + + float kv_col = warp_reduce_sum(kv_shard); + + // delta[col] = (v[col] - kv[col]) * beta + float delta_col = (vs[i] - kv_col) * beta_ts[i]; + + // fused: S[i][col] = g[i] * S[i][col] + k[i] * delta[col] + // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] + float attn_partial = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + s_shard[r] = g_vals[i][r] * s_shard[r] + ks[i][r] * delta_col; + attn_partial += s_shard[r] * qs[i][r]; + } + + float attn_col = warp_reduce_sum(attn_partial); + + if (lane == 0) { + attn_data[col] = attn_col * scale; + } } - float attn_col = warp_reduce_sum(attn_partial); + attn_data += S_v * H; + } + } - if (lane == 0) { - attn_data[col] = attn_col * scale; + if (tokens_to_process != n_tokens) { + // handle tail case +#pragma unroll + for (int i = 0; i < n_tokens_per_loop; i++) { + const int t = tokens_to_process + i; + if (t >= n_tokens) { + break; + } + const float * q_t = q + iq3 * sq3 + t * sq2 + iq1 * sq1; + const float * k_t = k + iq3 * sq3 + t * sq2 + iq1 * sq1; + const float * v_t = v + sequence * sv3 + t * sv2 + h_idx * sv1; + const int64_t gb_offset = sequence * sb3 + t * sb2 + h_idx * sb1; + const float * beta_t = beta + gb_offset; + const float * g_t = g + gb_offset * (KDA ? S_v : 1); + + beta_ts[i] = *beta_t; + vs[i] = v_t[col]; + if constexpr (!KDA) { + g_vals[i][0] = expf(*g_t); + } +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + qs[i][r] = q_t[idx[r]]; + ks[i][r] = k_t[idx[r]]; + if constexpr (KDA) { + g_vals[i][r] = expf(g_t[idx[r]]); + } } } +#pragma unroll + for (int i = 0; i < n_tokens_per_loop; i++) { + const int t = tokens_to_process + i; + if (t >= n_tokens) { + break; + } + if constexpr (!KDA) { + const float g_val = g_vals[i][0]; - attn_data += S_v * H; - } + // kv[col] = (S^T @ k)[col] = sum_i S[i][col] * k[i] + float kv_shard = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + kv_shard += s_shard[r] * ks[i][r]; + } + float kv_col = warp_reduce_sum(kv_shard); - // Write state back to global memory + // delta[col] = (v[col] - g * kv[col]) * beta + float delta_col = (vs[i] - g_val * kv_col) * beta_ts[i]; + + // fused: S[i][col] = g * S[i][col] + k[i] * delta[col] + // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] + float attn_partial = 0.0f; #pragma unroll - for (int r = 0; r < rows_per_lane; r++) { - const int i = r * warp_size + lane; - state[i * S_v + col] = s_shard[r]; + for (int r = 0; r < rows_per_lane; r++) { + s_shard[r] = g_val * s_shard[r] + ks[i][r] * delta_col; + attn_partial += s_shard[r] * qs[i][r]; + } + + float attn_col = warp_reduce_sum(attn_partial); + + if (lane == 0) { + attn_data[col] = attn_col * scale; + } + } else { + // kv[col] = sum_i g[i] * S[i][col] * k[i] + float kv_shard = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + kv_shard += g_vals[i][r] * s_shard[r] * ks[i][r]; + } + + float kv_col = warp_reduce_sum(kv_shard); + + // delta[col] = (v[col] - kv[col]) * beta + float delta_col = (vs[i] - kv_col) * beta_ts[i]; + + // fused: S[i][col] = g[i] * S[i][col] + k[i] * delta[col] + // attn[col] = (S^T @ q)[col] = sum_i S[i][col] * q[i] + float attn_partial = 0.0f; +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + s_shard[r] = g_vals[i][r] * s_shard[r] + ks[i][r] * delta_col; + attn_partial += s_shard[r] * qs[i][r]; + } + + float attn_col = warp_reduce_sum(attn_partial); + + if (lane == 0) { + attn_data[col] = attn_col * scale; + } + } + + attn_data += S_v * H; + } } -} -static size_t calculate_smem(const int sv, int cc) -{ - size_t smem = 0; - if ((GGML_CUDA_CC_IS_AMD(cc) && !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_RDNA4(cc)) || GGML_CUDA_CC_IS_MTHREADS(cc)) { - smem = sv * sv * sizeof(float); + // Write state back to global memory (transposed layout) +#pragma unroll + for (int r = 0; r < rows_per_lane; r++) { + state[col * S_v + idx[r]] = s_shard[r]; } - return smem; } template @@ -165,37 +440,74 @@ static void launch_gated_delta_net( int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + // for KDA we have to store one g per row, so we have higher register pressure + const int no_tail = (n_tokens % 32) == 0 && (n_tokens != 1); + constexpr int tail_thresh = KDA ? 14 : 20; switch (S_v) { case 16: - gated_delta_net_cuda<16, KDA><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, - n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); + if (no_tail) { + gated_delta_net_cuda_no_tail<16, KDA, 32><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, + sb3, neqk1_magic, rq3_magic, scale); + } else if (n_tokens >= tail_thresh) { + gated_delta_net_cuda_tail<16, KDA, tail_thresh><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, + sb3, neqk1_magic, rq3_magic, scale); + } else { + gated_delta_net_cuda_no_tail<16, KDA, 1><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, + sb3, neqk1_magic, rq3_magic, scale); + } break; case 32: - gated_delta_net_cuda<32, KDA><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, - n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); - break; - case 64: { - constexpr int sv = 64; - size_t smem = calculate_smem(sv, cc); - gated_delta_net_cuda<<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, - n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); - break; - } - case 128: { - constexpr int sv = 128; - size_t smem = calculate_smem(sv, cc); - gated_delta_net_cuda<<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, - n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); + if (no_tail) { + gated_delta_net_cuda_no_tail<32, KDA, 32><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, + sb3, neqk1_magic, rq3_magic, scale); + } else if (n_tokens >= tail_thresh) { + gated_delta_net_cuda_tail<32, KDA, tail_thresh><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, + sb3, neqk1_magic, rq3_magic, scale); + } else { + gated_delta_net_cuda_no_tail<32, KDA, 1><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, + sb3, neqk1_magic, rq3_magic, scale); + } break; - } + case 64: + { + if (no_tail) { + gated_delta_net_cuda_no_tail<64, KDA, 32><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, + sb2, sb3, neqk1_magic, rq3_magic, scale); + } else if (n_tokens >= tail_thresh) { + gated_delta_net_cuda_tail<64, KDA, tail_thresh><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, + sb2, sb3, neqk1_magic, rq3_magic, scale); + } else { + gated_delta_net_cuda_no_tail<64, KDA, 1><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, + sb2, sb3, neqk1_magic, rq3_magic, scale); + } + break; + } + case 128: + { + if (no_tail) { + gated_delta_net_cuda_no_tail<128, KDA, 32><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, + sb2, sb3, neqk1_magic, rq3_magic, scale); + } else if (n_tokens >= tail_thresh) { + gated_delta_net_cuda_tail<128, KDA, tail_thresh><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, + sb2, sb3, neqk1_magic, rq3_magic, scale); + } else { + gated_delta_net_cuda_no_tail<128, KDA, 1><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, + sb2, sb3, neqk1_magic, rq3_magic, scale); + } + break; + } default: GGML_ABORT("fatal error"); break; diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 0b77d5349b86..e06f266b5337 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -2466,13 +2466,14 @@ kernel void kernel_gated_delta_net_impl( const float scale = 1.0f / sqrt((float)S_v); - device const float * s_ptr = (device const float *) (s) + (i23*args.ne21 + i21)*S_v*S_v + i20; + // state is stored transposed: M[i20][is] = S[is][i20], so row i20 is contiguous + device const float * s_ptr = (device const float *) (s) + (i23*args.ne21 + i21)*S_v*S_v + i20*S_v; float ls[NSG]; FOR_UNROLL (short j = 0; j < NSG; j++) { const short is = tx*NSG + j; - ls[j] = s_ptr[is*S_v]; + ls[j] = s_ptr[is]; } device float * dst_attn = (device float *) (dst) + (i23*args.ne22*args.ne21 + i21)*S_v + i20; @@ -2533,11 +2534,11 @@ kernel void kernel_gated_delta_net_impl( g_ptr += args.ne21*G; } - device float * dst_state = (device float *) (dst) + args.ne23*args.ne22*args.ne21*S_v + (i23*args.ne21 + i21)*S_v*S_v + i20; + device float * dst_state = (device float *) (dst) + args.ne23*args.ne22*args.ne21*S_v + (i23*args.ne21 + i21)*S_v*S_v + i20*S_v; FOR_UNROLL (short j = 0; j < NSG; j++) { const short is = tx*NSG + j; - dst_state[is*S_v] = ls[j]; + dst_state[is] = ls[j]; } #undef S_v diff --git a/src/models/delta-net-base.cpp b/src/models/delta-net-base.cpp index a62dbc15dd05..6bc989c95099 100644 --- a/src/models/delta-net-base.cpp +++ b/src/models/delta-net-base.cpp @@ -225,9 +225,8 @@ std::pair llm_build_delta_net_base::build_delta_ne ggml_tensor * kg_t = ggml_cont(ctx0, ggml_transpose(ctx0, kg)); cb(kg_t, "key_gdiff_t", il); - ggml_tensor * s_t = ggml_transpose(ctx0, s); - s_t = ggml_cont_4d(ctx0, s_t, S_v, S_v, 1, H_v * n_seqs); - cb(s_t, "dnet_add_ch_state", il); + s = ggml_reshape_4d(ctx0, s, S_v, S_v, 1, H_v * n_seqs); + cb(s, "dnet_add_ch_state", il); // [CS, S_v, n_chunks, H_v * n_seqs] ggml_tensor * v_t = ggml_cont(ctx0, ggml_transpose(ctx0, v)); @@ -240,7 +239,7 @@ std::pair llm_build_delta_net_base::build_delta_ne ggml_tensor * ch_kg_t = get_slice_2d(ctx0, kg_t, chunk); // [ CS, S_k, 1, H_v * n_seqs] // [CS, S_v, 1, H_v * n_seqs] - ggml_tensor * v_t_p = ggml_mul_mat(ctx0, ch_k_cd, s_t); + ggml_tensor * v_t_p = ggml_mul_mat(ctx0, ch_k_cd, s); cb(v_t_p, "v_prime", il); // [CS, S_v, 1, H_v * n_seqs] @@ -252,7 +251,7 @@ std::pair llm_build_delta_net_base::build_delta_ne cb(v_attn, "v_attn", il); // [S_v, CS, 1, H_v * n_seqs] - ggml_tensor * attn_inter = ggml_mul_mat(ctx0, s_t, ch_q_g_exp); + ggml_tensor * attn_inter = ggml_mul_mat(ctx0, s, ch_q_g_exp); cb(attn_inter, "attn_inter", il); // [S_v, CS, 1, H_v * n_seqs] @@ -268,13 +267,11 @@ std::pair llm_build_delta_net_base::build_delta_ne // last_recurrent_state = last_recurrent_state * g_last + kgdmulvnew ggml_tensor * ch_g_last_exp_t = get_slice_2d(ctx0, g_last_exp_t, chunk); - s_t = ggml_mul(ctx0, s_t, ch_g_last_exp_t); - s_t = ggml_add(ctx0, s_t, kgv); - cb(s_t, "dnet_add_ch_state", il); + s = ggml_mul(ctx0, s, ch_g_last_exp_t); + s = ggml_add(ctx0, s, kgv); + cb(s, "dnet_add_ch_state", il); } - s_t = ggml_reshape_4d(ctx0, s_t, S_v, S_v, H_v, n_seqs); - // truncate padded tokens ggml_tensor * o = ggml_view_4d(ctx0, v, S_v, n_tokens, H_v, n_seqs, @@ -282,7 +279,7 @@ std::pair llm_build_delta_net_base::build_delta_ne ggml_row_size(v->type, S_v * CS * n_chunks), ggml_row_size(v->type, S_v * CS * n_chunks * H_v), 0); o = ggml_permute (ctx0, o, 0, 2, 1, 3); // [S_v, H_v, n_tokens, n_seqs] - s = ggml_transpose(ctx0, s_t); + s = ggml_reshape_4d(ctx0, s, S_v, S_v, H_v, n_seqs); cb(s, "output_state", il); return {o, s}; @@ -341,11 +338,9 @@ std::pair llm_build_delta_net_base::build_delta_ne g = ggml_exp(ctx0, g); s = ggml_mul(ctx0, s, g); - ggml_tensor * s_t = ggml_cont(ctx0, ggml_transpose(ctx0, s)); - // [1, S_v, H_v, n_seqs] ggml_tensor * sk; - sk = ggml_mul (ctx0, s_t, k); + sk = ggml_mul (ctx0, s, k); sk = ggml_sum_rows(ctx0, sk); // [S_v, 1, H_v, n_seqs] @@ -362,15 +357,14 @@ std::pair llm_build_delta_net_base::build_delta_ne k = ggml_repeat(ctx0, k, s); kd = ggml_mul (ctx0, k, d_t); - s_t = ggml_add(ctx0, s_t, kd); + s = ggml_add(ctx0, s, kd); - cb(s_t, "dnet_add_ar_state", il); + cb(s, "dnet_add_ar_state", il); - ggml_tensor * s_q = ggml_mul (ctx0, s_t, q); + ggml_tensor * s_q = ggml_mul (ctx0, s, q); ggml_tensor * o = ggml_sum_rows(ctx0, s_q); o = ggml_permute (ctx0, o, 2, 0, 1, 3); // [S_v, H_v, n_tokens, n_seqs] - s = ggml_transpose(ctx0, s_t); // [S_v, S_v, H_v, n_seqs] return {o, s}; } diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index a821655d10c7..e21eaeeee444 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -8456,6 +8456,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 8, 32, 4, 2, 2)); test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 4, 2, 1, true)); test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 4, 1, 1, true)); + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 33, 1, 1, true)); // KDA (vector gate) test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 1, 1, 1, false, true)); test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 1, 2, 1, false, true)); @@ -8464,6 +8465,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 4, 2, 1, false, true)); test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 8, 32, 4, 2, 2, false, true)); test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 4, 2, 1, true, true)); + test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64, 33, 1, 1, true, true)); test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 16, 4, 2, 1, true, true)); #if 0