From 83d5c9d12e9508dfc36840bb210e4daedc6b9889 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Thu, 8 Oct 2026 13:49:24 +0200 Subject: [PATCH 1/7] Remove lmem by indexing base_output with compile-time constants --- ggml/src/ggml-cuda/gated-delta-net-mma.cu | 27 +++++++++++++++++------ 1 file changed, 20 insertions(+), 7 deletions(-) diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu index 955647dc058b..101ce228d6b6 100644 --- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu +++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu @@ -229,14 +229,23 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m } // V32 overlaps Gram work with one output tile per projection warp. constexpr bool SPLIT_PROJECTION = C == 16 && V == 32; - acc base_output[((C / 16) * (V / N) + WARPS - 1) / WARPS]; + constexpr int OUTPUT_TILES = (C / 16) * (V / N); + constexpr int OUTPUT_STEP = SPLIT_PROJECTION ? 4 : WARPS; + constexpr int OUTPUT_SLOTS = (OUTPUT_TILES + OUTPUT_STEP - 1) / OUTPUT_STEP; + const int output_start = SPLIT_PROJECTION ? warp - 4 : warp; + acc base_output[OUTPUT_SLOTS]; auto project = [&]() { - if constexpr (V == 32) { + if constexpr (SPLIT_PROJECTION) { if (warp < 4) { return; } } - for (int ti = (V == 32 ? warp - 4 : warp); ti < (C / 16) * (V / N); ti += (V == 32 ? 4 : WARPS)) { +#pragma unroll + for (int slot = 0; slot < OUTPUT_SLOTS; ++slot) { + const int ti = output_start + slot * OUTPUT_STEP; + if (ti >= OUTPUT_TILES) { + continue; + } const int r = (ti / (V / N)) * 16, c = (ti % (V / N)) * N; acc qs, ks; product(qs, s.q, s.scratch.state, r, c); @@ -247,7 +256,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m const float decay = expf(s.prefix[t]); const float vv = t < valid ? a.v[voff + (t0 + t) * a.sv2 + v] : 0.f; s.delta.store(t, v, s.beta[t] * (vv - decay * ks.x[i])); - base_output[ti / WARPS].x[i] = decay * qs.x[i]; + base_output[slot].x[i] = decay * qs.x[i]; } } }; @@ -312,8 +321,12 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m store(s.scratch.solved, x, r, c); } __syncthreads(); - for (int ti = (SPLIT_PROJECTION ? warp - 4 : warp); ti < (C / 16) * (V / N) && (!SPLIT_PROJECTION || warp >= 4); - ti += WARPS) { +#pragma unroll + for (int slot = 0; slot < OUTPUT_SLOTS; ++slot) { + const int ti = output_start + slot * OUTPUT_STEP; + if (ti >= OUTPUT_TILES || (SPLIT_PROJECTION && warp < 4)) { + continue; + } const int r = (ti / (V / N)) * 16, c = (ti % (V / N)) * N; acc x; product(x, s.p, s.scratch.solved, r, c); @@ -321,7 +334,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m for (int i = 0; i < acc::ne; ++i) { const int t = r + acc::get_i(i), v = c + acc::get_j(i); if (t < valid) { - out[(t0 + t) * a.H * D + v] = base_output[ti / WARPS].x[i] + x.x[i]; + out[(t0 + t) * a.H * D + v] = base_output[slot].x[i] + x.x[i]; } } } From 5b1b63d3ff2f65049d61171747153886f5a2446c Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Thu, 8 Oct 2026 14:54:21 +0200 Subject: [PATCH 2/7] Increase arithmetic density by mocked mma.sync.m16n16k16 --- ggml/src/ggml-cuda/gated-delta-net-mma.cu | 8 ++------ ggml/src/ggml-cuda/mma.cuh | 15 +++++++++++++-- 2 files changed, 15 insertions(+), 8 deletions(-) diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu index 101ce228d6b6..14a1ce248da8 100644 --- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu +++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu @@ -63,9 +63,9 @@ using acc = tile<16, 16, float, DATA_LAYOUT_J_MAJOR>; using a_tile = tile<16, 8, nv_bfloat162, DATA_LAYOUT_I_MAJOR_MIRRORED>; using b_tile = a_tile; #else -using acc = tile<16, 8, float>; +using acc = tile<16, 16, float>; using a_tile = tile<16, 8, nv_bfloat162>; -using b_tile = tile<8, 8, nv_bfloat162>; +using b_tile = a_tile; #endif // GGML_USE_HIP template @@ -127,11 +127,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m #if defined(AMPERE_MMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)) static_assert(ggml_cuda_get_physical_warp_size() == 32, "GDN MMA requires 32-lane warps"); constexpr int WARPS = 8; -#ifdef GGML_USE_HIP constexpr int N = 16; -#else - constexpr int N = 8; -#endif // GGML_USE_HIP static_assert(C >= 16 && (C & (C - 1)) == 0 && V % N == 0 && D % V == 0, "invalid GDN tile"); extern __shared__ __align__(128) unsigned char bytes[]; auto & s = *reinterpret_cast *>(bytes); diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh index 3e4946b15c6a..56ae6b15f7bf 100644 --- a/ggml/src/ggml-cuda/mma.cuh +++ b/ggml/src/ggml-cuda/mma.cuh @@ -1337,7 +1337,18 @@ namespace ggml_cuda_mma { template static __device__ __forceinline__ void mma( tile<16, 16, float, dl_d> & D, const tile<16, 8, nv_bfloat162, dl_ab> & A, const tile<16, 8, nv_bfloat162, dl_ab> & B) { -#if defined(AMD_WMMA_AVAILABLE) +#if defined(AMPERE_MMA_AVAILABLE) + static_assert(dl_ab == DATA_LAYOUT_I_MAJOR && dl_d == DATA_LAYOUT_I_MAJOR, "bad data layout"); + const int * Axi = (const int *) A.x; + const int * Bxi = (const int *) B.x; + int * Dxi = (int *) D.x; + asm("mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, {%0, %1, %2, %3};" + : "+r"(Dxi[0]), "+r"(Dxi[1]), "+r"(Dxi[2]), "+r"(Dxi[3]) + : "r"(Axi[0]), "r"(Axi[1]), "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[0]), "r"(Bxi[2])); + asm("mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, {%0, %1, %2, %3};" + : "+r"(Dxi[4]), "+r"(Dxi[5]), "+r"(Dxi[6]), "+r"(Dxi[7]) + : "r"(Axi[0]), "r"(Axi[1]), "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[1]), "r"(Bxi[3])); +#elif defined(AMD_WMMA_AVAILABLE) #if defined(RDNA4) using bf16x8_t = __attribute__((ext_vector_type(8))) __bf16; using floatx8_t = __attribute__((ext_vector_type(8))) float; @@ -1379,7 +1390,7 @@ namespace ggml_cuda_mma { #else GGML_UNUSED_VARS(D, A, B); NO_DEVICE_CODE; -#endif // defined(AMD_WMMA_AVAILABLE) +#endif // defined(AMD_WMMA_AVAILABLE) || defined(AMPERE_MMA_AVAILABLE) } template From 4d2e4dc02a40443176639d515f48193600d8d5c8 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 9 Oct 2026 15:35:50 +0200 Subject: [PATCH 3/7] Hoist loop invariants --- ggml/src/ggml-cuda/gated-delta-net-mma.cu | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu index 14a1ce248da8..35ed5f5c39b6 100644 --- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu +++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu @@ -140,6 +140,11 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m const int64_t qoff = (seq / a.rq3) * a.sq3 + (h % a.H_k) * a.sq1; const int64_t voff = seq * a.sv3 + h * a.sv1 + v0; const int64_t goff = seq * a.sb3 + h * a.sb1; + const float * const g_base = a.g + goff; + const float * const beta_base = a.beta + goff; + const float * const v_base = a.v + voff; + const float * const state_base = a.state + soff; + float * const state_out_base = a.state_out + soff; float * out = a.dst + ((int64_t) seq * a.n_tokens * a.H + h) * D + v0; constexpr int STATE_TILES = (D / 16) * (V / N); constexpr int PER_WARP = STATE_TILES / WARPS; @@ -150,7 +155,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m const int tile_index = warp + j * WARPS, r = (tile_index / (V / N)) * 16, c = (tile_index % (V / N)) * N; #pragma unroll for (int i = 0; i < acc::ne; ++i) { - state[j].x[i] = a.state[soff + (v0 + c + acc::get_j(i)) * D + r + acc::get_i(i)]; + state[j].x[i] = state_base[(v0 + c + acc::get_j(i)) * D + r + acc::get_i(i)]; } } auto read_qk = [&](int64_t t, float4 & q, float4 & k) { @@ -195,8 +200,8 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m float carry = 0.f; for (int base = 0; base < C; base += 32) { const int t = base + lane; - float g = t < valid ? a.g[goff + (t0 + t) * a.sb2] : 0.f; - const float beta = t < valid ? a.beta[goff + (t0 + t) * a.sb2] : 0.f; + float g = t < valid ? g_base[(t0 + t) * a.sb2] : 0.f; + const float beta = t < valid ? beta_base[(t0 + t) * a.sb2] : 0.f; #pragma unroll for (int offset = 1; offset < 32; offset *= 2) { const float previous = __shfl_up_sync(0xffffffff, g, offset, 32); @@ -250,7 +255,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m for (int i = 0; i < acc::ne; ++i) { const int t = r + acc::get_i(i), v = c + acc::get_j(i); const float decay = expf(s.prefix[t]); - const float vv = t < valid ? a.v[voff + (t0 + t) * a.sv2 + v] : 0.f; + const float vv = t < valid ? v_base[(t0 + t) * a.sv2 + v] : 0.f; s.delta.store(t, v, s.beta[t] * (vv - decay * ks.x[i])); base_output[slot].x[i] = decay * qs.x[i]; } @@ -357,7 +362,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m const int ti = warp + j * WARPS, r = (ti / (V / N)) * 16, c = (ti % (V / N)) * N; #pragma unroll for (int i = 0; i < acc::ne; ++i) { - a.state_out[soff + (v0 + c + acc::get_j(i)) * D + r + acc::get_i(i)] = state[j].x[i]; + state_out_base[(v0 + c + acc::get_j(i)) * D + r + acc::get_i(i)] = state[j].x[i]; } } #else From 1c8e55ae533f0017ad32624b257e4d72e7192918 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 9 Oct 2026 15:48:50 +0200 Subject: [PATCH 4/7] Add store4 to avoid smem bank conflicts for qk storing --- ggml/src/ggml-cuda/gated-delta-net-mma.cu | 24 +++++++++++++++++++---- 1 file changed, 20 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu index 35ed5f5c39b6..520255e128b9 100644 --- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu +++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu @@ -33,6 +33,24 @@ template struct matrix { #endif // GGML_USE_MUSA } + __device__ __forceinline__ void store4(int r, int c, float4 x) { +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + store2(r, c, make_float2(x.x, x.y)); + store2(r, c + 2, make_float2(x.z, x.w)); +#else + static_assert(C % 4 == 0, "store4 needs aligned rows"); + const int i = index(r, c); + const auto h0 = __float22bfloat162_rn(make_float2(x.x, x.y)); + const auto h1 = __float22bfloat162_rn(make_float2(x.z, x.w)); + const float2 r0 = __bfloat1622float2(h0); + const float2 r1 = __bfloat1622float2(h1); + const auto l0 = __float22bfloat162_rn(make_float2(x.x - r0.x, x.y - r0.y)); + const auto l1 = __float22bfloat162_rn(make_float2(x.z - r1.x, x.w - r1.y)); + *reinterpret_cast(hi + i) = make_uint2(reinterpret_cast(h0), reinterpret_cast(h1)); + *reinterpret_cast(lo + i) = make_uint2(reinterpret_cast(l0), reinterpret_cast(l1)); +#endif // GGML_USE_HIP || GGML_USE_MUSA + } + __device__ __forceinline__ void store(int r, int c, float x) { const int i = index(r, c); const bf16 h = __float2bfloat16(x); @@ -190,10 +208,8 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m } else { read_qk(t0 + 8 * j, q, k); } - s.q.store2(t, d, make_float2(q.x * a.scale, q.y * a.scale)); - s.q.store2(t, d + 2, make_float2(q.z * a.scale, q.w * a.scale)); - s.k.store2(t, d, make_float2(k.x, k.y)); - s.k.store2(t, d + 2, make_float2(k.z, k.w)); + s.q.store4(t, d, make_float4(q.x * a.scale, q.y * a.scale, q.z * a.scale, q.w * a.scale)); + s.k.store4(t, d, k); } // Load gates in parallel before the FP32 prefix scan. if (warp == 0) { From 26de4efee68fcb8f862e984d3ff0edb8a11e6a0d Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 9 Oct 2026 16:43:26 +0200 Subject: [PATCH 5/7] Load float8 to get compiler to unroll q/k fetches --- ggml/src/ggml-cuda/gated-delta-net-mma.cu | 103 +++++++++++++--------- 1 file changed, 61 insertions(+), 42 deletions(-) diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu index 520255e128b9..f816b5ae07e8 100644 --- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu +++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu @@ -8,6 +8,12 @@ namespace { constexpr int D = 128; using bf16 = nv_bfloat16; +// TODO: Reuse from quantize.cu/extract into common.cuh or shared definition +struct alignas(32) float8 { + float x; float y; float z; float w; + float p; float q; float r; float s; +}; + template struct matrix { bf16 hi[R * C], lo[R * C]; @@ -33,21 +39,29 @@ template struct matrix { #endif // GGML_USE_MUSA } - __device__ __forceinline__ void store4(int r, int c, float4 x) { + __device__ __forceinline__ void store8(int r, int c, float8 x) { #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) store2(r, c, make_float2(x.x, x.y)); store2(r, c + 2, make_float2(x.z, x.w)); + store2(r, c + 4, make_float2(x.p, x.q)); + store2(r, c + 6, make_float2(x.r, x.s)); #else - static_assert(C % 4 == 0, "store4 needs aligned rows"); + static_assert(C % 8 == 0, "store8 needs aligned rows"); const int i = index(r, c); const auto h0 = __float22bfloat162_rn(make_float2(x.x, x.y)); const auto h1 = __float22bfloat162_rn(make_float2(x.z, x.w)); + const auto h2 = __float22bfloat162_rn(make_float2(x.p, x.q)); + const auto h3 = __float22bfloat162_rn(make_float2(x.r, x.s)); const float2 r0 = __bfloat1622float2(h0); const float2 r1 = __bfloat1622float2(h1); + const float2 r2 = __bfloat1622float2(h2); + const float2 r3 = __bfloat1622float2(h3); const auto l0 = __float22bfloat162_rn(make_float2(x.x - r0.x, x.y - r0.y)); const auto l1 = __float22bfloat162_rn(make_float2(x.z - r1.x, x.w - r1.y)); - *reinterpret_cast(hi + i) = make_uint2(reinterpret_cast(h0), reinterpret_cast(h1)); - *reinterpret_cast(lo + i) = make_uint2(reinterpret_cast(l0), reinterpret_cast(l1)); + const auto l2 = __float22bfloat162_rn(make_float2(x.p - r2.x, x.q - r2.y)); + const auto l3 = __float22bfloat162_rn(make_float2(x.r - r3.x, x.s - r3.y)); + *reinterpret_cast(hi + i) = make_uint4(reinterpret_cast(h0), reinterpret_cast(h1), reinterpret_cast(h2), reinterpret_cast(h3)); + *reinterpret_cast(lo + i) = make_uint4(reinterpret_cast(l0), reinterpret_cast(l1), reinterpret_cast(l2), reinterpret_cast(l3)); #endif // GGML_USE_HIP || GGML_USE_MUSA } @@ -140,13 +154,13 @@ template __device__ __forceinline__ void store(matrix & m, // Each block owns V independent value columns and advances by C tokens. // Keep recurrent state in FP32; pack high/residual BF16 operands inside the block. -template +template static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_mma_args a) { #if defined(AMPERE_MMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)) static_assert(ggml_cuda_get_physical_warp_size() == 32, "GDN MMA requires 32-lane warps"); constexpr int WARPS = 8; constexpr int N = 16; - static_assert(C >= 16 && (C & (C - 1)) == 0 && V % N == 0 && D % V == 0, "invalid GDN tile"); + static_assert(C == 16 && V % N == 0 && D % V == 0, "invalid GDN tile"); extern __shared__ __align__(128) unsigned char bytes[]; auto & s = *reinterpret_cast *>(bytes); const int lane = threadIdx.x, warp = threadIdx.y, tid = warp * 32 + lane; @@ -176,41 +190,39 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m state[j].x[i] = state_base[(v0 + c + acc::get_j(i)) * D + r + acc::get_i(i)]; } } - auto read_qk = [&](int64_t t, float4 & q, float4 & k) { - t += warp; - q = k = make_float4(0, 0, 0, 0); + // Each half warp loads one row, with eight contiguous floats per lane. + const int input_row = warp + 8 * (lane / 16), input_col = 8 * (lane % 16); + auto read_qk = [&](int64_t t, float8 & q, float8 & k) { + t += input_row; + q = k = {}; if (t < a.n_tokens) { - const int64_t off = qoff + t * a.sq2 + 4 * lane; - if ((((uintptr_t) (a.q + off) | (uintptr_t) (a.k + off)) & 15) == 0) { - q = *reinterpret_cast(a.q + off); - k = *reinterpret_cast(a.k + off); + const int64_t off = qoff + t * a.sq2 + input_col; + if constexpr (ALIGNED) { + q = *reinterpret_cast(a.q + off); + k = *reinterpret_cast(a.k + off); } else { - q = make_float4(a.q[off], a.q[off + 1], a.q[off + 2], a.q[off + 3]); - k = make_float4(a.k[off], a.k[off + 1], a.k[off + 2], a.k[off + 3]); + q = {a.q[off], a.q[off + 1], a.q[off + 2], a.q[off + 3], a.q[off + 4], a.q[off + 5], a.q[off + 6], a.q[off + 7]}; + k = {a.k[off], a.k[off + 1], a.k[off + 2], a.k[off + 3], a.k[off + 4], a.k[off + 5], a.k[off + 6], a.k[off + 7]}; } } }; - // Carry the first eight input rows across chunks to hide global-load latency. + // Carry input rows across chunks to hide global-load latency. constexpr bool PREFETCH = C == 16 && V == 32; - float4 next_q, next_k; + float8 next_q, next_k; if constexpr (PREFETCH) { read_qk(0, next_q, next_k); } for (int64_t t0 = 0; t0 < a.n_tokens; t0 += C) { const int valid = min((int64_t) C, a.n_tokens - t0); -#pragma unroll - for (int j = 0; j < C / 8; ++j) { - const int t = warp + 8 * j, d = 4 * lane; - float4 q, k; - if (PREFETCH && j == 0) { - q = next_q; - k = next_k; - } else { - read_qk(t0 + 8 * j, q, k); - } - s.q.store4(t, d, make_float4(q.x * a.scale, q.y * a.scale, q.z * a.scale, q.w * a.scale)); - s.k.store4(t, d, k); + float8 q, k; + if constexpr (PREFETCH) { + q = next_q; + k = next_k; + } else { + read_qk(t0, q, k); } + s.q.store8(input_row, input_col, {q.x * a.scale, q.y * a.scale, q.z * a.scale, q.w * a.scale, q.p * a.scale, q.q * a.scale, q.r * a.scale, q.s * a.scale}); + s.k.store8(input_row, input_col, k); // Load gates in parallel before the FP32 prefix scan. if (warp == 0) { float carry = 0.f; @@ -387,7 +399,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m #endif // defined(AMPERE_MMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)) } -template bool init(int device) { +template bool init(int device) { static std::once_flag once[GGML_CUDA_MAX_DEVICES]; static bool available[GGML_CUDA_MAX_DEVICES] = {}; std::call_once(once[device], [device] { @@ -395,7 +407,7 @@ template bool init(int device) { if (sizeof(shared) > ggml_cuda_info().devices[device].smpbo) { return; } - const auto status = cudaFuncSetAttribute((const void *) gdn_single, + const auto status = cudaFuncSetAttribute((const void *) gdn_single, cudaFuncAttributeMaxDynamicSharedMemorySize, sizeof(shared)); available[device] = status == cudaSuccess; if (status != cudaSuccess) { @@ -407,6 +419,7 @@ template bool init(int device) { enum class gdn_path { ar, slice32_one_block, slice32_two_blocks, slice64, whole_head }; +template static gdn_path select_gdn_mma_path(int device, const ggml_cuda_gdn_mma_args & a) { GGML_ASSERT(device >= 0 && device < GGML_CUDA_MAX_DEVICES); if (!a.eligible || a.H_k != 16 || (a.H != 16 && a.H != 32 && a.H != 48 && a.H != 64) || a.n_tokens < 64 || @@ -418,7 +431,7 @@ static gdn_path select_gdn_mma_path(int device, const ggml_cuda_gdn_mma_args & a if (!GGML_CUDA_CC_IS_RDNA3_5(info.cc) || a.n_tokens < 2048) { return gdn_path::ar; } - return init<16, 64>(device) ? gdn_path::slice64 : gdn_path::ar; + return init<16, 64, 1, ALIGNED>(device) ? gdn_path::slice64 : gdn_path::ar; #elif defined(GGML_USE_MUSA) GGML_UNUSED_VARS(info); return gdn_path::ar; @@ -433,34 +446,40 @@ static gdn_path select_gdn_mma_path(int device, const ggml_cuda_gdn_mma_args & a if (a.H * a.n_seqs * 2 <= info.nsm) { // A grid that already fits in one wave benefits from a larger register budget. if (a.H * a.n_seqs * 4 <= info.nsm) { - return init<16, 32, 1>(device) ? gdn_path::slice32_one_block : gdn_path::ar; + return init<16, 32, 1, ALIGNED>(device) ? gdn_path::slice32_one_block : gdn_path::ar; } - return init<16, 32, 2>(device) ? gdn_path::slice32_two_blocks : gdn_path::ar; + return init<16, 32, 2, ALIGNED>(device) ? gdn_path::slice32_two_blocks : gdn_path::ar; } - return init<16, 128>(device) ? gdn_path::whole_head : gdn_path::ar; + return init<16, 128, 1, ALIGNED>(device) ? gdn_path::whole_head : gdn_path::ar; #endif // GGML_USE_HIP } -} // namespace -bool ggml_cuda_gdn_mma_launch(int device, const ggml_cuda_gdn_mma_args & a, cudaStream_t stream) { +template +static bool launch_gdn_mma(int device, const ggml_cuda_gdn_mma_args & a, cudaStream_t stream) { if (!a.state_out) { return false; } - const auto path = select_gdn_mma_path(device, a); + const auto path = select_gdn_mma_path(device, a); if (path == gdn_path::ar) { return false; } #ifdef GGML_USE_HIP - gdn_single<16, 64><<), stream>>>(a); + gdn_single<16, 64, 1, ALIGNED><<), stream>>>(a); #else if (path == gdn_path::slice32_two_blocks) { - gdn_single<16, 32, 2><<), stream>>>(a); + gdn_single<16, 32, 2, ALIGNED><<), stream>>>(a); } else if (path == gdn_path::slice32_one_block) { - gdn_single<16, 32, 1><<), stream>>>(a); + gdn_single<16, 32, 1, ALIGNED><<), stream>>>(a); } else { - gdn_single<16, 128><<), stream>>>(a); + gdn_single<16, 128, 1, ALIGNED><<), stream>>>(a); } #endif // GGML_USE_HIP CUDA_CHECK(cudaGetLastError()); return true; } +} // namespace + +bool ggml_cuda_gdn_mma_launch(int device, const ggml_cuda_gdn_mma_args & a, cudaStream_t stream) { + const bool aligned = (((uintptr_t) a.q | (uintptr_t) a.k) & 31) == 0 && ((a.sq1 | a.sq2 | a.sq3) & 7) == 0; + return aligned ? launch_gdn_mma(device, a, stream) : launch_gdn_mma(device, a, stream); +} From 9991bac895cceaf7d0b319d1aa5329f1645226f6 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 9 Oct 2026 17:23:44 +0200 Subject: [PATCH 6/7] Hide latency of vv by prefetching/doing manual scheduling This is an alternative to warp specialization obviously. TODO: see if local memory can be gotten rid off --- ggml/src/ggml-cuda/gated-delta-net-mma.cu | 49 ++++++++++++++++++++++- 1 file changed, 47 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu index f816b5ae07e8..bad78c2b6de6 100644 --- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu +++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu @@ -138,6 +138,29 @@ __device__ __forceinline__ void product(acc & c, const matrix & a, const } } +#ifndef GGML_USE_HIP +template +__device__ __forceinline__ acc load_values(const float * base, int64_t stride, int64_t t0, int row, int col, int valid) { + acc values; +#pragma unroll + for (int i = 0; i < acc::ne; i += 2) { + const int t = row + acc::get_i(i), v = col + acc::get_j(i); + float2 vv = make_float2(0.f, 0.f); + if (FULL || t < valid) { + const float * p = base + (t0 + t) * stride + v; + if constexpr (ALIGNED) { + vv = *reinterpret_cast(p); + } else { + vv = make_float2(p[0], p[1]); + } + } + values.x[i] = vv.x; + values.x[i + 1] = vv.y; + } + return values; +} +#endif // GGML_USE_HIP + template __device__ __forceinline__ void store(matrix & m, const acc & x, int r, int c) { #pragma unroll #ifdef GGML_USE_HIP @@ -278,12 +301,33 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m const int r = (ti / (V / N)) * 16, c = (ti % (V / N)) * N; acc qs, ks; product(qs, s.q, s.scratch.state, r, c); +#ifndef GGML_USE_HIP + auto read_v = [&]() { + return valid == C ? load_values(v_base, a.sv2, t0, r, c, valid) : + load_values(v_base, a.sv2, t0, r, c, valid); + }; + acc values; + // Prefetch V for whole heads; value-split blocks need a shorter register lifetime. + constexpr bool PREFETCH_V = V == 128 && ALIGNED; + if constexpr (PREFETCH_V) { + values = read_v(); + } +#endif // GGML_USE_HIP product(ks, s.k, s.scratch.state, r, c); +#ifndef GGML_USE_HIP + if constexpr (!PREFETCH_V) { + values = read_v(); + } +#endif // GGML_USE_HIP #pragma unroll for (int i = 0; i < acc::ne; ++i) { const int t = r + acc::get_i(i), v = c + acc::get_j(i); const float decay = expf(s.prefix[t]); - const float vv = t < valid ? v_base[(t0 + t) * a.sv2 + v] : 0.f; +#ifdef GGML_USE_HIP + const float vv = t < valid ? v_base[(t0 + t) * a.sv2 + v] : 0.f; +#else + const float vv = values.x[i]; +#endif // GGML_USE_HIP s.delta.store(t, v, s.beta[t] * (vv - decay * ks.x[i])); base_output[slot].x[i] = decay * qs.x[i]; } @@ -480,6 +524,7 @@ static bool launch_gdn_mma(int device, const ggml_cuda_gdn_mma_args & a, cudaStr } // namespace bool ggml_cuda_gdn_mma_launch(int device, const ggml_cuda_gdn_mma_args & a, cudaStream_t stream) { - const bool aligned = (((uintptr_t) a.q | (uintptr_t) a.k) & 31) == 0 && ((a.sq1 | a.sq2 | a.sq3) & 7) == 0; + const bool aligned = (((uintptr_t) a.q | (uintptr_t) a.k) & 31) == 0 && ((a.sq1 | a.sq2 | a.sq3) & 7) == 0 && + ((uintptr_t) a.v & 7) == 0 && ((a.sv1 | a.sv2 | a.sv3) & 1) == 0; return aligned ? launch_gdn_mma(device, a, stream) : launch_gdn_mma(device, a, stream); } From ce37862f844d90e96794e5f2e88f566af3c18422 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 9 Oct 2026 17:28:29 +0200 Subject: [PATCH 7/7] Recompute state_out to avoid lmem Should benefit meomries that have to write spills to DRAM instead of hosting it in L2 --- ggml/src/ggml-cuda/gated-delta-net-mma.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu index bad78c2b6de6..d3b43182ab4c 100644 --- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu +++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu @@ -199,7 +199,6 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m const float * const beta_base = a.beta + goff; const float * const v_base = a.v + voff; const float * const state_base = a.state + soff; - float * const state_out_base = a.state_out + soff; float * out = a.dst + ((int64_t) seq * a.n_tokens * a.H + h) * D + v0; constexpr int STATE_TILES = (D / 16) * (V / N); constexpr int PER_WARP = STATE_TILES / WARPS; @@ -429,6 +428,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m } __syncthreads(); } + float * const state_out_base = a.state_out + (int64_t) (blockIdx.x / (D / V)) * D * D; #pragma unroll for (int j = 0; j < PER_WARP; ++j) { const int ti = warp + j * WARPS, r = (ti / (V / N)) * 16, c = (ti % (V / N)) * N;