Hugging Face's logo

Kernels:
webgpu-kernels
/
com.microsoft.GatedAdd

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
b5d19c4
·
verified ·
1 Parent(s): fb9dcfe

sync 2e7068faf55e

Browse files
Files changed (6) hide show
  1. README.md +62 -0
  2. build/webgpu/bench.json +42 -0
  3. build/webgpu/gated-add.wgsl.jinja +63 -0
  4. build/webgpu/manifest.json +147 -0
  5. build/webgpu/metadata.json +18 -0
  6. build/webgpu/test.json +185 -0
README.md CHANGED
@@ -1,3 +1,65 @@
1
  ---
 
2
  license: apache-2.0
 
 
 
 
3
  ---
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
  ---
2
+ library_name: kernels
3
  license: apache-2.0
4
+ tags:
5
+ - kernel
6
+ - webgpu
7
+ - wgsl
8
  ---
9
+ # com.microsoft.GatedAdd
10
+
11
+ `com.microsoft` · ONNX Runtime contrib operator · contrib since_version 1
12
+
13
+ ## Description
14
+
15
+ 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.
16
+
17
+ See the [ONNX Runtime `GatedAdd` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.GatedAdd) for the reference semantics.
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
+
35
+ | Variable | Allowed dtypes |
36
+ | --- | --- |
37
+ | `T` | `float32`, `float16` |
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
45
+ - [`gated-add.wgsl.jinja`](build/webgpu/gated-add.wgsl.jinja)
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
+
56
+ ```js
57
+ import { getKernel } from "@huggingface/kernels";
58
+
59
+ const kernel = await getKernel("webgpu-kernels/com.microsoft.GatedAdd", { version: 1 });
60
+ const { output } = await kernel({
61
+ X: { data: XData, shape: [2, 3] },
62
+ Y: { data: YData, shape: [2, 3] },
63
+ gate: { data: gateData, shape: [2, 1] },
64
+ });
65
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,42 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.GatedAdd",
3
+ "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
+ "cases": [
5
+ {
6
+ "name": "gatedadd-f32-t2048-h2048",
7
+ "preset": "smoke",
8
+ "vars": { "dtype": "float32" },
9
+ "inputs": {
10
+ "X": { "shape": [2048, 2048], "dtype": "float32", "dist": "normal", "seed": 7101, "scale": 2 },
11
+ "Y": { "shape": [2048, 2048], "dtype": "float32", "dist": "normal", "seed": 7102, "scale": 2 },
12
+ "gate": { "shape": [2048, 1], "dtype": "float32", "dist": "normal", "seed": 7103, "scale": 1 }
13
+ },
14
+ "outputs": { "output": { "shape": [2048, 2048], "dtype": "float32" } },
15
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "2048 * 2048 * 3 * dtypeBytes(args.dtype)" }] }
16
+ },
17
+ {
18
+ "name": "gatedadd-f16-t2048-h2048",
19
+ "preset": "model",
20
+ "vars": { "dtype": "float16" },
21
+ "inputs": {
22
+ "X": { "shape": [2048, 2048], "dtype": "float16", "dist": "normal", "seed": 7104, "scale": 2 },
23
+ "Y": { "shape": [2048, 2048], "dtype": "float16", "dist": "normal", "seed": 7105, "scale": 2 },
24
+ "gate": { "shape": [2048, 1], "dtype": "float16", "dist": "normal", "seed": 7106, "scale": 1 }
25
+ },
26
+ "outputs": { "output": { "shape": [2048, 2048], "dtype": "float16" } },
27
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "2048 * 2048 * 3 * dtypeBytes(args.dtype)" }] }
28
+ },
29
+ {
30
+ "name": "gatedadd-f32-channels-odd-h1023",
31
+ "preset": "stress",
32
+ "vars": { "dtype": "float32" },
33
+ "inputs": {
34
+ "X": { "shape": [4096, 1023], "dtype": "float32", "dist": "normal", "seed": 7107, "scale": 2 },
35
+ "Y": { "shape": [4096, 1023], "dtype": "float32", "dist": "normal", "seed": 7108, "scale": 2 },
36
+ "gate": { "shape": [4096, 1], "dtype": "float32", "dist": "normal", "seed": 7109, "scale": 1 }
37
+ },
38
+ "outputs": { "output": { "shape": [4096, 1023], "dtype": "float32" } },
39
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "4096 * 1023 * 3 * dtypeBytes(args.dtype)" }] }
40
+ }
41
+ ]
42
+ }
build/webgpu/gated-add.wgsl.jinja ADDED
@@ -0,0 +1,63 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
2
+ {% if note == "dispatch-limit" %}
3
+ // 2D-folded flat index: gid.y carries the high bits past the
4
+ // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
5
+ {% elif note == "limit" %}
6
+ // 2D-folded flat index: gid.y carries the high bits past the
7
+ // maxComputeWorkgroupsPerDimension limit.
8
+ {% elif note == "device-axis" %}
9
+ // The flat dispatch is folded across x/y at the device's per-axis workgroup
10
+ // limit; gid.y carries the high portion of the output index.
11
+ {% elif note == "vec4-limit" %}
12
+ // 2D-folded flat vec4 index: gid.y carries the high bits past the
13
+ // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
14
+ {% elif note == "element-limit" %}
15
+ // 2D-folded flat element index: gid.y carries the high bits past the
16
+ // maxComputeWorkgroupsPerDimension limit.
17
+ {% elif note == "dispatch" %}
18
+ // 2D-folded flat index: gid.y carries the high bits past the
19
+ // maxComputeWorkgroupsPerDimension dispatch limit.
20
+ {% endif %}
21
+ {% if bound == "" %}
22
+ let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
23
+ {%- elif guardInline %}
24
+ let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
25
+ if ({{ name }} >= {{ bound }}) { return; }
26
+ {%- else %}
27
+ let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
28
+ if ({{ name }} >= {{ bound }}) {
29
+ return;
30
+ }
31
+ {%- endif %}
32
+ {% endmacro %}
33
+
34
+ {% if usesF16 %}
35
+ enable f16;
36
+ {% endif %}
37
+ {{ env.wgsl.resourceDeclarations }}
38
+
39
+ // com.microsoft.GatedAdd : output = X + round_to_T(Y * gate)
40
+ // X, Y, output : the same shape (..., C).
41
+ // gate : (..., 1) -- one value per row of C channels, so a row index
42
+ // of element_index / HIDDEN selects it.
43
+ // The product is rounded to T before the add, so the fusion agrees with a
44
+ // separate Mul followed by Add. fma(y, gate, 0.0) is that single rounding
45
+ // (adding zero cannot move the product) and, unlike a bare y * gate, it cannot
46
+ // be contracted into the following add by a backend that permits floating-point
47
+ // contraction -- contraction would keep an unrounded wider product and silently
48
+ // make this op more accurate than the graph it replaces.
49
+ const HIDDEN: u32 = {{ hidden }}u;
50
+
51
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
52
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
53
+ {{ flat_index_2d() }}
54
+ {% if vec4 %}
55
+ // A vec4 group is four consecutive channels of one row: HIDDEN % 4 == 0 stops
56
+ // it from ever straddling two rows, so the whole group shares one gate value.
57
+ let g = vec4<{{ scalar }}>(gate[i * 4u / HIDDEN]);
58
+ output[i] = x[i] + fma(y[i], g, vec4<{{ scalar }}>(0.0));
59
+ {% else %}
60
+ let g = gate[i / HIDDEN];
61
+ output[i] = x[i] + fma(y[i], g, {{ scalar }}(0.0));
62
+ {% endif %}
63
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,147 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
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
+ "bindings": [
107
+ {
108
+ "name": "x",
109
+ "arg": "X",
110
+ "semantic": "X",
111
+ "buffer": { "type": "read-only-storage" },
112
+ "elementType": "$scalar"
113
+ },
114
+ {
115
+ "name": "y",
116
+ "arg": "Y",
117
+ "semantic": "Y",
118
+ "buffer": { "type": "read-only-storage" },
119
+ "elementType": "$scalar"
120
+ },
121
+ {
122
+ "name": "gate",
123
+ "arg": "gate",
124
+ "semantic": "gate",
125
+ "buffer": { "type": "read-only-storage" },
126
+ "elementType": "$scalar"
127
+ },
128
+ {
129
+ "name": "output",
130
+ "arg": "output",
131
+ "semantic": "output",
132
+ "buffer": { "type": "storage" },
133
+ "elementType": "$scalar"
134
+ },
135
+ {
136
+ "name": "params",
137
+ "semantic": "kernel.params",
138
+ "buffer": { "type": "uniform" },
139
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
140
+ }
141
+ ],
142
+ "dispatch": { "threads": "numel(shapes.X)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
143
+ }
144
+ ]
145
+ }
146
+ ]
147
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,18 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "com.microsoft.GatedAdd",
3
+ "id": "_com_microsoft_gatedadd_webgpu_c9e5b9d",
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": "GiWIIFkb2GXB3/MDWr4Buyc/kKQrH5xr2tcee1RCOog=",
12
+ "manifest.json": "eGSFQTnmYL4/Zqwfz8zkKFJgFT1Fn5IjP+slAJPCuTU=",
13
+ "test.json": "5Gsy9SVd5u5t8bVgKOLTcNMfO4wqI7rFSjHIhPpLEgk="
14
+ }
15
+ },
16
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
17
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.GatedAdd" }
18
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,185 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.GatedAdd",
3
+ "cases": [
4
+ {
5
+ "name": "rank3_rows_f32",
6
+ "inputs": {
7
+ "X": {
8
+ "dtype": "float32",
9
+ "shape": [2, 3, 8],
10
+ "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
11
+ },
12
+ "Y": {
13
+ "dtype": "float32",
14
+ "shape": [2, 3, 8],
15
+ "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07 }
16
+ },
17
+ "gate": {
18
+ "dtype": "float32",
19
+ "shape": [2, 3, 1],
20
+ "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.17, "offset": 0.75 }
21
+ }
22
+ },
23
+ "outputs": {
24
+ "output": { "dtype": "float32", "shape": [2, 3, 8], "tolerance": 0.000001, "relTolerance": 0.000001 }
25
+ }
26
+ },
27
+ {
28
+ "name": "rank2_channels_odd_f32",
29
+ "inputs": {
30
+ "X": {
31
+ "dtype": "float32",
32
+ "shape": [5, 7],
33
+ "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.23 }
34
+ },
35
+ "Y": {
36
+ "dtype": "float32",
37
+ "shape": [5, 7],
38
+ "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11 }
39
+ },
40
+ "gate": {
41
+ "dtype": "float32",
42
+ "shape": [5, 1],
43
+ "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.29, "offset": -1.25 }
44
+ }
45
+ },
46
+ "outputs": { "output": { "dtype": "float32", "shape": [5, 7], "tolerance": 0.000001, "relTolerance": 0.000001 } }
47
+ },
48
+ {
49
+ "name": "rank2_channels_not_vec4_aligned_f32",
50
+ "inputs": {
51
+ "X": {
52
+ "dtype": "float32",
53
+ "shape": [4, 6],
54
+ "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.43 }
55
+ },
56
+ "Y": {
57
+ "dtype": "float32",
58
+ "shape": [4, 6],
59
+ "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.13 }
60
+ },
61
+ "gate": { "dtype": "float32", "shape": [4, 1], "data": { "kind": "values", "values": [2.0, -3.0, 0.5, 7.0] } }
62
+ },
63
+ "outputs": { "output": { "dtype": "float32", "shape": [4, 6], "tolerance": 0.000001, "relTolerance": 0.000001 } }
64
+ },
65
+ {
66
+ "name": "rank1_single_row_f32",
67
+ "inputs": {
68
+ "X": { "dtype": "float32", "shape": [16], "data": { "kind": "fillFloat32", "sinStep": 0.27, "cosStep": 0.09 } },
69
+ "Y": { "dtype": "float32", "shape": [16], "data": { "kind": "fillFloat32", "sinStep": 0.47, "cosStep": 0.21 } },
70
+ "gate": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-1.75] } }
71
+ },
72
+ "outputs": { "output": { "dtype": "float32", "shape": [16], "tolerance": 0.000001, "relTolerance": 0.000001 } }
73
+ },
74
+ {
75
+ "name": "rank4_rows_f32",
76
+ "inputs": {
77
+ "X": {
78
+ "dtype": "float32",
79
+ "shape": [2, 2, 3, 4],
80
+ "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.37 }
81
+ },
82
+ "Y": {
83
+ "dtype": "float32",
84
+ "shape": [2, 2, 3, 4],
85
+ "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.59 }
86
+ },
87
+ "gate": {
88
+ "dtype": "float32",
89
+ "shape": [2, 2, 3, 1],
90
+ "data": { "kind": "fillFloat32", "sinStep": 0.71, "cosStep": 0.19, "offset": 1.5 }
91
+ }
92
+ },
93
+ "outputs": {
94
+ "output": { "dtype": "float32", "shape": [2, 2, 3, 4], "tolerance": 0.000001, "relTolerance": 0.000001 }
95
+ }
96
+ },
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": {
104
+ "dtype": "float32",
105
+ "shape": [2, 3],
106
+ "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
107
+ },
108
+ "Y": {
109
+ "dtype": "float32",
110
+ "shape": [2, 3],
111
+ "data": { "kind": "values", "values": [0.5, -1.0, 2.0, 10.0, -0.25, 0.125] }
112
+ },
113
+ "gate": { "dtype": "float32", "shape": [2, 1], "data": { "kind": "values", "values": [2.0, -0.5] } }
114
+ },
115
+ "outputs": {
116
+ "output": {
117
+ "dtype": "float32",
118
+ "shape": [2, 3],
119
+ "data": { "kind": "values", "values": [2.0, 0.0, 7.0, -1.0, 5.125, 5.9375] }
120
+ }
121
+ }
122
+ },
123
+ {
124
+ "name": "f16_product_rounds_to_type_pinned",
125
+ "provenance": {
126
+ "notes": "Pins the round_to_T rule that separates this op from a wider-precision fused multiply-add. Row 0 uses gate = 1 + 2^-10, so Y = 1025 gives a real product of 1026.0009765625 that float16 rounds to 1026.0; X = -1026 then cancels it exactly to 0. A kernel that let the product stay unrounded -- by contracting the multiply into the add -- would return 2^-10 there instead. Every other value is exact in float16, so the whole expectation is hand-computable."
127
+ },
128
+ "inputs": {
129
+ "X": {
130
+ "dtype": "float16",
131
+ "shape": [2, 4],
132
+ "data": { "kind": "values", "values": [-1026.0, -1000.0, 0.5, -8.0, 0.5, -1.0, 0.25, 3.0] }
133
+ },
134
+ "Y": {
135
+ "dtype": "float16",
136
+ "shape": [2, 4],
137
+ "data": { "kind": "values", "values": [1025.0, 1024.0, 512.0, 8.0, 1.0, 2.0, 3.0, 4.0] }
138
+ },
139
+ "gate": { "dtype": "float16", "shape": [2, 1], "data": { "kind": "values", "values": [1.0009765625, 2.0] } }
140
+ },
141
+ "outputs": {
142
+ "output": {
143
+ "dtype": "float16",
144
+ "shape": [2, 4],
145
+ "data": { "kind": "values", "values": [0.0, 25.0, 513.0, 0.0078125, 2.5, 3.0, 6.25, 11.0] },
146
+ "tolerance": 0.0005
147
+ }
148
+ }
149
+ },
150
+ {
151
+ "name": "f16_rows",
152
+ "inputs": {
153
+ "X": {
154
+ "dtype": "float16",
155
+ "shape": [3, 16],
156
+ "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "offset": 1.25 }
157
+ },
158
+ "Y": {
159
+ "dtype": "float16",
160
+ "shape": [3, 16],
161
+ "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07, "offset": -0.75 }
162
+ },
163
+ "gate": { "dtype": "float16", "shape": [3, 1], "data": { "kind": "values", "values": [1.5, -0.5, 2.25] } }
164
+ },
165
+ "outputs": { "output": { "dtype": "float16", "shape": [3, 16], "tolerance": 0.0005, "relTolerance": 0.002 } }
166
+ },
167
+ {
168
+ "name": "f16_channels_odd",
169
+ "inputs": {
170
+ "X": {
171
+ "dtype": "float16",
172
+ "shape": [4, 5],
173
+ "data": { "kind": "fillFloat32", "sinStep": 0.43, "cosStep": 0.17, "offset": -1.5 }
174
+ },
175
+ "Y": {
176
+ "dtype": "float16",
177
+ "shape": [4, 5],
178
+ "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.53, "offset": 0.875 }
179
+ },
180
+ "gate": { "dtype": "float16", "shape": [4, 1], "data": { "kind": "values", "values": [0.75, -1.25, 3.0, 0.5] } }
181
+ },
182
+ "outputs": { "output": { "dtype": "float16", "shape": [4, 5], "tolerance": 0.0005, "relTolerance": 0.002 } }
183
+ }
184
+ ]
185
+ }