Hugging Face's logo

Kernels:
webgpu-kernels
/
com.microsoft.GatedAdd

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
5e78980
·
verified ·
1 Parent(s): d27ea5a

sync 91d990483a17

Browse files
Files changed (6) hide show
  1. README.md +15 -11
  2. build/webgpu/bench.json +0 -1
  3. build/webgpu/gated-add.wgsl.jinja +24 -35
  4. build/webgpu/manifest.json +31 -110
  5. build/webgpu/metadata.json +10 -7
  6. build/webgpu/test.json +1 -2
README.md CHANGED
@@ -18,17 +18,17 @@ See the [ONNX Runtime `GatedAdd` contrib-operator spec](https://github.com/micro
18
 
19
  ## Inputs
20
 
21
- | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
- | --- | --- | --- | --- | --- | --- | --- |
23
- | `X` | `X` | `T` | — | — | Unscaled input with shape `(..., C)`. Any rank of at least 1 is accepted; only the trailing channel axis is distinguished. | required |
24
- | `Y` | `Y` | `T` | — | — | Input scaled by the gate, with the same shape as `X`. | required |
25
- | `gate` | `gate` | `T` | — | — | Per-row gate with shape `(..., 1)`: the same rank and leading dimensions as `X`, with a trailing dimension of 1 that broadcasts over the `C` channels. | required |
26
 
27
  ## Outputs
28
 
29
- | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
30
- | --- | --- | --- | --- | --- | --- | --- |
31
- | `output` | `output` | `T` | same as `X` | same as `X` | Gated sum `X + round_to_T(Y * gate)`, with the same shape as `X`. | required |
32
 
33
  ## Type constraints
34
 
@@ -38,7 +38,7 @@ See the [ONNX Runtime `GatedAdd` contrib-operator spec](https://github.com/micro
38
 
39
  ## Files
40
 
41
- - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
42
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
43
  - [`test.json`](build/webgpu/test.json) — correctness cases
44
  - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
@@ -46,10 +46,14 @@ See the [ONNX Runtime `GatedAdd` contrib-operator spec](https://github.com/micro
46
 
47
  ## Use with `@huggingface/kernels`
48
 
49
- The loader derives every required output's shape and logical dtype from the manifest contract and this call.
50
- It then allocates the result tensors automatically.
 
 
 
51
 
52
  The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 
53
 
54
  Replace each `*Data` placeholder with a typed array containing the corresponding input data.
55
 
 
18
 
19
  ## Inputs
20
 
21
+ | Name | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- |
23
+ | `X` | `T` | — | — | Unscaled input with shape `(..., C)`. Any rank of at least 1 is accepted; only the trailing channel axis is distinguished. | required |
24
+ | `Y` | `T` | — | — | Input scaled by the gate, with the same shape as `X`. | required |
25
+ | `gate` | `T` | — | — | Per-row gate with shape `(..., 1)`: the same rank and leading dimensions as `X`, with a trailing dimension of 1 that broadcasts over the `C` channels. | required |
26
 
27
  ## Outputs
28
 
29
+ | Name | Logical dtype | Rank | Shape | Description | Presence |
30
+ | --- | --- | --- | --- | --- | --- |
31
+ | `output` | `T` | same as `X` | same as `X` | Gated sum `X + round_to_T(Y * gate)`, with the same shape as `X`. | required |
32
 
33
  ## Type constraints
34
 
 
38
 
39
  ## Files
40
 
41
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
42
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
43
  - [`test.json`](build/webgpu/test.json) — correctness cases
44
  - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 
46
 
47
  ## Use with `@huggingface/kernels`
48
 
49
+ ```sh
50
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.2
51
+ ```
52
+
53
+ Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
54
 
55
  The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
56
+ It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
57
 
58
  Replace each `*Data` placeholder with a typed array containing the corresponding input data.
59
 
build/webgpu/bench.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "com.microsoft.GatedAdd",
3
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
  "cases": [
5
  {
 
1
  {
 
2
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
3
  "cases": [
4
  {
build/webgpu/gated-add.wgsl.jinja CHANGED
@@ -1,64 +1,53 @@
1
  {% macro flat_tail_open() %}
2
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
3
- fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
4
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
5
- // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
6
- let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
7
- {% if source.itemsPerInvocation is defined %}
8
  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
9
  // the logical tensor length (or its storage binding) to be vec4 aligned.
10
- let begin = invocation * {{ source.itemsPerInvocation }}u;
11
- let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
 
12
  for (var i = begin; i < end; i = i + 1u) {
13
- {%- else %}
14
- let i = invocation;
15
- if (i >= params.count) {
16
- return;
17
- }
18
- {%- endif %}
19
- {% endmacro %}
20
  {% macro flat_tail_close() %}
21
- {% if source.itemsPerInvocation is defined %}
22
  }
23
- {% endif %}
24
  {% endmacro %}
25
 
26
  {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
27
  {% if note == "dispatch-limit" %}
28
- // 2D-folded flat index: gid.y carries the high bits past the
29
- // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
30
  {% elif note == "limit" %}
31
- // 2D-folded flat index: gid.y carries the high bits past the
32
- // maxComputeWorkgroupsPerDimension limit.
33
  {% elif note == "device-axis" %}
34
- // The flat dispatch is folded across x/y at the device's per-axis workgroup
35
- // limit; gid.y carries the high portion of the output index.
36
  {% elif note == "vec4-limit" %}
37
- // 2D-folded flat vec4 index: gid.y carries the high bits past the
38
- // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
39
  {% elif note == "element-limit" %}
40
  // 2D-folded flat element index: gid.y carries the high bits past the
41
- // maxComputeWorkgroupsPerDimension limit.
42
  {% elif note == "dispatch" %}
43
- // 2D-folded flat index: gid.y carries the high bits past the
44
- // maxComputeWorkgroupsPerDimension dispatch limit.
45
  {% endif %}
46
  {% if bound == "" %}
47
- let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
48
  {%- elif guardInline %}
49
- let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
50
  if ({{ name }} >= {{ bound }}) { return; }
51
  {%- else %}
52
- let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
53
  if ({{ name }} >= {{ bound }}) {
54
  return;
55
  }
56
  {%- endif %}
57
  {% endmacro %}
58
 
59
- {% if usesF16 %}
60
- enable f16;
61
- {% endif %}
62
  {{ env.wgsl.resourceDeclarations }}
63
 
64
  // com.microsoft.GatedAdd : output = X + round_to_T(Y * gate)
@@ -69,13 +58,13 @@ enable f16;
69
  // separate Mul followed by Add. fma(y, gate, 0.0) is that single rounding
70
  // (adding zero cannot move the product) and, unlike a bare y * gate, it cannot
71
  // be contracted into the following add by a backend that permits floating-point
72
- // contraction -- contraction would keep an unrounded wider product and silently
73
- // make this op more accurate than the graph it replaces.
74
  const HIDDEN: u32 = {{ hidden }}u;
75
 
76
  {% if vec4 %}
77
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
78
- fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
79
  {{ flat_index_2d() }}
80
  // A vec4 group is four consecutive channels of one row: HIDDEN % 4 == 0 stops
81
  // it from ever straddling two rows, so the whole group shares one gate value.
 
1
  {% macro flat_tail_open() %}
2
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
3
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
4
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
5
+ // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
6
+ let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 
7
  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
8
  // the logical tensor length (or its storage binding) to be vec4 aligned.
9
+ {% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
10
+ let begin = invocation * {{ itemsPerInvocation }}u;
11
+ let end = min(begin + {{ itemsPerInvocation }}u, params.count);
12
  for (var i = begin; i < end; i = i + 1u) {
13
+ {%- endmacro %}
 
 
 
 
 
 
14
  {% macro flat_tail_close() %}
 
15
  }
 
16
  {% endmacro %}
17
 
18
  {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
19
  {% if note == "dispatch-limit" %}
20
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
21
+ // per-axis workgroup fold width (outputs > 16.7M elements).
22
  {% elif note == "limit" %}
23
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
24
+ // per-axis workgroup fold width.
25
  {% elif note == "device-axis" %}
26
+ // The flat dispatch is folded across x/y at a fixed per-axis workgroup
27
+ // width; gid.y carries the high portion of the output index.
28
  {% elif note == "vec4-limit" %}
29
+ // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
30
+ // per-axis workgroup fold width (the dispatch caps x and spills into y).
31
  {% elif note == "element-limit" %}
32
  // 2D-folded flat element index: gid.y carries the high bits past the
33
+ // dispatch's per-axis workgroup fold width.
34
  {% elif note == "dispatch" %}
35
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
36
+ // per-axis workgroup fold width.
37
  {% endif %}
38
  {% if bound == "" %}
39
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
40
  {%- elif guardInline %}
41
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
42
  if ({{ name }} >= {{ bound }}) { return; }
43
  {%- else %}
44
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
45
  if ({{ name }} >= {{ bound }}) {
46
  return;
47
  }
48
  {%- endif %}
49
  {% endmacro %}
50
 
 
 
 
51
  {{ env.wgsl.resourceDeclarations }}
52
 
53
  // com.microsoft.GatedAdd : output = X + round_to_T(Y * gate)
 
58
  // separate Mul followed by Add. fma(y, gate, 0.0) is that single rounding
59
  // (adding zero cannot move the product) and, unlike a bare y * gate, it cannot
60
  // be contracted into the following add by a backend that permits floating-point
61
+ // contraction. This preserves the staged precision of separate multiplication
62
+ // and addition across the storage-type boundary.
63
  const HIDDEN: u32 = {{ hidden }}u;
64
 
65
  {% if vec4 %}
66
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
67
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
68
  {{ flat_index_2d() }}
69
  // A vec4 group is four consecutive channels of one row: HIDDEN % 4 == 0 stops
70
  // it from ever straddling two rows, so the whole group shares one gate value.
build/webgpu/manifest.json CHANGED
@@ -2,145 +2,66 @@
2
  "domain": "com.microsoft",
3
  "name": "GatedAdd",
4
  "sinceVersion": 1,
5
- "description": "Adds `Y`, scaled by a per-row `gate`, to `X`: `output = X + round_to_T(Y * gate)`. `X` and `Y` have shape `(..., C)`; `gate` has the same rank with a trailing dimension of 1, so one value covers each row of `C` channels. Rounding the product to `T` before the addition preserves the semantics of a separate `Mul` followed by `Add`. Bfloat16 is not implemented.",
6
- "inputs": [
7
- {
8
- "role": "X",
9
- "dtype": "T",
10
- "description": "Unscaled input with shape `(..., C)`. Any rank of at least 1 is accepted; only the trailing channel axis is distinguished."
11
- },
12
- { "role": "Y", "dtype": "T", "description": "Input scaled by the gate, with the same shape as `X`." },
13
- {
14
- "role": "gate",
15
- "dtype": "T",
16
- "description": "Per-row gate with shape `(..., 1)`: the same rank and leading dimensions as `X`, with a trailing dimension of 1 that broadcasts over the `C` channels."
17
- }
18
- ],
19
- "outputs": [
20
- {
21
- "role": "output",
22
- "dtype": "T",
23
- "rank": "ranks.X",
24
- "shape": "shapes.X",
25
- "description": "Gated sum `X + round_to_T(Y * gate)`, with the same shape as `X`."
26
- }
27
- ],
28
  "typeConstraints": { "T": ["float32", "float16"] },
29
- "tunables": { "WORKGROUP_SIZE": 256 },
30
- "args": {
31
- "X": { "kind": "tensor", "semantic": "X", "role": "input" },
32
- "Y": { "kind": "tensor", "semantic": "Y", "role": "input" },
33
- "gate": { "kind": "tensor", "semantic": "gate", "role": "input" },
34
- "output": { "kind": "tensor", "semantic": "output", "role": "output" }
35
- },
36
  "derive": {
37
  "channels": "dim(shapes.X, ranks.X - 1)",
38
  "gateContract": "ranks.X >= 1 and channels > 0 and ranks.Y == ranks.X and ranks.gate == ranks.X and sameShape(shapes.Y, shapes.X) and sameShape(shapes.output, shapes.X) and dim(shapes.gate, ranks.gate - 1) == 1 and sameShape(prefix(shapes.gate, ranks.gate - 1), prefix(shapes.X, ranks.X - 1)) and f16Ok(dtypes.T)",
39
- "vec4Rows": "channels % 4 == 0 and numel(shapes.X) % 4 == 0"
 
 
40
  },
41
- "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"", "hidden": "channels if channels > 0 else 1" },
42
  "variants": [
43
  {
44
  "id": "vec4",
45
  "priority": 30,
46
- "when": ["gateContract", "vec4Rows", "numel(shapes.X) > 0"],
47
- "constants": { "vec4": true, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
48
  "passes": [
49
  {
50
  "id": "main",
51
  "name": "GatedAdd.vec4",
52
  "shader": "gated-add.wgsl.jinja",
53
  "bindings": [
54
- {
55
- "name": "x",
56
- "arg": "X",
57
- "semantic": "X",
58
- "buffer": { "type": "read-only-storage" },
59
- "elementType": "$vectorScalar"
60
- },
61
- {
62
- "name": "y",
63
- "arg": "Y",
64
- "semantic": "Y",
65
- "buffer": { "type": "read-only-storage" },
66
- "elementType": "$vectorScalar"
67
- },
68
- {
69
- "name": "gate",
70
- "arg": "gate",
71
- "semantic": "gate",
72
- "buffer": { "type": "read-only-storage" },
73
- "elementType": "$scalar"
74
- },
75
- {
76
- "name": "output",
77
- "arg": "output",
78
- "semantic": "output",
79
- "buffer": { "type": "storage" },
80
- "elementType": "$vectorScalar"
81
- },
82
- {
83
- "name": "params",
84
- "semantic": "kernel.params",
85
- "buffer": { "type": "uniform" },
86
- "struct": {
87
- "name": "Params",
88
- "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }]
89
- }
90
- }
91
  ],
92
- "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
 
 
 
 
93
  }
94
  ]
95
  },
96
  {
97
  "id": "scalar",
98
  "priority": 0,
99
- "when": ["gateContract"],
100
- "constants": { "vec4": false },
101
  "passes": [
102
  {
103
  "id": "main",
104
  "name": "GatedAdd.scalar",
105
  "shader": "gated-add.wgsl.jinja",
106
- "source": { "inputs": { "itemsPerInvocation": 4 } },
107
  "bindings": [
108
- {
109
- "name": "x",
110
- "arg": "X",
111
- "semantic": "X",
112
- "buffer": { "type": "read-only-storage" },
113
- "elementType": "$scalar"
114
- },
115
- {
116
- "name": "y",
117
- "arg": "Y",
118
- "semantic": "Y",
119
- "buffer": { "type": "read-only-storage" },
120
- "elementType": "$scalar"
121
- },
122
- {
123
- "name": "gate",
124
- "arg": "gate",
125
- "semantic": "gate",
126
- "buffer": { "type": "read-only-storage" },
127
- "elementType": "$scalar"
128
- },
129
- {
130
- "name": "output",
131
- "arg": "output",
132
- "semantic": "output",
133
- "buffer": { "type": "storage" },
134
- "elementType": "$scalar"
135
- },
136
- {
137
- "name": "params",
138
- "semantic": "kernel.params",
139
- "buffer": { "type": "uniform" },
140
- "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
141
- }
142
  ],
143
- "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
 
 
 
 
144
  }
145
  ]
146
  }
 
2
  "domain": "com.microsoft",
3
  "name": "GatedAdd",
4
  "sinceVersion": 1,
5
+ "inputs": { "X": { "dtype": "T" }, "Y": { "dtype": "T" }, "gate": { "dtype": "T" } },
6
+ "outputs": { "output": { "dtype": "T", "rank": "ranks.X", "shape": "shapes.X" } },
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
7
  "typeConstraints": { "T": ["float32", "float16"] },
8
+ "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
 
 
 
 
 
 
9
  "derive": {
10
  "channels": "dim(shapes.X, ranks.X - 1)",
11
  "gateContract": "ranks.X >= 1 and channels > 0 and ranks.Y == ranks.X and ranks.gate == ranks.X and sameShape(shapes.Y, shapes.X) and sameShape(shapes.output, shapes.X) and dim(shapes.gate, ranks.gate - 1) == 1 and sameShape(prefix(shapes.gate, ranks.gate - 1), prefix(shapes.X, ranks.X - 1)) and f16Ok(dtypes.T)",
12
+ "vec4Rows": "channels % 4 == 0 and numel(shapes.X) % 4 == 0",
13
+ "scalar": "dtypes.T",
14
+ "hidden": "channels if channels > 0 else 1"
15
  },
16
+ "when": ["gateContract"],
17
  "variants": [
18
  {
19
  "id": "vec4",
20
  "priority": 30,
21
+ "when": ["vec4Rows", "numel(shapes.X) > 0"],
22
+ "derive": { "vec4": true, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
23
  "passes": [
24
  {
25
  "id": "main",
26
  "name": "GatedAdd.vec4",
27
  "shader": "gated-add.wgsl.jinja",
28
  "bindings": [
29
+ { "arg": "X", "name": "x", "elementType": "$vectorScalar" },
30
+ { "arg": "Y", "name": "y", "elementType": "$vectorScalar" },
31
+ "gate",
32
+ { "arg": "output", "elementType": "$vectorScalar" },
33
+ { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }] }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
34
  ],
35
+ "dispatch": {
36
+ "x": "min(ceilDiv((numel(shapes.X) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
37
+ "y": "ceilDiv(ceilDiv((numel(shapes.X) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
38
+ "z": 1
39
+ }
40
  }
41
  ]
42
  },
43
  {
44
  "id": "scalar",
45
  "priority": 0,
46
+ "derive": { "vec4": false },
 
47
  "passes": [
48
  {
49
  "id": "main",
50
  "name": "GatedAdd.scalar",
51
  "shader": "gated-add.wgsl.jinja",
52
+ "derive": { "itemsPerInvocation": 4 },
53
  "bindings": [
54
+ { "arg": "X", "name": "x", "elementType": "$scalar" },
55
+ { "arg": "Y", "name": "y", "elementType": "$scalar" },
56
+ "gate",
57
+ "output",
58
+ { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
59
  ],
60
+ "dispatch": {
61
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.X), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
62
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.X), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
63
+ "z": 1
64
+ }
65
  }
66
  ]
67
  }
build/webgpu/metadata.json CHANGED
@@ -1,18 +1,21 @@
1
  {
2
  "name": "com.microsoft.GatedAdd",
3
- "id": "_com_microsoft_gatedadd_webgpu_bd62a87",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "Yo5szY3Ccb8NOxk0YLKXGApxHMFCENocZt9z2/GaJew=",
11
- "gated-add.wgsl.jinja": "1jbbfo5jhq8VJQHFSPM20gaFGAman4LWBgrssi9fQ5E=",
12
- "manifest.json": "t9s9nEQJMkL8JEPMCsIRRZhZlohyMFsAkyIW7GIwTRw=",
13
- "test.json": "5Gsy9SVd5u5t8bVgKOLTcNMfO4wqI7rFSjHIhPpLEgk="
14
  }
15
  },
16
- "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
17
- "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.GatedAdd" }
 
 
 
18
  }
 
1
  {
2
  "name": "com.microsoft.GatedAdd",
3
+ "id": "_com_microsoft_gatedadd_webgpu_913bb83",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "f/TmVS5V2Nom2d9yiZ8HU0gLA+HtwVePXaGO4EmsmoY=",
11
+ "gated-add.wgsl.jinja": "6xWRekBrMasrrGGXBN6cHj1NeycAppVTPGz1QPRDnUQ=",
12
+ "manifest.json": "a3y5tgF4lud3fMIhIIyomyBnUdsTqCeCaWuGJOSdW1U=",
13
+ "test.json": "CgQChPZDm5zqD3shwXcy5E57YHqJlw8wSwSyTQeEvN4="
14
  }
15
  },
16
+ "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
17
+ "webgpu": {
18
+ "manifestSpec": "2.0",
19
+ "variants": { "vec4": ["gated-add.wgsl.jinja"], "scalar": ["gated-add.wgsl.jinja"] }
20
+ }
21
  }
build/webgpu/test.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "com.microsoft.GatedAdd",
3
  "cases": [
4
  {
5
  "name": "rank3_rows_f32",
@@ -97,7 +96,7 @@
97
  {
98
  "name": "gate_broadcast_rows_pinned",
99
  "provenance": {
100
- "notes": "Hand-computed from the schema formula output = X + round_to_T(Y * gate); every value is exact in float32, so the expectation is independent of the reference."
101
  },
102
  "inputs": {
103
  "X": {
 
1
  {
 
2
  "cases": [
3
  {
4
  "name": "rank3_rows_f32",
 
96
  {
97
  "name": "gate_broadcast_rows_pinned",
98
  "provenance": {
99
+ "notes": "Expected values follow `output = X + round_to_T(Y * gate)` directly, and every intermediate is exactly representable in float32."
100
  },
101
  "inputs": {
102
  "X": {