diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp
index a8eab90a6f50..c5ff7265d704 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-device.cpp
@@ -692,6 +692,25 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_rwkv(ggml_metal_
     return res;
 }
 
+ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_gated_delta_net_back(ggml_metal_library_t lib, const ggml_tensor * op) {
+    char name[256];
+    snprintf(name, 256, "kernel_gated_delta_net_back_f32_%d", (int) op->src[2]->ne[0]);
+    ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
+    if (!res.pipeline) {
+        res = ggml_metal_library_compile_pipeline(lib, name, name, nullptr);
+    }
+    return res;
+}
+
+ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_gated_delta_net_back_zero(ggml_metal_library_t lib) {
+    const char * name = "kernel_gated_delta_net_back_zero";
+    ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
+    if (!res.pipeline) {
+        res = ggml_metal_library_compile_pipeline(lib, name, name, nullptr);
+    }
+    return res;
+}
+
 ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_gated_delta_net(ggml_metal_library_t lib, const ggml_tensor * op) {
     char base[256];
     char name[256];
diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h
index 007895e451ed..51f563fa2020 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.h
+++ b/ggml/src/ggml-metal/ggml-metal-device.h
@@ -135,6 +135,8 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_ssm_scan
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_ssm_scan_ssd_mma  (ggml_metal_library_t lib, const struct ggml_tensor * op);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_rwkv              (ggml_metal_library_t lib, const struct ggml_tensor * op);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_gated_delta_net   (ggml_metal_library_t lib, const struct ggml_tensor * op);
+struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_gated_delta_net_back(ggml_metal_library_t lib, const struct ggml_tensor * op);
+struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_gated_delta_net_back_zero(ggml_metal_library_t lib);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_solve_tri         (ggml_metal_library_t lib, const struct ggml_tensor * op);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_ext        (ggml_metal_library_t lib, const struct ggml_tensor * op, int nsg, int nxpsg, int r1ptg);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm            (ggml_metal_library_t lib, const struct ggml_tensor * op);
diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m
index e5556fa6989d..82215837c8fa 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.m
+++ b/ggml/src/ggml-metal/ggml-metal-device.m
@@ -1736,6 +1736,8 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
             return true;
         case GGML_OP_GATED_DELTA_NET:
             return has_simdgroup_reduction && op->src[2]->ne[0] % 32 == 0;
+        case GGML_OP_GATED_DELTA_NET_BACK:
+            return has_simdgroup_reduction && ggml_metal_op_gated_delta_net_back_supported(op);
         case GGML_OP_SOLVE_TRI:
         case GGML_OP_MUL_MAT:
         case GGML_OP_MUL_MAT_ID:
diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h
index 6359d2bf0175..adc883e20f1d 100644
--- a/ggml/src/ggml-metal/ggml-metal-impl.h
+++ b/ggml/src/ggml-metal/ggml-metal-impl.h
@@ -991,6 +991,25 @@ typedef struct {
     uint64_t nb3;
 } ggml_metal_kargs_gated_delta_net;
 
+// GATED_DELTA_NET_BACK: element strides (floats) of q/k/v, the shape, and the segment size
+typedef struct {
+    int64_t H;
+    int64_t T;
+    int64_t n_seqs;
+    int64_t neq1;
+    int64_t neq3;
+    int64_t nek1;
+    int64_t nek3;
+    int64_t sq1, sq2, sq3;
+    int64_t sk1, sk2, sk3;
+    int64_t sv1, sv2, sv3;
+    int64_t C;       // checkpoint segment length (~sqrt(T))
+    int64_t off_dk, off_dv, off_dg, off_db, off_ds; // float offsets of the packed output blocks
+    int32_t kda;
+    int32_t K;
+    float   scale;
+} ggml_metal_kargs_gated_delta_net_back;
+
 typedef struct {
     int32_t  ne00;
     int32_t  ne01;
diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp
index 579631609563..4f41e69d422a 100644
--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp
@@ -369,6 +369,10 @@ static int ggml_metal_op_encode_impl(ggml_metal_op_t ctx, int idx) {
             {
                 n_fuse = ggml_metal_op_rwkv(ctx, idx);
             } break;
+        case GGML_OP_GATED_DELTA_NET_BACK:
+            {
+                n_fuse = ggml_metal_op_gated_delta_net_back(ctx, idx);
+            } break;
         case GGML_OP_GATED_DELTA_NET:
             {
                 n_fuse = ggml_metal_op_gated_delta_net(ctx, idx);
@@ -1870,6 +1874,115 @@ int ggml_metal_op_rwkv(ggml_metal_op_t ctx, int idx) {
     return 1;
 }
 
+// GATED_DELTA_NET_BACK: the checkpoint length and the scratch after dst (see the kernel)
+static int64_t ggml_metal_gdn_back_C(int64_t T) {
+    int64_t C = 1;
+    while (C*C < T) {
+        ++C;
+    }
+    return C;
+}
+
+bool ggml_metal_op_gated_delta_net_back_supported(const ggml_tensor * dst) {
+    for (int s = 0; s < 7; ++s) {
+        if (dst->src[s] == nullptr || dst->src[s]->type != GGML_TYPE_F32) {
+            return false;
+        }
+    }
+    const ggml_tensor * q = dst->src[0], * k = dst->src[1], * v = dst->src[2],
+                      * g = dst->src[3], * b = dst->src[4], * st = dst->src[5], * gr = dst->src[6];
+    const int64_t S_v = v->ne[0];
+    if (S_v != 32 && S_v != 64 && S_v != 128) {
+        return false;
+    }
+    if (q->ne[0] != S_v || k->ne[0] != S_v || !(g->ne[0] == 1 || g->ne[0] == S_v)) {
+        return false;
+    }
+    if (v->ne[3] % q->ne[3] != 0 || v->ne[3] % k->ne[3] != 0 || v->ne[1] % q->ne[1] != 0 || v->ne[1] % k->ne[1] != 0) {
+        return false;
+    }
+    return dst->type == GGML_TYPE_F32 && ggml_is_contiguous(dst) &&
+           ggml_is_contiguous_rows(q) && ggml_is_contiguous_rows(k) && ggml_is_contiguous_rows(v) &&
+           ggml_is_contiguous(g) && ggml_is_contiguous(b) && ggml_is_contiguous(st) && ggml_is_contiguous(gr);
+}
+
+size_t ggml_metal_op_gated_delta_net_back_extra(const ggml_tensor * dst) {
+    const ggml_tensor * v = dst->src[2];
+    const int64_t S_v = v->ne[0], H = v->ne[1], T = v->ne[2], n_seqs = v->ne[3];
+    const int64_t C = ggml_metal_gdn_back_C(T);
+    const int64_t n_ck = (T + C - 1)/C;
+    return (size_t) (n_seqs*H*n_ck*S_v*S_v + n_seqs*H*(C + 3)*S_v*S_v)*sizeof(float) + 256;
+}
+
+int ggml_metal_op_gated_delta_net_back(ggml_metal_op_t ctx, int idx) {
+    ggml_tensor * op = ctx->node(idx);
+    ggml_metal_library_t lib = ctx->lib;
+    ggml_metal_encoder_t enc = ctx->enc;
+
+    GGML_ASSERT(ggml_metal_op_gated_delta_net_back_supported(op));
+
+    const ggml_tensor * q = op->src[0], * k = op->src[1], * v = op->src[2], * g = op->src[3], * b = op->src[4];
+
+    const int64_t S_v = v->ne[0], H = v->ne[1], T = v->ne[2], n_seqs = v->ne[3];
+    const int64_t C    = ggml_metal_gdn_back_C(T);
+    const int64_t n_ck = (T + C - 1)/C;
+
+    ggml_metal_kargs_gated_delta_net_back args = {
+        /*.H      =*/ H,
+        /*.T      =*/ T,
+        /*.n_seqs =*/ n_seqs,
+        /*.neq1   =*/ q->ne[1],
+        /*.neq3   =*/ q->ne[3],
+        /*.nek1   =*/ k->ne[1],
+        /*.nek3   =*/ k->ne[3],
+        /*.sq1    =*/ (int64_t) (q->nb[1]/sizeof(float)), (int64_t) (q->nb[2]/sizeof(float)), (int64_t) (q->nb[3]/sizeof(float)),
+        /*.sk1    =*/ (int64_t) (k->nb[1]/sizeof(float)), (int64_t) (k->nb[2]/sizeof(float)), (int64_t) (k->nb[3]/sizeof(float)),
+        /*.sv1    =*/ (int64_t) (v->nb[1]/sizeof(float)), (int64_t) (v->nb[2]/sizeof(float)), (int64_t) (v->nb[3]/sizeof(float)),
+        /*.C      =*/ C,
+        /*.off_dk =*/ ggml_nelements(q),
+        /*.off_dv =*/ ggml_nelements(q) + ggml_nelements(k),
+        /*.off_dg =*/ ggml_nelements(q) + ggml_nelements(k) + ggml_nelements(v),
+        /*.off_db =*/ ggml_nelements(q) + ggml_nelements(k) + ggml_nelements(v) + ggml_nelements(g),
+        /*.off_ds =*/ ggml_nelements(q) + ggml_nelements(k) + ggml_nelements(v) + ggml_nelements(g) + ggml_nelements(b),
+        /*.kda    =*/ g->ne[0] == S_v ? 1 : 0,
+        /*.K      =*/ ggml_get_op_params_i32(op, 0),
+        /*.scale  =*/ 1.0f/sqrtf((float) S_v),
+    };
+
+    ggml_metal_buffer_id bid_dst  = ggml_metal_get_buffer_id(op);
+    ggml_metal_buffer_id bid_ckpt = bid_dst;
+    bid_ckpt.offs += GGML_PAD(ggml_nbytes(op), 16);
+    ggml_metal_buffer_id bid_seg  = bid_ckpt;
+    bid_seg.offs  += (size_t) n_seqs*H*n_ck*S_v*S_v*sizeof(float);
+
+    // dq/dk accumulate across broadcast heads: zero the packed output first
+    {
+        const uint64_t n = (uint64_t) ggml_nelements(op);
+        auto pz = ggml_metal_library_get_pipeline_gated_delta_net_back_zero(lib);
+        ggml_metal_encoder_set_pipeline(enc, pz);
+        ggml_metal_encoder_set_buffer  (enc, bid_dst, 0);
+        ggml_metal_encoder_set_bytes   (enc, (void *) &n, sizeof(n), 1);
+        const int nth = 256;
+        ggml_metal_encoder_dispatch_threadgroups(enc, (int) ((n + nth - 1)/nth), 1, 1, nth, 1, 1);
+        ggml_metal_op_concurrency_reset(ctx);
+    }
+
+    auto pipeline = ggml_metal_library_get_pipeline_gated_delta_net_back(lib, op);
+    int ida = 0;
+    ggml_metal_encoder_set_pipeline(enc, pipeline);
+    ggml_metal_encoder_set_bytes   (enc, &args, sizeof(args), ida++);
+    for (int s = 0; s < 7; ++s) {
+        ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(op->src[s]), ida++);
+    }
+    ggml_metal_encoder_set_buffer  (enc, bid_dst,  ida++);
+    ggml_metal_encoder_set_buffer  (enc, bid_ckpt, ida++);
+    ggml_metal_encoder_set_buffer  (enc, bid_seg,  ida++);
+    ggml_metal_encoder_dispatch_threadgroups(enc, (int) H, (int) n_seqs, 1, (int) S_v, 1, 1);
+    ggml_metal_op_concurrency_reset(ctx);
+
+    return 1;
+}
+
 int ggml_metal_op_gated_delta_net(ggml_metal_op_t ctx, int idx) {
     ggml_tensor * op = ctx->node(idx);
 
diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h
index a8a89693aa0c..e76003f71bca 100644
--- a/ggml/src/ggml-metal/ggml-metal-ops.h
+++ b/ggml/src/ggml-metal/ggml-metal-ops.h
@@ -62,6 +62,9 @@ int ggml_metal_op_ssm_conv          (ggml_metal_op_t ctx, int idx);
 int ggml_metal_op_ssm_scan          (ggml_metal_op_t ctx, int idx);
 int ggml_metal_op_rwkv              (ggml_metal_op_t ctx, int idx);
 int ggml_metal_op_gated_delta_net   (ggml_metal_op_t ctx, int idx);
+int    ggml_metal_op_gated_delta_net_back(ggml_metal_op_t ctx, int idx);
+bool   ggml_metal_op_gated_delta_net_back_supported(const struct ggml_tensor * op);
+size_t ggml_metal_op_gated_delta_net_back_extra(const struct ggml_tensor * op);
 int ggml_metal_op_solve_tri         (ggml_metal_op_t ctx, int idx);
 int ggml_metal_op_set               (ggml_metal_op_t ctx, int idx);
 int ggml_metal_op_cpy               (ggml_metal_op_t ctx, int idx);
diff --git a/ggml/src/ggml-metal/ggml-metal.cpp b/ggml/src/ggml-metal/ggml-metal.cpp
index fb829adb2605..0ae45e3d7bc3 100644
--- a/ggml/src/ggml-metal/ggml-metal.cpp
+++ b/ggml/src/ggml-metal/ggml-metal.cpp
@@ -237,6 +237,10 @@ static size_t ggml_backend_metal_buffer_type_get_alloc_size(ggml_backend_buffer_
             {
                 res += ggml_metal_op_out_prod_extra(tensor);
             } break;
+        case GGML_OP_GATED_DELTA_NET_BACK:
+            {
+                res += ggml_metal_op_gated_delta_net_back_extra(tensor);
+            } break;
         case GGML_OP_CROSS_ENTROPY_LOSS:
             {
                 res += ggml_nrows(tensor->src[0])*sizeof(float); // the per-row losses before the sum
diff --git a/ggml/src/ggml-metal/kernels/gated_delta_net.metal b/ggml/src/ggml-metal/kernels/gated_delta_net.metal
index 8422d8e29f8b..aa2ba4172f88 100644
--- a/ggml/src/ggml-metal/kernels/gated_delta_net.metal
+++ b/ggml/src/ggml-metal/kernels/gated_delta_net.metal
@@ -248,3 +248,275 @@ template [[host_name("kernel_gated_delta_net_f32_1")]] kernel kernel_gated_delta
 template [[host_name("kernel_gated_delta_net_f32_2")]] kernel kernel_gated_delta_net_t kernel_gated_delta_net_impl<float2, 2>;
 template [[host_name("kernel_gated_delta_net_f32_4")]] kernel kernel_gated_delta_net_t kernel_gated_delta_net_impl<float4, 4>;
 #endif
+
+
+// GATED_DELTA_NET_BACK — the Metal twin of the CPU reference (ggml-cpu ops.cpp) and the CUDA
+// kernel (ggml-cuda/gated_delta_net_back.cu), same math, same layout, same checkpointing.
+// One threadgroup per (head, sequence), S_v threads; thread j owns row j of the state
+// (kept transposed, M[j][i] = S[i][j]):
+//   S' = a (.) S_prev (a = exp(g), scalar or per-row KDA),  u = S'^T k,  delta = beta (v - u),
+//   S = S' + k delta^T,  o = scale S^T q
+// The forward saves a checkpoint every C ~ sqrt(T) tokens; the backward walks the segments
+// last to first, re-running each forward from its checkpoint, carrying dS. Row state lives in
+// device scratch (after dst: 2 x S_v^2 floats is 128 KB at S_v = 128, above threadgroup
+// memory); only per-token vectors and the column accumulators are in threadgroup memory.
+// Output packed as [dq | dk | dv | dg | dbeta | dstate], each dense in its input's order;
+// dq/dk accumulate across the heads that share them (broadcast), so they are zeroed first
+// and added with device atomics. Column sums inside a threadgroup (dq, dk, the KDA dg) are a
+// column pass: thread i walks column i of the rows the other threads wrote, after a device
+// barrier (Metal has no threadgroup float atomics).
+
+kernel void kernel_gated_delta_net_back_zero(
+        device float * dst,
+        constant uint64_t & n,
+        uint tpig[[thread_position_in_grid]]) {
+    if (tpig < n) {
+        dst[tpig] = 0.0f;
+    }
+}
+
+template<short S_v>
+kernel void kernel_gated_delta_net_back_impl(
+        constant ggml_metal_kargs_gated_delta_net_back & args,
+        device const float * q,
+        device const float * k,
+        device const float * v,
+        device const float * g,
+        device const float * beta,
+        device const float * s0,
+        device const float * grad,
+        device       float * out,
+        device       float * ckpt,
+        device       float * seg,
+        uint3 tgpig[[threadgroup_position_in_grid]],
+        ushort tiitg[[thread_index_in_threadgroup]],
+        ushort sgitg[[simdgroup_index_in_threadgroup]],
+        ushort tiisg[[thread_index_in_simdgroup]]) {
+    constexpr int SS  = S_v*S_v;
+    constexpr int NSG = S_v/32 > 0 ? S_v/32 : 1;
+
+    threadgroup float a_s[S_v], q_s[S_v], k_s[S_v], v_s[S_v], gy_s[S_v];
+    threadgroup float delta_s[S_v], u_s[S_v], dd_s[S_v], du_s[S_v];
+    threadgroup float red[NSG];
+
+    const int     j   = tiitg;
+    const int64_t iv1 = tgpig.x; // head
+    const int64_t iv3 = tgpig.y; // sequence
+
+    const int64_t H = args.H, T = args.T, C = args.C;
+    const int64_t rq3 = args.n_seqs/args.neq3;
+    const int64_t rk3 = args.n_seqs/args.nek3;
+    const bool    kda = args.kda != 0;
+    const int64_t neg0 = kda ? S_v : 1;
+
+    device float * dq = out;
+    device float * dk = out + args.off_dk;
+    device float * dv = out + args.off_dv;
+    device float * dg = out + args.off_dg;
+    device float * db = out + args.off_db;
+    device float * ds = out + args.off_ds;
+
+    device atomic_float * dq_at = (device atomic_float *) dq;
+    device atomic_float * dk_at = (device atomic_float *) dk;
+
+    const int64_t n_ck = (T + C - 1)/C;
+    ckpt += (iv3*H + iv1)*n_ck*SS;
+    seg  += (iv3*H + iv1)*(C + 3)*SS;
+    device float * Mp = seg + (C + 1)*SS;
+    device float * dM = Mp + SS;
+
+    device const float * s0_blk = s0 + iv3*H*SS + iv1*SS;
+    const int64_t snap = (int64_t) SS*H*args.n_seqs;
+    device const float * gy_all = grad;
+    device const float * gs_all = grad + (int64_t) S_v*H*T*args.n_seqs;
+
+    // a sum over the threadgroup, valid in every thread
+    auto tg_sum = [&](float x) {
+        float r = simd_sum(x);
+        if (NSG > 1) {
+            threadgroup_barrier(mem_flags::mem_threadgroup);
+            if (tiisg == 0) {
+                red[sgitg] = r;
+            }
+            threadgroup_barrier(mem_flags::mem_threadgroup);
+            r = 0.0f;
+            for (int s = 0; s < NSG; ++s) {
+                r += red[s];
+            }
+        }
+        return r;
+    };
+
+    auto q_at = [&](int64_t t) { return q + (iv3/rq3)*args.sq3 + t*args.sq2 + (iv1 % args.neq1)*args.sq1; };
+    auto k_at = [&](int64_t t) { return k + (iv3/rk3)*args.sk3 + t*args.sk2 + (iv1 % args.nek1)*args.sk1; };
+    auto v_at = [&](int64_t t) { return v + iv3*args.sv3 + t*args.sv2 + iv1*args.sv1; };
+    auto g_at = [&](int64_t t) { return g + (iv3*T + t)*H*neg0 + iv1*neg0; };
+    auto b_at = [&](int64_t t) { return beta[(iv3*T + t)*H + iv1]; };
+
+    // one forward step of row j: Mnext_j = a (.) Mprev_j + k delta_j
+    auto fwd_step = [&](device const float * Mprev, device float * Mnext, int64_t t) {
+        device const float * g_t = g_at(t);
+        a_s[j] = exp(kda ? g_t[j] : g_t[0]);
+        k_s[j] = k_at(t)[j];
+        v_s[j] = v_at(t)[j];
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        const float b = b_at(t);
+        float u = 0.0f;
+        for (int i = 0; i < S_v; ++i) {
+            u += Mprev[i]*a_s[i]*k_s[i];
+        }
+        const float delta = (v_s[j] - u)*b;
+        for (int i = 0; i < S_v; ++i) {
+            Mnext[i] = Mprev[i]*a_s[i] + k_s[i]*delta;
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+    };
+
+    // ---- forward, checkpointing every C tokens
+    for (int i = 0; i < S_v; ++i) {
+        seg[j*S_v + i] = s0_blk[j*S_v + i];
+    }
+    for (int64_t t = 0; t < T; ++t) {
+        device float * cur = seg + (t & 1)*SS + j*S_v;
+        device float * nxt = seg + ((t + 1) & 1)*SS + j*S_v;
+        if (t % C == 0) {
+            device float * ck = ckpt + (t/C)*SS + j*S_v;
+            for (int i = 0; i < S_v; ++i) {
+                ck[i] = cur[i];
+            }
+        }
+        fwd_step(cur, nxt, t);
+    }
+
+    // ---- backward, carrying dM
+    for (int i = 0; i < S_v; ++i) {
+        dM[j*S_v + i] = 0.0f;
+    }
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+
+    for (int64_t s = n_ck - 1; s >= 0; --s) {
+        const int64_t t0 = s*C;
+        const int64_t t1 = t0 + C < T ? t0 + C : T;
+        {
+            device const float * ck = ckpt + s*SS + j*S_v;
+            for (int i = 0; i < S_v; ++i) {
+                seg[j*S_v + i] = ck[i];
+            }
+            for (int64_t tt = t0; tt < t1; ++tt) {
+                fwd_step(seg + (tt - t0)*SS + j*S_v, seg + (tt - t0 + 1)*SS + j*S_v, tt);
+            }
+        }
+        threadgroup_barrier(mem_flags::mem_device);
+        for (int64_t t = t1 - 1; t >= t0; --t) {
+            device const float * g_t = g_at(t);
+            a_s[j]  = exp(kda ? g_t[j] : g_t[0]);
+            q_s[j]  = q_at(t)[j];
+            k_s[j]  = k_at(t)[j];
+            v_s[j]  = v_at(t)[j];
+            gy_s[j] = gy_all[((iv3*T + t)*H + iv1)*S_v + j];
+            threadgroup_barrier(mem_flags::mem_threadgroup);
+
+            const float b = b_at(t);
+            device float       * dMj = dM + j*S_v;
+            device float       * Mpj = Mp + j*S_v;
+            device const float * M1  = seg + (t - t0 + 1)*SS;      // S after token t (all rows)
+            device const float * M0j = seg + (t - t0)*SS + j*S_v;  // S before token t, row j
+
+            // row pass A: the snapshot gradient, the output's dS, and S', u, delta
+            const int64_t slot = T - 1 - t;
+            if (slot < args.K) {
+                device const float * gs = gs_all + slot*snap + (iv3*H + iv1)*SS + j*S_v;
+                for (int i = 0; i < S_v; ++i) {
+                    dMj[i] += gs[i];
+                }
+            }
+            {
+                const float gyj = gy_s[j];
+                for (int i = 0; i < S_v; ++i) {
+                    dMj[i] += args.scale*gyj*q_s[i];
+                }
+            }
+            float u = 0.0f;
+            for (int i = 0; i < S_v; ++i) {
+                const float mp = M0j[i]*a_s[i];
+                Mpj[i] = mp;
+                u += mp*k_s[i];
+            }
+            u_s[j]     = u;
+            delta_s[j] = (v_s[j] - u)*b;
+
+            // row pass B: dd_j = sum_i dM[j][i] k_i ; dv ; du
+            {
+                float dd = 0.0f;
+                for (int i = 0; i < S_v; ++i) {
+                    dd += dMj[i]*k_s[i];
+                }
+                dd_s[j] = dd;
+                dv[((iv3*T + t)*H + iv1)*S_v + j] = b*dd;
+                du_s[j] = -b*dd;
+            }
+            threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup);
+
+            // column pass C (thread j = column j): dq_j = scale sum_r S[r][j] gy_r ;
+            // dk_j = sum_r dM[r][j] delta_r + S'[r][j] du_r   (dM before the du update)
+            {
+                float cq = 0.0f, ck = 0.0f;
+                for (int r = 0; r < S_v; ++r) {
+                    cq += M1[r*S_v + j]*gy_s[r];
+                    ck += dM[r*S_v + j]*delta_s[r] + Mp[r*S_v + j]*du_s[r];
+                }
+                atomic_fetch_add_explicit(&dq_at[(((iv3/rq3)*T + t)*args.neq1 + iv1 % args.neq1)*S_v + j], args.scale*cq, memory_order_relaxed);
+                atomic_fetch_add_explicit(&dk_at[(((iv3/rk3)*T + t)*args.nek1 + iv1 % args.nek1)*S_v + j], ck, memory_order_relaxed);
+            }
+            {
+                const float dbeta = tg_sum(dd_s[j]*(v_s[j] - u_s[j]));
+                if (j == 0) {
+                    db[(iv3*T + t)*H + iv1] = dbeta;
+                }
+            }
+            threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup);
+
+            // row pass D: u = S'^T k -> dM[j][i] += k_i du_j ; the scalar gate's row sum
+            float dg_local = 0.0f;
+            {
+                const float duj = du_s[j];
+                for (int i = 0; i < S_v; ++i) {
+                    dMj[i] += k_s[i]*duj;
+                    dg_local += dMj[i]*Mpj[i];
+                }
+            }
+            threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup);
+
+            // S' = a (.) S_prev: dg = sum dS' (.) S' (column j for KDA, a scalar otherwise)
+            if (kda) {
+                float cg = 0.0f;
+                for (int r = 0; r < S_v; ++r) {
+                    cg += dM[r*S_v + j]*Mp[r*S_v + j];
+                }
+                dg[((iv3*T + t)*H + iv1)*S_v + j] = cg;
+            } else {
+                const float dg0 = tg_sum(dg_local);
+                if (j == 0) {
+                    dg[(iv3*T + t)*H + iv1] = dg0;
+                }
+            }
+            threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup);
+
+            // row pass E: dS_prev = a (.) dS'
+            for (int i = 0; i < S_v; ++i) {
+                dMj[i] *= a_s[i];
+            }
+            threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup);
+        }
+    }
+
+    for (int i = 0; i < S_v; ++i) {
+        ds[(iv3*H + iv1)*SS + j*S_v + i] = dM[j*S_v + i];
+    }
+}
+
+typedef decltype(kernel_gated_delta_net_back_impl<32>) kernel_gated_delta_net_back_t;
+
+template [[host_name("kernel_gated_delta_net_back_f32_32")]]  kernel kernel_gated_delta_net_back_t kernel_gated_delta_net_back_impl<32>;
+template [[host_name("kernel_gated_delta_net_back_f32_64")]]  kernel kernel_gated_delta_net_back_t kernel_gated_delta_net_back_impl<64>;
+template [[host_name("kernel_gated_delta_net_back_f32_128")]] kernel kernel_gated_delta_net_back_t kernel_gated_delta_net_back_impl<128>;