sync 6fdf6301e2bb

Browse files

Files changed (7) hide show

README.md +93 -0
build/webgpu/bench.json +0 -0
build/webgpu/gated-delta-net-gate.wgsl.jinja +170 -0
build/webgpu/gated-delta-net.wgsl.jinja +384 -0
build/webgpu/manifest.json +1093 -0
build/webgpu/metadata.json +47 -0
build/webgpu/test.json +0 -0

README.md CHANGED Viewed

@@ -1,3 +1,96 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.GatedDeltaNet
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Token-major gated delta network with an explicit recurrent state: `S_t = exp(g_t) * S_{t-1} + k_t * (beta_t * (v_t - exp(g_t) * S_{t-1}^T k_t))^T` and `o_t = scale * S_t^T q_t` per value head; `update_rule` selects which terms survive. Query, key and value are `(total_tokens, heads, head_size)` or `(batch, sequence, heads, head_size)`. `num_heads_v` is a positive multiple of `num_heads_q`, and value head `h` reads query head `h * num_heads_q / num_heads_v`. The V-major state and the gate fusions are float32. Bfloat16, per-key-dimension decay and compact capture are not implemented.
+See the [ONNX Runtime `GatedDeltaNet` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.GatedDeltaNet) for the reference semantics.
+## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `queryT` | `query` | `T` | — | — | Query vectors, `(total_tokens, num_heads_q, head_size_qk)` or `(batch_size, sequence_length, num_heads_q, head_size_qk)`. The two spellings share one memory layout; the rank-4 form lets an exporter keep a static reshape target. | required |
+| `keyT` | `key` | `T` | — | — | Key vectors with exactly the shape of `query`. The delta rules need L2-normalized keys, either from upstream or from `qk_l2_norm=1`; without them the per-chunk system is ill-conditioned and the recurrence diverges. | required |
+| `valueT` | `value` | `T` | — | — | Value vectors, `(total_tokens, num_heads_v, head_size_v)`, sharing the leading token axes of `query`. | required |
+| `cuSeqlensT` | `cu_seqlens` | `TI` | `1` | — | Exclusive prefix sums of the per-request token counts, `(batch_size + 1)`. Requires the rank-3 spelling. Offsets are clamped on device to `[0, total_tokens]`, so a decreasing or out-of-range entry yields an empty request rather than an out-of-bounds read. Absent means uniform packing. | optional |
+| `decayT` | `decay` | `TS` | — | — | Log-space decay, `(..., num_heads_v)` over the same leading token axes as `query`, always float32. Present exactly for `update_rule` `gated` and `gated_delta`. The per-key-dimension spelling `(..., num_heads_v, head_size_qk)` is not implemented. | optional |
+| `betaT` | `beta` | `TS` | — | — | Update rate, `(..., num_heads_v)` over the same leading token axes as `query`, always float32. Present exactly for `update_rule` `delta` and `gated_delta`. | optional |
+| `initialStateT` | `initial_state` | `TS` | `4` | — | Incoming recurrent state, V-major `(batch_size, num_heads_v, head_size_v, head_size_qk)`, always float32. Absent means a zero state. It is a separate allocation from `final_state`: WebGPU usage-scope validation tracks whole buffers, so binding one buffer as both read-only and read-write invalidates the command buffer. | optional |
+| `aLogT` | `a_log` | `TS` | `1` | — | Per-head `A_log`, `(num_heads_v)`, always float32. Present exactly when `gate_activation` is `qwen`. | optional |
+| `dtBiasT` | `dt_bias` | `TS` | `1` | — | Per-head gate bias, `(num_heads_v)`, always float32. Present exactly when `gate_activation` is `qwen`. | optional |
+## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `outputT` | `output` | `T` | same as `queryT` | derived | Attention output, `(total_tokens, num_heads_v, head_size_v)` or `(batch_size, sequence_length, num_heads_v, head_size_v)`. `num_heads_v` is the larger head count because it is a positive multiple of `num_heads_q`. | required |
+| `finalStateT` | `final_state` | `TS` | `4` | derived | State after the last token of each request, V-major `(batch_size, num_heads_v, head_size_v, head_size_qk)`, always float32. A request with no tokens copies its incoming state through. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `beta_activation` | `"none"` | `none` treats `beta` as the effective update rate; `sigmoid` applies a logistic sigmoid to it in float32. Default is `none`. |
+| `chunk_size` | `64` | Accepted scheduling hint for a chunk-parallel prefill algorithm. Every value is accepted and none changes the result. Default is `64`. |
+| `gate_activation` | `"none"` | `none` treats `decay` as the effective log-space decay. `qwen` computes `-exp(a_log) * Softplus(decay + dt_bias)` in float32 and requires both `a_log` and `dt_bias`. Default is `none`. |
+| `qk_l2_norm` | `0` | Set to `1` to L2-normalize each query and key head vector, with epsilon `1e-12`, before the recurrence. Default is `0`. |
+| `scale` | `0` | Output scaling factor applied to `S_t^T q_t`. The default `0.0` selects `1 / sqrt(head_size_qk)`. |
+| `state_update_capacity` | `0` | Capacity for compact transition capture. The default and only supported value is `0`; compact `state_update` capture is not implemented. |
+| `update_rule` | `"gated_delta"` | Which terms of the recurrence survive: `linear` drops the decay and the delta retrieval, `gated` keeps only the decay, `delta` keeps only the retrieval, and `gated_delta` keeps both. Default is `gated_delta`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+| `TS` | `float32` |
+| `TI` | `int32` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
+- [`gated-delta-net-gate.wgsl.jinja`](build/webgpu/gated-delta-net-gate.wgsl.jinja)
+- [`gated-delta-net.wgsl.jinja`](build/webgpu/gated-delta-net.wgsl.jinja)
+## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.GatedDeltaNet", { version: 1 });
+const { outputT, finalStateT } = await kernel({
+  queryT: { data: queryTData, shape: [4, 1, 12] },
+  keyT: { data: keyTData, shape: [4, 1, 12] },
+  valueT: { data: valueTData, shape: [4, 2, 9] },
+  decayT: { data: decayTData, shape: [4, 2] },
+  betaT: { data: betaTData, shape: [4, 2] },
+  initialStateT: { data: initialStateTData, shape: [1, 2, 9, 12] },
+});
+```

build/webgpu/bench.json ADDED Viewed

The diff for this file is too large to render. See raw diff

build/webgpu/gated-delta-net-gate.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,170 @@

+{{ env.wgsl.resourceDeclarations }}
+// Per-token/head gate and q/k statistics shared by every value tile.
+// Paired updates also read previous-key cross products. All statistics are f32.
+{% if hasDecayInput or hasBetaInput %}
+const NUM_Q_HEADS: u32 = {{ numQHeads }}u;
+const NUM_V_HEADS: u32 = {{ numVHeads }}u;
+const HEADS_PER_Q: u32 = {{ headsPerQ }}u;
+{% endif %}
+// Words, not elements: vec4 where the key dimension is a multiple of four.
+const DK_WORDS: u32 = {{ dkWords }}u;
+const WG: u32 = {{ gateLanes }}u;
+{% if useL2Norm %}
+// The upstream L2 normalization epsilon is 1e-12, added to the sum of squares.
+const L2_EPS: f32 = 1e-12;
+{% endif %}
+// One lane's contribution to a key-dimension reduction. The vec4 form is reached
+// wherever the key dimension is a multiple of four; the scalar form carries every
+// other head size with the same schedule.
+fn word_dot(a: {{ wordF32 }}, b: {{ wordF32 }}) -> f32 {
+{% if useVec4 %}
+  return dot(a, b);
+{% else %}
+  return a * b;
+{% endif %}
+}
+// The three key-dimension sums share one staging array and one fold.
+var<workgroup> partials: array<vec3<f32>, WG>;
+{% if pairTokens %}
+var<workgroup> crossPartials: array<vec2<f32>, WG>;
+{% endif %}
+{% if sigmoidBeta and hasBetaInput %}
+fn sigmoid_safe(x: f32) -> f32 {
+  if (x >= 0.0) {
+    let z = exp(-x);
+    return 1.0 / (1.0 + z);
+  }
+  let z = exp(x);
+  return z / (1.0 + z);
+}
+{% endif %}
+{% if qwen %}
+// WGSL has no log1p builtin, and log(1.0 + y) loses relative accuracy for small y: rounding
+// 1.0 + y costs up to half an ulp of 1.0, which is enormous next to y itself. Below the
+// crossover the series is the accurate form -- its truncation is y^7 / 8, under 6e-9 relative
+// at y = 0.125 -- and above it the same half-ulp is under 6e-7 relative because log1p(y) has
+// grown past 0.118. The algebraic (1.0 + y) - 1.0 correction is not usable here: a backend is
+// free to fold that subtraction back to y, which silently restores the error it removes.
+fn log1p_pos(y: f32) -> f32 {
+  if (y < 0.125) {
+    return y * (1.0 + y * (-0.5 + y * (0.33333334 + y * (-0.25 + y * (0.2 + y * (-0.16666667 + y * 0.14285714))))));
+  }
+  return log(1.0 + y);
+}
+fn softplus(x: f32) -> f32 {
+  // Both arms feed log1p an exp() of a non-positive argument, so the operand stays in (0, 1]
+  // and never overflows; for large x the increment rounds away and the result is exactly x,
+  // which is the correct limit.
+  if (x > 0.0) {
+    return x + log1p_pos(exp(-x));
+  }
+  return log1p_pos(exp(x));
+}
+{% endif %}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+) {
+  // 2D-folded (token * query head) index: wg.y carries the high bits past the
+  // per-axis dispatch fold width.
+  let unit = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
+  if (unit >= params.units) {
+    return;
+  }
+  let qk_row = unit * DK_WORDS;
+  var acc = vec3<f32>(0.0);
+{% if pairTokens %}
+  var crossAcc = vec2<f32>(0.0);
+{% endif %}
+  for (var i = lid.x; i < DK_WORDS; i = i + WG) {
+    let kw = {{ wordF32 }}(key[qk_row + i]);
+    let qw = {{ wordF32 }}(query[qk_row + i]);
+    acc = acc + vec3<f32>(word_dot(kw, qw), word_dot(kw, kw), word_dot(qw, qw));
+{% if pairTokens %}
+    if (unit >= NUM_Q_HEADS) {
+      let previous = {{ wordF32 }}(key[qk_row - NUM_Q_HEADS * DK_WORDS + i]);
+      crossAcc = crossAcc + vec2<f32>(word_dot(previous, kw), word_dot(previous, qw));
+    }
+{% endif %}
+  }
+  partials[lid.x] = acc;
+{% if pairTokens %}
+  crossPartials[lid.x] = crossAcc;
+{% endif %}
+  workgroupBarrier();
+  // WG is a power of two, so the halving stride is uniform across the workgroup and
+  // every lane reaches every barrier.
+  var stride = WG / 2u;
+  loop {
+    if (stride == 0u) {
+      break;
+    }
+    if (lid.x < stride) {
+      partials[lid.x] = partials[lid.x] + partials[lid.x + stride];
+{% if pairTokens %}
+      crossPartials[lid.x] = crossPartials[lid.x] + crossPartials[lid.x + stride];
+{% endif %}
+    }
+    workgroupBarrier();
+    stride = stride / 2u;
+  }
+  let sums = partials[0];
+{% if useL2Norm %}
+  let inv_k = inverseSqrt(sums.y + L2_EPS);
+  let inv_q = inverseSqrt(sums.z + L2_EPS);
+{% else %}
+  let inv_k = 1.0;
+  let inv_q = 1.0;
+{% endif %}
+  if (lid.x == 0u) {
+    let qk_out = params.qkBase + unit * {{ 5 if pairTokens else 3 }}u;
+    gate[qk_out] = sums.x * inv_k * inv_q;
+    gate[qk_out + 1u] = inv_k;
+    gate[qk_out + 2u] = inv_q;
+{% if pairTokens %}
+    gate[qk_out + 3u] = crossPartials[0].x;
+    gate[qk_out + 4u] = crossPartials[0].y;
+{% endif %}
+  }
+{% if hasDecayInput or hasBetaInput %}
+  // Each query head serves HEADS_PER_Q consecutive value heads, so this unit owns
+  // exactly that block of region A.
+  let head_q = unit % NUM_Q_HEADS;
+  let token = unit / NUM_Q_HEADS;
+  for (var h = lid.x; h < HEADS_PER_Q; h = h + WG) {
+    let head_v = head_q * HEADS_PER_Q + h;
+    let item = token * NUM_V_HEADS + head_v;
+{% if hasDecayInput %}
+    var raw_decay = decay[item];
+{% if qwen %}
+    raw_decay = -exp(a_log[head_v]) * softplus(raw_decay + dt_bias[head_v]);
+{% endif %}
+    gate[item * 2u] = exp(raw_decay);
+{% else %}
+    // Rules without a decay term contract by exactly one per step.
+    gate[item * 2u] = 1.0;
+{% endif %}
+{% if hasBetaInput %}
+    let raw_beta = beta[item];
+{% if sigmoidBeta %}
+    gate[item * 2u + 1u] = sigmoid_safe(raw_beta);
+{% else %}
+    gate[item * 2u + 1u] = raw_beta;
+{% endif %}
+{% else %}
+    // Rules without a delta retrieval take the value vector whole.
+    gate[item * 2u + 1u] = 1.0;
+{% endif %}
+  }
+{% endif %}
+}

build/webgpu/gated-delta-net.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,384 @@

+{% set comps = ["x", "y", "z", "w"] %}
+{% set skBase = tileV %}
+{% set exactWords = wordsPerLane * lanes == dkWords %}
+{% set foldSteps = log2ceil(lanes) %}
+{% macro fold_values(foldSlots) %}
+{% if lanes > 1 %}
+{% if useSubgroups %}
+    // Segmented butterfly over the low lane bits only, so each aligned group of
+    // LANES lanes folds its own value tile and one subgroup carries several tiles
+    // with no barrier and no workgroup memory. Selection admits this arm only when
+    // LANES divides every subgroup width the adapter may use.
+{% for step in range(foldSteps) %}
+{% for s in range(foldSlots) %}
+    f{{ s }} = f{{ s }} + subgroupShuffleXor(f{{ s }}, 1u << {{ step }}u);
+{% endfor %}
+{% endfor %}
+{% else %}
+    // The previous token's results were read by every lane of the group, so the
+    // staging array cannot be rewritten until all of them have passed this point.
+    workgroupBarrier();
+{% for s in range(foldSlots) %}
+    fold[{{ s }}u * WG + lid.x] = f{{ s }};
+{% endfor %}
+    workgroupBarrier();
+    // LANES is a power of two, so the halving stride is uniform across the whole
+    // workgroup and every lane reaches every barrier.
+    var stride = LANES / 2u;
+    loop {
+      if (stride == 0u) {
+        break;
+      }
+      if (lane < stride) {
+{% for s in range(foldSlots) %}
+        fold[{{ s }}u * WG + lid.x] = fold[{{ s }}u * WG + lid.x] + fold[{{ s }}u * WG + lid.x + stride];
+{% endfor %}
+      }
+      workgroupBarrier();
+      stride = stride / 2u;
+    }
+{% for s in range(foldSlots) %}
+    f{{ s }} = fold[{{ s }}u * WG + v_group * LANES];
+{% endfor %}
+{% endif %}
+{% endif %}
+{% endmacro %}
+{% if pairTokens %}
+{% macro paired_recurrence() %}
+  // Two consecutive updates share one cross-lane fold. Both projections and
+  // retrievals are formed against the incoming state; the second token applies
+  // the first token's rank-one correction through precomputed cross products.
+  for (var t = token_begin; t + 1u < token_end; t = t + 2u) {
+{% for p in range(2) %}
+    let qk_row{{ p }} = ((t + {{ p }}u) * NUM_Q_HEADS + head_q) * DK_WORDS;
+{% for w in range(wordsPerLane) %}
+{% if exactWords %}
+    let kw{{ p }}_{{ w }} = {{ wordF32 }}(key[qk_row{{ p }} + r{{ w }}]);
+    let qw{{ p }}_{{ w }} = {{ wordF32 }}(query[qk_row{{ p }} + r{{ w }}]);
+{% else %}
+    var kw{{ p }}_{{ w }} = {{ wordF32 }}(0.0);
+    var qw{{ p }}_{{ w }} = {{ wordF32 }}(0.0);
+    if (word_ok{{ w }}) {
+      kw{{ p }}_{{ w }} = {{ wordF32 }}(key[qk_row{{ p }} + r{{ w }}]);
+      qw{{ p }}_{{ w }} = {{ wordF32 }}(query[qk_row{{ p }} + r{{ w }}]);
+    }
+{% endif %}
+{% endfor %}
+{% endfor %}
+{% for s in range(foldSlots) %}
+    var f{{ s }} = vec4<f32>(0.0);
+{% endfor %}
+{% for p in range(2) %}
+{% for j in range(tileV) %}
+{% set qi = p * 2 * tileV + j %}
+{% set ki = qi + skBase %}
+    f{{ (qi / 4) | int }}.{{ comps[qi % 4] }} = {% for w in range(wordsPerLane) %}{{ "+ " if not loop.first else "" }}word_dot(s{{ j }}_{{ w }}, qw{{ p }}_{{ w }}) {% endfor %};
+    f{{ (ki / 4) | int }}.{{ comps[ki % 4] }} = {% for w in range(wordsPerLane) %}{{ "+ " if not loop.first else "" }}word_dot(s{{ j }}_{{ w }}, kw{{ p }}_{{ w }}) {% endfor %};
+{% endfor %}
+{% endfor %}
+{{ fold_values(foldSlots) }}
+{% for p in range(2) %}
+    let qk_base{{ p }} = params.qkBase + ((t + {{ p }}u) * NUM_Q_HEADS + head_q) * 5u;
+    let kq{{ p }} = gate[qk_base{{ p }}];
+{% if useL2Norm %}
+    let inv_k{{ p }} = gate[qk_base{{ p }} + 1u];
+    let inv_q{{ p }} = gate[qk_base{{ p }} + 2u];
+{% else %}
+    let inv_k{{ p }} = 1.0;
+    let inv_q{{ p }} = 1.0;
+{% endif %}
+    let gate_base{{ p }} = ((t + {{ p }}u) * NUM_V_HEADS + head_v) * 2u;
+    let decay{{ p }} = gate[gate_base{{ p }}];
+    let beta{{ p }} = gate[gate_base{{ p }} + 1u];
+    let value_row{{ p }} = ((t + {{ p }}u) * NUM_V_HEADS + head_v) * HEAD_DIM_V + dv_start;
+{% if vectorValueIO %}
+    var output_values{{ p }} = vec{{ tileV }}<f32>(0.0);
+    var value_values{{ p }} = vec{{ tileV }}<f32>(0.0);
+    if (dv_start < HEAD_DIM_V) { value_values{{ p }} = vec{{ tileV }}<f32>(value[value_row{{ p }} / TILE_V]); }
+{% endif %}
+{% endfor %}
+    let cross_kk = gate[qk_base1 + 3u];
+    let cross_kq = gate[qk_base1 + 4u];
+{% for j in range(tileV) %}
+    var d0_{{ j }} = 0.0;
+    var d1_{{ j }} = 0.0;
+    if (dv_start + {{ j }}u < HEAD_DIM_V) {
+{% for p in range(2) %}
+{% set qi = p * 2 * tileV + j %}
+{% set ki = qi + skBase %}
+      let retrieved{{ p }}_{{ j }} = decay{{ p }} * inv_k{{ p }} * {% if p == 1 %}(decay0 * {% endif %}f{{ (ki / 4) | int }}.{{ comps[ki % 4] }}{% if p == 1 %} + d0_{{ j }} * cross_kk){% endif %};
+      let delta{{ p }}_{{ j }} = beta{{ p }} * ({% if vectorValueIO %}value_values{{ p }}[{{ j }}]{% else %}f32(value[value_row{{ p }} + {{ j }}u]){% endif %} - retrieved{{ p }}_{{ j }});
+      d{{ p }}_{{ j }} = inv_k{{ p }} * delta{{ p }}_{{ j }};
+      let projected{{ p }}_{{ j }} = decay{{ p }} * inv_q{{ p }} * {% if p == 1 %}(decay0 * {% endif %}f{{ (qi / 4) | int }}.{{ comps[qi % 4] }}{% if p == 1 %} + d0_{{ j }} * cross_kq){% endif %};
+{% if vectorValueIO %}
+      output_values{{ p }}[{{ j }}] = (projected{{ p }}_{{ j }} + delta{{ p }}_{{ j }} * kq{{ p }}) * SCALE;
+{% else %}
+      if (lane == {{ j % lanes }}u) {
+        output[value_row{{ p }} + {{ j }}u] = {{ T }}((projected{{ p }}_{{ j }} + delta{{ p }}_{{ j }} * kq{{ p }}) * SCALE);
+      }
+{% endif %}
+{% endfor %}
+    }
+{% endfor %}
+{% if vectorValueIO %}
+    if (dv_start < HEAD_DIM_V && lane == 0u) {
+      output[value_row0 / TILE_V] = vec{{ tileV }}<{{ T }}>(output_values0);
+      output[value_row1 / TILE_V] = vec{{ tileV }}<{{ T }}>(output_values1);
+    }
+{% endif %}
+{% for j in range(tileV) %}
+{% for w in range(wordsPerLane) %}
+    s{{ j }}_{{ w }} = s{{ j }}_{{ w }} * {{ wordF32 }}(decay0) + kw0_{{ w }} * {{ wordF32 }}(d0_{{ j }});
+    s{{ j }}_{{ w }} = s{{ j }}_{{ w }} * {{ wordF32 }}(decay1) + kw1_{{ w }} * {{ wordF32 }}(d1_{{ j }});
+{% endfor %}
+{% endfor %}
+  }
+{% endmacro %}
+{% endif %}
+{% macro recurrence_step(foldSlots, qkStride) %}
+    let qk_row = (t * NUM_Q_HEADS + head_q) * DK_WORDS;
+{% for w in range(wordsPerLane) %}
+{% if exactWords %}
+    let kw{{ w }} = {{ wordF32 }}(key[qk_row + r{{ w }}]);
+    let qw{{ w }} = {{ wordF32 }}(query[qk_row + r{{ w }}]);
+{% else %}
+    var kw{{ w }} = {{ wordF32 }}(0.0);
+    var qw{{ w }} = {{ wordF32 }}(0.0);
+    if (word_ok{{ w }}) {
+      kw{{ w }} = {{ wordF32 }}(key[qk_row + r{{ w }}]);
+      qw{{ w }} = {{ wordF32 }}(query[qk_row + r{{ w }}]);
+    }
+{% endif %}
+{% endfor %}
+{% for s in range(foldSlots) %}
+    var f{{ s }} = vec4<f32>(0.0);
+{% endfor %}
+{% for j in range(tileV) %}
+    f{{ (j / 4) | int }}.{{ comps[j % 4] }} = {% for w in range(wordsPerLane) %}{{ "+ " if not loop.first else "" }}word_dot(s{{ j }}_{{ w }}, qw{{ w }}) {% endfor %};
+{% endfor %}
+{% if useDelta %}
+{% for j in range(tileV) %}
+    f{{ ((skBase + j) / 4) | int }}.{{ comps[(skBase + j) % 4] }} = {% for w in range(wordsPerLane) %}{{ "+ " if not loop.first else "" }}word_dot(s{{ j }}_{{ w }}, kw{{ w }}) {% endfor %};
+{% endfor %}
+{% endif %}
+{{ fold_values(foldSlots) }}
+    // Normalizing q and k is a per-vector SCALE, so the state reductions above run on
+    // the raw vectors and are rescaled here. The state update then absorbs the key's
+    // factor into the delta, and `kq` arrives with both factors already applied.
+    let qk_row_base = params.qkBase + (t * NUM_Q_HEADS + head_q) * {{ qkStride }}u;
+    let kq = gate[qk_row_base];
+{% if useL2Norm %}
+    let inv_k = gate[qk_row_base + 1u];
+    let inv_q = gate[qk_row_base + 2u];
+{% else %}
+    let inv_k = 1.0;
+    let inv_q = 1.0;
+{% endif %}
+{% if hasGate %}
+    let gate_base = (t * NUM_V_HEADS + head_v) * 2u;
+    let decay = gate[gate_base];
+    let beta = gate[gate_base + 1u];
+{% else %}
+    let decay = 1.0;
+    let beta = 1.0;
+{% endif %}
+    let value_row = (t * NUM_V_HEADS + head_v) * HEAD_DIM_V + dv_start;
+{% if vectorValueIO %}
+    // Fold a complete value tile into one contiguous store. The template keeps
+    // scalar I/O for tails and for float32 groups below the vectorization
+    // threshold.
+    var output_values = vec{{ tileV }}<f32>(0.0);
+    var value_values = vec{{ tileV }}<f32>(0.0);
+    if (dv_start < HEAD_DIM_V) { value_values = vec{{ tileV }}<f32>(value[value_row / TILE_V]); }
+{% endif %}
+{% for j in range(tileV) %}
+    var d{{ j }} = 0.0;
+    if (dv_start + {{ j }}u < HEAD_DIM_V) {
+{% if useDelta %}
+      let retrieved{{ j }} = decay * inv_k * f{{ ((skBase + j) / 4) | int }}.{{ comps[(skBase + j) % 4] }};
+{% else %}
+      let retrieved{{ j }} = 0.0;
+{% endif %}
+      let delta{{ j }} = beta * ({% if vectorValueIO %}value_values[{{ j }}]{% else %}f32(value[value_row + {{ j }}u]){% endif %} - retrieved{{ j }});
+      d{{ j }} = inv_k * delta{{ j }};
+{% if vectorValueIO %}
+      let projected{{ j }} = decay * inv_q * f{{ (j / 4) | int }}.{{ comps[j % 4] }};
+      output_values[{{ j }}] = (projected{{ j }} + delta{{ j }} * kq) * SCALE;
+{% else %}
+      if (lane == {{ j % lanes }}u) {
+        let projected{{ j }} = decay * inv_q * f{{ (j / 4) | int }}.{{ comps[j % 4] }};
+        output[value_row + {{ j }}u] = {{ T }}((projected{{ j }} + delta{{ j }} * kq) * SCALE);
+      }
+{% endif %}
+    }
+{% endfor %}
+{% if vectorValueIO %}
+    if (dv_start < HEAD_DIM_V && lane == 0u) {
+      output[value_row / TILE_V] = vec{{ tileV }}<{{ T }}>(output_values);
+    }
+{% endif %}
+{% for j in range(tileV) %}
+{% for w in range(wordsPerLane) %}
+    s{{ j }}_{{ w }} = s{{ j }}_{{ w }} * {{ wordF32 }}(decay) + kw{{ w }} * {{ wordF32 }}(d{{ j }});
+{% endfor %}
+{% endfor %}
+{% endmacro %}
+{% if useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// One sequential pass of the gated delta recurrence, per (request, value head,
+// value tile):
+//
+//   S_t = d_t * S_{t-1} + k_t * (beta_t * (v_t - d_t * S_{t-1}^T k_t))^T
+//   o_t = SCALE * S_t^T q_t
+//
+// The state boundary is V-major, (batch, num_heads_v, head_size_v, head_size_qk)
+// with the key dimension contiguous, so one value channel's whole key slice is a
+// contiguous run. A lane group of LANES lanes owns TILE_V such channels and holds
+// WORDS_PER_LANE words of each in registers; the reductions therefore run along
+// the key dimension, across the lanes of one group.
+//
+// The output is folded into the same pass instead of re-reading the updated
+// state: S_t^T q = d * (S_{t-1}^T q) + delta * (k . q). The retrieval and the
+// output projection are therefore both products with the state BEFORE the update
+// and reduce together, while k . q carries no state and arrives as a scalar.
+const HEAD_DIM_V: u32 = {{ headDimV }}u;
+// Words, not elements: vec4 where the key dimension is a multiple of four.
+const DK_WORDS: u32 = {{ dkWords }}u;
+const NUM_Q_HEADS: u32 = {{ numQHeads }}u;
+const NUM_V_HEADS: u32 = {{ numVHeads }}u;
+// num_heads_v is a positive multiple of num_heads_q, so the query head a value
+// head reads is a BLOCK map: value heads 0..HEADS_PER_Q-1 all read query head 0.
+const HEADS_PER_Q: u32 = {{ headsPerQ }}u;
+const LANES: u32 = {{ lanes }}u;
+const V_GROUPS: u32 = {{ vGroups }}u;
+const WG: u32 = LANES * V_GROUPS;
+const TILE_V: u32 = {{ tileV }}u;
+const DV_BLOCKS: u32 = {{ dvBlocks }}u;
+const SCALE: f32 = {{ outScale }};
+// One lane's contribution to a key-dimension reduction. The vec4 form is reached
+// wherever the key dimension is a multiple of four; the scalar form carries every
+// other head size with the same schedule.
+fn word_dot(a: {{ wordF32 }}, b: {{ wordF32 }}) -> f32 {
+{% if useVec4 %}
+  return dot(a, b);
+{% else %}
+  return a * b;
+{% endif %}
+}
+{% if not useSubgroups and lanes > 1 %}
+// Portable fold: every reduced value of one token is staged in the same array and
+// folded once, so the barrier count is a property of the token rather than of how
+// many quantities it reduces. A lane group owns one contiguous LANES-wide slice of
+// each slot, so groups sharing this workgroup never read one another's partials.
+var<workgroup> fold: array<vec4<f32>, WG * {{ foldSlots }}u>;
+{% endif %}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+{% if useSubgroups %}
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_id) sg_id: u32,
+{% if variableSubgroupFold %}
+  @builtin(subgroup_size) sg_size: u32,
+{% endif %}
+{% else %}
+  @builtin(local_invocation_id) lid: vec3<u32>,
+{% endif %}
+) {
+{% if useSubgroups %}
+  // Subgroup-local coordinates define the segmented fold; local invocation
+  // ordering is not a portable substitute. The manifest admits only full
+  // subgroups and masks any extra value tiles after rounding a short group up.
+  let lane = sg_lane % LANES;
+{% if variableSubgroupFold %}
+  let v_group = sg_id * (sg_size / LANES) + sg_lane / LANES;
+{% else %}
+  let v_group = sg_id * {{ (subgroupWidth / lanes) | int }}u + sg_lane / LANES;
+{% endif %}
+{% else %}
+  let lane = lid.x % LANES;
+  let v_group = lid.x / LANES;
+{% endif %}
+  // 2D-folded (batch * value head * value-tile block) index: wg.y carries the high
+  // bits past the per-axis dispatch fold width.
+  let unit = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
+  let block = unit % DV_BLOCKS;
+  let bh = unit / DV_BLOCKS;
+  let head_v = bh % NUM_V_HEADS;
+  let batch = bh / NUM_V_HEADS;
+  if (batch >= params.batchSize) {
+    return;
+  }
+  let head_q = head_v / HEADS_PER_Q;
+  let dv_start = (block * V_GROUPS + v_group) * TILE_V;
+{% if hasCu %}
+  // Device-supplied offsets are clamped, not validated: a negative, decreasing or
+  // past-the-end entry must yield an empty or truncated request rather than steer
+  // an out-of-bounds access.
+  let token_limit = i32(params.totalTokens);
+  let token_begin = u32(clamp(cu_seqlens[batch], 0, token_limit));
+  let token_end = u32(clamp(cu_seqlens[batch + 1u], 0, token_limit));
+{% else %}
+  let token_begin = batch * params.seqLength;
+  let token_end = token_begin + params.seqLength;
+{% endif %}
+  let state_base = (batch * NUM_V_HEADS + head_v) * HEAD_DIM_V * DK_WORDS;
+{% for j in range(tileV) %}
+{% for w in range(wordsPerLane) %}
+  var s{{ j }}_{{ w }} = {{ wordF32 }}(0.0);
+{% endfor %}
+{% endfor %}
+{% for w in range(wordsPerLane) %}
+  let r{{ w }} = lane + {{ w }}u * LANES;
+{% if not exactWords %}
+  let word_ok{{ w }} = r{{ w }} < DK_WORDS;
+{% endif %}
+{% endfor %}
+{% if hasState %}
+{% for j in range(tileV) %}
+  if (dv_start + {{ j }}u < HEAD_DIM_V) {
+    let row{{ j }} = state_base + (dv_start + {{ j }}u) * DK_WORDS;
+{% for w in range(wordsPerLane) %}
+{% if exactWords %}
+    s{{ j }}_{{ w }} = initial_state[row{{ j }} + r{{ w }}];
+{% else %}
+    if (word_ok{{ w }}) { s{{ j }}_{{ w }} = initial_state[row{{ j }} + r{{ w }}]; }
+{% endif %}
+{% endfor %}
+  }
+{% endfor %}
+{% endif %}
+{% if pairTokens %}
+{{ paired_recurrence() }}
+  if (token_end > token_begin && (token_end - token_begin) % 2u != 0u) {
+    let t = token_end - 1u;
+{{ recurrence_step(foldSlotsDelta, 5) }}
+  }
+{% else %}
+  for (var t = token_begin; t < token_end; t = t + 1u) {
+{{ recurrence_step(foldSlots, 3) }}
+  }
+{% endif %}
+{% for j in range(tileV) %}
+  if (dv_start + {{ j }}u < HEAD_DIM_V) {
+    let out_state{{ j }} = state_base + (dv_start + {{ j }}u) * DK_WORDS;
+{% for w in range(wordsPerLane) %}
+{% if exactWords %}
+    final_state[out_state{{ j }} + r{{ w }}] = s{{ j }}_{{ w }};
+{% else %}
+    if (word_ok{{ w }}) { final_state[out_state{{ j }} + r{{ w }}] = s{{ j }}_{{ w }}; }
+{% endif %}
+{% endfor %}
+  }
+{% endfor %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1093 @@

+{
+  "domain": "com.microsoft",
+  "name": "GatedDeltaNet",
+  "sinceVersion": 1,
+  "inputs": {
+    "queryT": { "onnx": "query", "dtype": "T" },
+    "keyT": { "onnx": "key", "dtype": "T" },
+    "valueT": { "onnx": "value", "dtype": "T" },
+    "cuSeqlensT": { "onnx": "cu_seqlens", "dtype": "TI", "rank": 1, "optional": true },
+    "decayT": { "onnx": "decay", "dtype": "TS", "optional": true },
+    "betaT": { "onnx": "beta", "dtype": "TS", "optional": true },
+    "initialStateT": { "onnx": "initial_state", "dtype": "TS", "rank": 4, "optional": true },
+    "aLogT": { "onnx": "a_log", "dtype": "TS", "rank": 1, "optional": true },
+    "dtBiasT": { "onnx": "dt_bias", "dtype": "TS", "rank": 1, "optional": true }
+  },
+  "outputs": {
+    "outputT": {
+      "onnx": "output",
+      "dtype": "T",
+      "rank": "ranks.queryT",
+      "shape": "[dim(shapes.queryT, 0), dim(shapes.queryT, 1), numVHeads, headDimV] if rank4 else [dim(shapes.queryT, 0), numVHeads, headDimV]"
+    },
+    "finalStateT": {
+      "onnx": "final_state",
+      "dtype": "TS",
+      "rank": 4,
+      "shape": "[batchCount, numVHeads, headDimV, headDimK]"
+    }
+  },
+  "attributes": {
+    "update_rule": { "default": "gated_delta" },
+    "scale": { "default": 0 },
+    "gate_activation": { "default": "none" },
+    "beta_activation": { "default": "none" },
+    "qk_l2_norm": { "default": 0 },
+    "chunk_size": { "default": 64 },
+    "state_update_capacity": { "default": 0 }
+  },
+  "attributeConstraints": {
+    "update_rule": { "values": ["linear", "gated", "delta", "gated_delta"] },
+    "gate_activation": { "values": ["none", "qwen"] },
+    "beta_activation": { "values": ["none", "sigmoid"] },
+    "qk_l2_norm": { "values": [0, 1] },
+    "state_update_capacity": { "values": [0] }
+  },
+  "typeConstraints": { "T": ["float32", "float16"], "TS": ["float32"], "TI": ["int32"] },
+  "tunables": {
+    "TILE_V": { "default": 0 },
+    "LANE_WORDS": { "default": 4 },
+    "V_GROUPS": { "default": 8 },
+    "MIN_VECTOR_VALUE_GROUPS": { "default": 4 }
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "foldedDispatchCapacity": "min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "rank4": "ranks.queryT == 4",
+    "numQHeads": "dim(shapes.queryT, ranks.queryT - 2)",
+    "headDimK": "dim(shapes.queryT, ranks.queryT - 1)",
+    "numVHeads": "dim(shapes.valueT, ranks.valueT - 2)",
+    "headDimV": "dim(shapes.valueT, ranks.valueT - 1)",
+    "totalTokens": "numel(shapes.queryT) / max(1, numQHeads * headDimK)",
+    "batchCount": "(dim(shapes.cuSeqlensT, 0) - 1) if present.cuSeqlensT else (dim(shapes.queryT, 0) if rank4 else (dim(shapes.initialStateT, 0) if present.initialStateT else 0))",
+    "seqLength": "totalTokens / max(1, batchCount)",
+    "headsPerQ": "numVHeads / max(1, numQHeads)",
+    "rankOk": "(ranks.queryT == 3 or ranks.queryT == 4) and ranks.keyT == ranks.queryT and ranks.valueT == ranks.queryT and ranks.outputT == ranks.queryT",
+    "headsOk": "numQHeads > 0 and numVHeads > 0 and headDimK > 0 and headDimV > 0 and numVHeads % max(1, numQHeads) == 0",
+    "tokenAxesOk": "sameShape(shapes.keyT, shapes.queryT) and sameShape(prefix(shapes.valueT, ranks.queryT - 2), prefix(shapes.queryT, ranks.queryT - 2))",
+    "batchOk": "batchCount > 0 and (true if present.cuSeqlensT else totalTokens % max(1, batchCount) == 0)",
+    "outputOk": "sameShape(prefix(shapes.outputT, ranks.queryT - 2), prefix(shapes.queryT, ranks.queryT - 2)) and dim(shapes.outputT, ranks.outputT - 2) == numVHeads and dim(shapes.outputT, ranks.outputT - 1) == headDimV",
+    "stateOk": "ranks.finalStateT == 4 and dim(shapes.finalStateT, 0) == batchCount and dim(shapes.finalStateT, 1) == numVHeads and dim(shapes.finalStateT, 2) == headDimV and dim(shapes.finalStateT, 3) == headDimK and (sameShape(shapes.initialStateT, shapes.finalStateT) if present.initialStateT else true)",
+    "gateShapeOk": "(ranks.decayT == ranks.queryT - 1 and sameShape(prefix(shapes.decayT, ranks.queryT - 2), prefix(shapes.queryT, ranks.queryT - 2)) and dim(shapes.decayT, ranks.decayT - 1) == numVHeads if present.decayT else true) and (ranks.betaT == ranks.queryT - 1 and sameShape(prefix(shapes.betaT, ranks.queryT - 2), prefix(shapes.queryT, ranks.queryT - 2)) and dim(shapes.betaT, ranks.betaT - 1) == numVHeads if present.betaT else true) and (ranks.aLogT == 1 and dim(shapes.aLogT, 0) == numVHeads if present.aLogT else true) and (ranks.dtBiasT == 1 and dim(shapes.dtBiasT, 0) == numVHeads if present.dtBiasT else true)",
+    "cuOk": "ranks.cuSeqlensT == 1 and ranks.queryT == 3 and dim(shapes.cuSeqlensT, 0) == batchCount + 1 if present.cuSeqlensT else true",
+    "dtypeOk": "tensorDtypes.keyT == tensorDtypes.queryT and tensorDtypes.valueT == tensorDtypes.queryT and tensorDtypes.outputT == tensorDtypes.queryT and f16Ok(dtypes.T)",
+    "contractOk": "rankOk and headsOk and tokenAxesOk and batchOk and outputOk and stateOk and gateShapeOk and cuOk and dtypeOk",
+    "useVec4": "headDimK % 4 == 0",
+    "dkWords": "headDimK / 4 if useVec4 else headDimK",
+    "laneCap": "pow(2, log2ceil(deviceWorkgroupCap + 1) - 1)",
+    "baseLanes": "max(1, min(laneCap, pow2ceil(ceilDiv(dkWords, max(1, tunables.LANE_WORDS)))))",
+    "subgroupVariable": "device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize < device.adapterInfo.subgroupMaxSize",
+    "wideTileV": "max(1, min(4, headDimV, floor(16 / max(1, ceilDiv(dkWords, baseLanes)))))",
+    "wideVGroups": "max(1, min(tunables.V_GROUPS, ceilDiv(headDimV, wideTileV), max(1, floor(deviceWorkgroupCap / baseLanes))))",
+    "wideWorkgroupSize": "baseLanes * wideVGroups",
+    "variableSubgroupTile": "subgroupVariable and tensorDtypes.queryT == \"float16\" and baseLanes > 1 and baseLanes <= device.adapterInfo.subgroupMinSize and device.adapterInfo.subgroupMinSize % baseLanes == 0 and device.adapterInfo.subgroupMaxSize % baseLanes == 0 and wideWorkgroupSize % device.adapterInfo.subgroupMinSize == 0 and wideWorkgroupSize % device.adapterInfo.subgroupMaxSize == 0 and seqLength >= baseLanes and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "tileV": "max(1, min(tunables.TILE_V if tunables.TILE_V > 0 else (wideTileV if variableSubgroupTile else 2), headDimV))",
+    "subgroupFixed": "device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == device.adapterInfo.subgroupMaxSize",
+    "subgroupWidth": "device.adapterInfo.subgroupMinSize if subgroupFixed else 0",
+    "subgroupLaneFit": "not useVec4 and subgroupFixed and subgroupWidth > 1 and subgroupWidth < baseLanes and subgroupWidth <= deviceWorkgroupCap and device.wgslLanguageFeatures.has(\"subgroup_id\") and tileV * ceilDiv(dkWords, subgroupWidth) <= 64 and ceilDiv(dkWords, subgroupWidth) <= 4 * max(1, tunables.LANE_WORDS)",
+    "lanes": "subgroupWidth if subgroupLaneFit else baseLanes",
+    "wordsPerLane": "ceilDiv(dkWords, max(1, lanes))",
+    "dvTiles": "ceilDiv(headDimV, max(1, tileV))",
+    "baseVGroups": "max(1, min(tunables.V_GROUPS, dvTiles, max(1, floor(deviceWorkgroupCap / max(1, lanes)))))",
+    "vGroups": "min(ceilDiv(baseVGroups, subgroupWidth / lanes), floor(deviceWorkgroupCap / subgroupWidth)) * (subgroupWidth / lanes) if subgroupFixed and lanes > 1 and lanes <= subgroupWidth and subgroupWidth <= deviceWorkgroupCap and device.wgslLanguageFeatures.has(\"subgroup_id\") else baseVGroups",
+    "workgroupSize": "lanes * vGroups",
+    "variableSubgroupFold": "subgroupVariable and lanes > 1 and lanes <= device.adapterInfo.subgroupMinSize and device.adapterInfo.subgroupMinSize % lanes == 0 and device.adapterInfo.subgroupMaxSize % lanes == 0 and workgroupSize % device.adapterInfo.subgroupMinSize == 0 and workgroupSize % device.adapterInfo.subgroupMaxSize == 0 and seqLength >= lanes and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "dvBlocks": "ceilDiv(dvTiles, max(1, vGroups))",
+    "dispatchUnits": "batchCount * numVHeads * dvBlocks",
+    "useL2Norm": "attrs.qk_l2_norm == 1",
+    "useSubgroups": "(lanes > 1 and subgroupFixed and lanes <= subgroupWidth and workgroupSize % max(1, subgroupWidth) == 0 and device.wgslLanguageFeatures.has(\"subgroup_id\")) or variableSubgroupFold",
+    "foldSlotsPlain": "ceilDiv(tileV, 4)",
+    "foldSlotsDelta": "ceilDiv(2 * tileV, 4)",
+    "foldFitsPlain": "useSubgroups or lanes == 1 or workgroupSize * foldSlotsPlain * 16 <= device.limits.maxComputeWorkgroupStorageSize",
+    "foldFitsDelta": "useSubgroups or lanes == 1 or workgroupSize * foldSlotsDelta * 16 <= device.limits.maxComputeWorkgroupStorageSize",
+    "registerWordsOk": "tileV * wordsPerLane <= (64 if subgroupLaneFit else 16)",
+    "geometryOk": "registerWordsOk and workgroupSize <= deviceWorkgroupCap and dispatchUnits > 0 and dispatchUnits <= foldedDispatchCapacity",
+    "outScale": "attrs.scale if attrs.scale != 0.0 else 1.0 / pow(headDimK, 0.5)",
+    "qkElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if useVec4 else dtypes.T",
+    "wordF32": "\"vec4<f32>\" if useVec4 else \"f32\"",
+    "gateItems": "totalTokens * numVHeads",
+    "qkUnits": "totalTokens * numQHeads",
+    "gateLanes": "max(1, min(laneCap, 64, pow2ceil(dkWords)))",
+    "qwenGate": "attrs.gate_activation == \"qwen\"",
+    "betaSigmoid": "attrs.beta_activation == \"sigmoid\"",
+    "valueIOWidth": "device.adapterInfo.subgroupMinSize if has(device.adapterInfo, \"subgroupMinSize\") else 0",
+    "vectorValueIO": "(tileV == 2 or tileV == 4) and headDimV % tileV == 0 and (tensorDtypes.valueT == \"float16\" or lanes * tunables.MIN_VECTOR_VALUE_GROUPS <= valueIOWidth)",
+    "valueIOElement": "(\"vec\" ~ tileV ~ \"<\" ~ dtypes.T ~ \">\") if vectorValueIO else dtypes.T"
+  },
+  "when": ["contractOk", "geometryOk"],
+  "bindings": {
+    "query": { "arg": "queryT", "name": "query", "elementType": "$qkElement" },
+    "key": { "arg": "keyT", "name": "key", "elementType": "$qkElement" },
+    "value": { "arg": "valueT", "name": "value", "elementType": "$valueIOElement" },
+    "cu_seqlens": { "arg": "cuSeqlensT", "name": "cu_seqlens" },
+    "gate": { "scratch": "gate", "name": "gate", "buffer": "read-only-storage", "elementType": "f32" },
+    "initial_state": { "arg": "initialStateT", "name": "initial_state", "elementType": "$wordF32" },
+    "output": { "arg": "outputT", "name": "output", "elementType": "$valueIOElement" },
+    "final_state": { "arg": "finalStateT", "name": "final_state", "elementType": "$wordF32" },
+    "params": {
+      "struct": [
+        { "name": "batchSize", "type": "u32", "value": "batchCount" },
+        { "name": "seqLength", "type": "u32", "value": "seqLength" },
+        { "name": "qkBase", "type": "u32", "value": "qkBase" }
+      ]
+    },
+    "params_ragged": {
+      "name": "params",
+      "struct": [
+        { "name": "batchSize", "type": "u32", "value": "batchCount" },
+        { "name": "totalTokens", "type": "u32", "value": "totalTokens" },
+        { "name": "qkBase", "type": "u32", "value": "qkBase" }
+      ]
+    },
+    "decay": { "arg": "decayT", "name": "decay" },
+    "beta": { "arg": "betaT", "name": "beta" },
+    "a_log": { "arg": "aLogT", "name": "a_log" },
+    "dt_bias": { "arg": "dtBiasT", "name": "dt_bias" },
+    "gate_out": { "scratch": "gate", "name": "gate", "elementType": "f32" },
+    "params_gate": {
+      "name": "params",
+      "struct": [
+        { "name": "units", "type": "u32", "value": "qkUnits" },
+        { "name": "qkBase", "type": "u32", "value": "qkBase" }
+      ]
+    }
+  },
+  "variants": [
+    {
+      "id": "linear_state_ragged",
+      "when": ["attrs.update_rule == \"linear\"", "not present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "initial_state", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "linear_state_packed",
+      "when": ["attrs.update_rule == \"linear\"", "not present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "initial_state", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "linear_zero_ragged",
+      "when": ["attrs.update_rule == \"linear\"", "not present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "linear_zero_packed",
+      "when": ["attrs.update_rule == \"linear\"", "not present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_state_ragged",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "initial_state", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_state_packed",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "initial_state", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_zero_ragged",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_zero_packed",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_qwen_state_ragged",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "initial_state", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_qwen_state_packed",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "initial_state", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_qwen_zero_ragged",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "not present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_qwen_zero_packed",
+      "when": ["attrs.update_rule == \"gated\"", "present.decayT", "not present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "not present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "delta_state_ragged",
+      "when": ["attrs.update_rule == \"delta\"", "not present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "initial_state", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "delta_state_packed",
+      "when": ["attrs.update_rule == \"delta\"", "not present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "initial_state", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "delta_zero_ragged",
+      "when": ["attrs.update_rule == \"delta\"", "not present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "delta_zero_packed",
+      "when": ["attrs.update_rule == \"delta\"", "not present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_state_ragged",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "initial_state", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_state_packed",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "initial_state", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_zero_ragged",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_zero_packed",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "not present.aLogT", "not present.dtBiasT", "not qwenGate", "not present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_qwen_state_ragged",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "initial_state", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_qwen_state_packed",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "initial_state", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_qwen_zero_ragged",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "not present.initialStateT", "present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "cu_seqlens", "gate", "output", "final_state", "params_ragged"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gated_delta_qwen_zero_packed",
+      "when": ["attrs.update_rule == \"gated_delta\"", "present.decayT", "present.betaT", "present.aLogT", "present.dtBiasT", "qwenGate", "not present.initialStateT", "not present.cuSeqlensT", "foldFitsDelta if present.betaT else foldFitsPlain"],
+      "intermediates": [{ "id": "gate", "dtype": "float32", "shape": "[qkBase + (5 if pairTokens else 3) * qkUnits]" }],
+      "derive": {
+        "hasGate": "present.decayT or present.betaT",
+        "qkBase": "2 * gateItems if hasGate else 0",
+        "useDelta": "present.betaT",
+        "hasState": "present.initialStateT",
+        "hasCu": "present.cuSeqlensT",
+        "pairTokens": "useDelta and not useSubgroups and tensorDtypes.queryT == \"float16\" and lanes > 1 and seqLength >= 2 * lanes and workgroupSize * tileV * 16 <= device.limits.maxComputeWorkgroupStorageSize and 4 * (qkBase + 5 * qkUnits) <= min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+        "foldSlots": "tileV if pairTokens else (foldSlotsDelta if useDelta else foldSlotsPlain)"
+      },
+      "passes": [
+        {
+          "id": "gate",
+          "name": "GatedDeltaNet.Gate",
+          "shader": "gated-delta-net-gate.wgsl.jinja",
+          "derive": {
+            "hasDecayInput": "present.decayT",
+            "hasBetaInput": "present.betaT",
+            "qwen": "qwenGate",
+            "sigmoidBeta": "betaSigmoid"
+          },
+          "bindings": ["query", "key", "decay", "beta", "a_log", "dt_bias", "gate_out", "params_gate"],
+          "dispatch": { "x": "min(qkUnits, 65535)", "y": "ceilDiv(qkUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "main",
+          "name": "GatedDeltaNet.Recurrent",
+          "shader": "gated-delta-net.wgsl.jinja",
+          "bindings": ["query", "key", "value", "gate", "output", "final_state", "params"],
+          "dispatch": {
+            "x": "min(dispatchUnits, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(dispatchUnits, DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,47 @@

+{
+  "name": "com.microsoft.GatedDeltaNet",
+  "id": "_com_microsoft_gateddeltanet_webgpu_327a002",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "JjlCHBCgwhVtN5uX68wTtRvl28U3uxij+8CtQr6ihus=",
+      "gated-delta-net-gate.wgsl.jinja": "F2VfVvKYHtF5V+WJYEFCyTIoLtqOZScZjglN7Ld/qZk=",
+      "gated-delta-net.wgsl.jinja": "afS/AoT0XbtYWu3HVPbuzBMEij1yt+oz71UGqaUlplU=",
+      "manifest.json": "TQScCwpV1Wu5pvm5/XBsrQdcq9np2crCD+QN/pbaUC8=",
+      "test.json": "I7KVc4VWBUm1VG7XYAc/1ntOGoAmvG6eWc2/7pW9VVU="
+    }
+  },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.1",
+    "variants": {
+      "linear_state_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "linear_state_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "linear_zero_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "linear_zero_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_state_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_state_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_zero_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_zero_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_qwen_state_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_qwen_state_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_qwen_zero_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_qwen_zero_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "delta_state_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "delta_state_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "delta_zero_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "delta_zero_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_state_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_state_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_zero_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_zero_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_qwen_state_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_qwen_state_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_qwen_zero_ragged": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"],
+      "gated_delta_qwen_zero_packed": ["gated-delta-net-gate.wgsl.jinja", "gated-delta-net.wgsl.jinja"]
+    }
+  }
+}

build/webgpu/test.json ADDED Viewed

The diff for this file is too large to render. See raw diff