Hugging Face's logo

Kernels:
webgpu-kernels
/
com.microsoft.BiasAdd

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
a44211c
·
verified ·
1 Parent(s): 0ef25e5

sync 2e7068faf55e

Browse files
Files changed (6) hide show
  1. README.md +62 -0
  2. build/webgpu/bench.json +42 -0
  3. build/webgpu/bias-add.wgsl.jinja +68 -0
  4. build/webgpu/manifest.json +148 -0
  5. build/webgpu/metadata.json +18 -0
  6. build/webgpu/test.json +221 -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.BiasAdd
10
+
11
+ `com.microsoft` · ONNX Runtime contrib operator · contrib since_version 1
12
+
13
+ ## Description
14
+
15
+ Adds a 1-D `bias` (broadcast over the channel dimension) to input `X`, then adds the residual tensor `skip` elementwise. All three tensors share the same channel count `C`; `X` and `skip` have shape `(N, S, C)`.
16
+
17
+ See the [ONNX Runtime `BiasAdd` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.BiasAdd) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `X` | `X` | `T` | `3` | — | Input tensor of shape `(N, S, C)`: batch size `N`, spatial size `S`, and `C` channels. | required |
24
+ | `bias` | `bias` | `T` | `1` | — | 1-D bias vector of length C, broadcast-added along the channel dimension. | required |
25
+ | `skip` | `skip` | `T` | `3` | — | Residual tensor with the same `(N, S, C)` shape as `X`, added after the bias. | required |
26
+
27
+ ## Outputs
28
+
29
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
30
+ | --- | --- | --- | --- | --- | --- | --- |
31
+ | `Y` | `Y` | `T` | `3` | same as `X` | Output tensor of shape `(N, S, C)`: the elementwise sum `X + bias + skip`. | 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
+ - [`bias-add.wgsl.jinja`](build/webgpu/bias-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.BiasAdd", { version: 1 });
60
+ const { Y } = await kernel({
61
+ X: { data: XData, shape: [1, 2, 4] },
62
+ bias: { data: biasData, shape: [4] },
63
+ skip: { data: skipData, shape: [1, 2, 4] },
64
+ });
65
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,42 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.BiasAdd",
3
+ "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
+ "cases": [
5
+ {
6
+ "name": "biasadd-f32-2x4096x320",
7
+ "preset": "smoke",
8
+ "vars": { "dtype": "float32" },
9
+ "inputs": {
10
+ "X": { "shape": [2, 4096, 320], "dtype": "float32", "dist": "normal", "seed": 510, "scale": 2 },
11
+ "bias": { "shape": [320], "dtype": "float32", "dist": "normal", "seed": 511, "scale": 1 },
12
+ "skip": { "shape": [2, 4096, 320], "dtype": "float32", "dist": "normal", "seed": 512, "scale": 2 }
13
+ },
14
+ "outputs": { "Y": { "shape": [2, 4096, 320], "dtype": "float32" } },
15
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4096 * 320 * 3 * dtypeBytes(args.dtype)" }] }
16
+ },
17
+ {
18
+ "name": "biasadd-f16-2x4096x320",
19
+ "preset": "model",
20
+ "vars": { "dtype": "float16" },
21
+ "inputs": {
22
+ "X": { "shape": [2, 4096, 320], "dtype": "float16", "dist": "normal", "seed": 513, "scale": 2 },
23
+ "bias": { "shape": [320], "dtype": "float16", "dist": "normal", "seed": 514, "scale": 1 },
24
+ "skip": { "shape": [2, 4096, 320], "dtype": "float16", "dist": "normal", "seed": 515, "scale": 2 }
25
+ },
26
+ "outputs": { "Y": { "shape": [2, 4096, 320], "dtype": "float16" } },
27
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4096 * 320 * 3 * dtypeBytes(args.dtype)" }] }
28
+ },
29
+ {
30
+ "name": "biasadd-f32-scalar-cliff-hidden3",
31
+ "preset": "stress",
32
+ "vars": { "dtype": "float32" },
33
+ "inputs": {
34
+ "X": { "shape": [1, 1398101, 3], "dtype": "float32", "dist": "normal", "seed": 5301, "scale": 2 },
35
+ "bias": { "shape": [3], "dtype": "float32", "dist": "normal", "seed": 5302, "scale": 1 },
36
+ "skip": { "shape": [1, 1398101, 3], "dtype": "float32", "dist": "normal", "seed": 5303, "scale": 2 }
37
+ },
38
+ "outputs": { "Y": { "shape": [1, 1398101, 3], "dtype": "float32" } },
39
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "1398101 * 3 * 3 * dtypeBytes(args.dtype)" }] }
40
+ }
41
+ ]
42
+ }
build/webgpu/bias-add.wgsl.jinja ADDED
@@ -0,0 +1,68 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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.BiasAdd : Y = X + bias + skip
40
+ // X, skip : same shape (commonly rank-3 [N,S,C]), same numel as Y.
41
+ // bias : rank-1 [C] broadcast over the innermost (channel) axis:
42
+ // bias index = element_index % HIDDEN.
43
+ // The `vec4` path requires HIDDEN % 4 == 0 and numel % 4 == 0 so a vec4 group of
44
+ // 4 consecutive elements never crosses the channel axis (the bias slice is then
45
+ // contiguous). x and skip retain scalar read-only bindings; y uses the vec4
46
+ // output binding shared with the vectorized broadcast path. y is a distinct
47
+ // storage buffer and never aliases skip; mixing read-only and read_write usage
48
+ // for one buffer would invalidate the command buffer.
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
+ let base = i * 4u;
56
+ let xv = vec4<f32>(f32(x[base]), f32(x[base + 1u]), f32(x[base + 2u]), f32(x[base + 3u]));
57
+ let sv = vec4<f32>(f32(skip[base]), f32(skip[base + 1u]), f32(skip[base + 2u]), f32(skip[base + 3u]));
58
+ let bcol = base % HIDDEN;
59
+ let bv = vec4<f32>(f32(bias[bcol]), f32(bias[bcol + 1u]), f32(bias[bcol + 2u]), f32(bias[bcol + 3u]));
60
+ let v = xv + bv + sv;
61
+ y[i] = vec4<{{ scalar }}>(v);
62
+ {% else %}
63
+ let xv = f32(x[i]);
64
+ let sv = f32(skip[i]);
65
+ let v = xv + f32(bias[i % HIDDEN]) + sv;
66
+ y[i] = {{ scalar }}(v);
67
+ {% endif %}
68
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,148 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "com.microsoft",
3
+ "name": "BiasAdd",
4
+ "sinceVersion": 1,
5
+ "description": "Adds a 1-D `bias` (broadcast over the channel dimension) to input `X`, then adds the residual tensor `skip` elementwise. All three tensors share the same channel count `C`; `X` and `skip` have shape `(N, S, C)`.",
6
+ "inputs": [
7
+ {
8
+ "role": "X",
9
+ "dtype": "T",
10
+ "rank": 3,
11
+ "description": "Input tensor of shape `(N, S, C)`: batch size `N`, spatial size `S`, and `C` channels."
12
+ },
13
+ {
14
+ "role": "bias",
15
+ "dtype": "T",
16
+ "rank": 1,
17
+ "description": "1-D bias vector of length C, broadcast-added along the channel dimension."
18
+ },
19
+ {
20
+ "role": "skip",
21
+ "dtype": "T",
22
+ "rank": 3,
23
+ "description": "Residual tensor with the same `(N, S, C)` shape as `X`, added after the bias."
24
+ }
25
+ ],
26
+ "outputs": [
27
+ {
28
+ "role": "Y",
29
+ "dtype": "T",
30
+ "rank": 3,
31
+ "shape": "shapes.X",
32
+ "description": "Output tensor of shape `(N, S, C)`: the elementwise sum `X + bias + skip`."
33
+ }
34
+ ],
35
+ "typeConstraints": { "T": ["float32", "float16"] },
36
+ "args": {
37
+ "X": { "kind": "tensor", "semantic": "X", "role": "input" },
38
+ "bias": { "kind": "tensor", "semantic": "bias", "role": "input" },
39
+ "skip": { "kind": "tensor", "semantic": "skip", "role": "input" },
40
+ "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
41
+ },
42
+ "tunables": { "WORKGROUP_SIZE": 256 },
43
+ "constants": {
44
+ "scalar": "dtypes.T",
45
+ "usesF16": "dtypes.T == \"f16\"",
46
+ "hidden": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 1"
47
+ },
48
+ "variants": [
49
+ {
50
+ "id": "vec4",
51
+ "priority": 30,
52
+ "when": ["ranks.X == 3", "ranks.skip == 3", "ranks.Y == 3", "sameShape(shapes.X, shapes.skip)", "sameShape(shapes.X, shapes.Y)", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "numel(shapes.X) > 0", "numel(shapes.X) % 4 == 0", "dim(shapes.X, 2) % 4 == 0"],
53
+ "constants": { "vec4": true, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
54
+ "passes": [
55
+ {
56
+ "id": "main",
57
+ "name": "BiasAdd.vec4",
58
+ "shader": "bias-add.wgsl.jinja",
59
+ "bindings": [
60
+ {
61
+ "name": "x",
62
+ "arg": "X",
63
+ "semantic": "X",
64
+ "buffer": { "type": "read-only-storage" },
65
+ "elementType": "$scalar"
66
+ },
67
+ {
68
+ "name": "bias",
69
+ "arg": "bias",
70
+ "semantic": "bias",
71
+ "buffer": { "type": "read-only-storage" },
72
+ "elementType": "$scalar",
73
+ "length": "$hidden"
74
+ },
75
+ {
76
+ "name": "skip",
77
+ "arg": "skip",
78
+ "semantic": "skip",
79
+ "buffer": { "type": "read-only-storage" },
80
+ "elementType": "$scalar"
81
+ },
82
+ {
83
+ "name": "y",
84
+ "arg": "Y",
85
+ "semantic": "Y",
86
+ "buffer": { "type": "storage" },
87
+ "elementType": "$vectorScalar"
88
+ },
89
+ {
90
+ "name": "params",
91
+ "semantic": "kernel.params",
92
+ "buffer": { "type": "uniform" },
93
+ "struct": {
94
+ "name": "Params",
95
+ "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }]
96
+ }
97
+ }
98
+ ],
99
+ "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
100
+ }
101
+ ]
102
+ },
103
+ {
104
+ "id": "scalar",
105
+ "priority": 0,
106
+ "when": ["ranks.X == 3", "ranks.skip == 3", "ranks.Y == 3", "sameShape(shapes.X, shapes.skip)", "sameShape(shapes.X, shapes.Y)", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)"],
107
+ "constants": { "vec4": false },
108
+ "passes": [
109
+ {
110
+ "id": "main",
111
+ "name": "BiasAdd.scalar",
112
+ "shader": "bias-add.wgsl.jinja",
113
+ "bindings": [
114
+ {
115
+ "name": "x",
116
+ "arg": "X",
117
+ "semantic": "X",
118
+ "buffer": { "type": "read-only-storage" },
119
+ "elementType": "$scalar"
120
+ },
121
+ {
122
+ "name": "bias",
123
+ "arg": "bias",
124
+ "semantic": "bias",
125
+ "buffer": { "type": "read-only-storage" },
126
+ "elementType": "$scalar"
127
+ },
128
+ {
129
+ "name": "skip",
130
+ "arg": "skip",
131
+ "semantic": "skip",
132
+ "buffer": { "type": "read-only-storage" },
133
+ "elementType": "$scalar"
134
+ },
135
+ { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
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": "numel(shapes.X)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
144
+ }
145
+ ]
146
+ }
147
+ ]
148
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,18 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "com.microsoft.BiasAdd",
3
+ "id": "_com_microsoft_biasadd_webgpu_e97b808",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "W2cQSIIm7hI8MbHD6Cro9DGB/ByB035ELwNigWwGHYw=",
11
+ "bias-add.wgsl.jinja": "VJnmyJ3GrZrzADYN5uP73XMs4i+YXvC39ECRQ9j40Ow=",
12
+ "manifest.json": "pAy9IP3Uvl8QAonrSEw9nzg8kJxnbXeRZXoQql+sxe4=",
13
+ "test.json": "PLZQbRu0fhmrRSakHOzPjTiZL0LYBLFxmBdw7gd31m4="
14
+ }
15
+ },
16
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
17
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.BiasAdd" }
18
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,221 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.BiasAdd",
3
+ "cases": [
4
+ {
5
+ "name": "dispatch_cliff_scalar_f32",
6
+ "inputs": {
7
+ "X": { "dtype": "float32", "shape": [1, 5592321, 3], "data": { "kind": "linspace", "start": -1.0, "end": 1.0 } },
8
+ "bias": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, -1.0, 2.0] } },
9
+ "skip": { "dtype": "float32", "shape": [1, 5592321, 3], "data": { "kind": "constant", "value": 0.25 } }
10
+ },
11
+ "outputs": { "Y": { "dtype": "float32", "shape": [1, 5592321, 3], "tolerance": 0.00001 } }
12
+ },
13
+ {
14
+ "name": "ort_fused_skip_bias_channels",
15
+ "provenance": {
16
+ "source": "onnxruntime/test/contrib_ops/bias_add_op_test.cc",
17
+ "test": "BiasAddTest.BiasAddTest_HiddenSize_320",
18
+ "notes": "Small deterministic projection of the ORT fused skip+bias channel pattern."
19
+ },
20
+ "inputs": {
21
+ "X": {
22
+ "dtype": "float32",
23
+ "shape": [2, 2, 3],
24
+ "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, -2.0, -3.0, -4.0] }
25
+ },
26
+ "bias": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, -1.0, 2.0] } },
27
+ "skip": {
28
+ "dtype": "float32",
29
+ "shape": [2, 2, 3],
30
+ "data": { "kind": "values", "values": [10.0, 10.0, 10.0, -1.0, -1.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
31
+ }
32
+ },
33
+ "outputs": {
34
+ "Y": {
35
+ "dtype": "float32",
36
+ "shape": [2, 2, 3],
37
+ "data": { "kind": "values", "values": [9.5, 9.0, 13.0, 1.5, 1.0, 5.0, 5.5, 6.0, 11.0, 1.5, 0.0, 3.0] }
38
+ }
39
+ }
40
+ },
41
+ {
42
+ "name": "ort_fused_skip_bias_singleton_image_f16",
43
+ "provenance": {
44
+ "source": "onnxruntime/test/contrib_ops/bias_add_op_test.cc",
45
+ "test": "BiasAddTest.BiasAddTest_HiddenSize_640",
46
+ "notes": "Small deterministic float16 projection of the ORT singleton image-size path."
47
+ },
48
+ "inputs": {
49
+ "X": {
50
+ "dtype": "float16",
51
+ "shape": [2, 1, 4],
52
+ "data": { "kind": "values", "values": [-1.0, -0.5, 0.5, 1.0, 2.0, -2.0, 0.25, -0.25] }
53
+ },
54
+ "bias": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.5, -1.0, 1.5, -0.5] } },
55
+ "skip": {
56
+ "dtype": "float16",
57
+ "shape": [2, 1, 4],
58
+ "data": { "kind": "values", "values": [1.0, 2.0, -1.0, -2.0, -0.5, 0.5, 1.0, -1.0] }
59
+ }
60
+ },
61
+ "outputs": {
62
+ "Y": {
63
+ "dtype": "float16",
64
+ "shape": [2, 1, 4],
65
+ "data": { "kind": "values", "values": [0.5, 0.5, 1.0, -1.5, 2.0, -2.5, 2.75, -1.75] }
66
+ }
67
+ }
68
+ },
69
+ {
70
+ "name": "ort_fused_skip_bias_vec4_f32",
71
+ "provenance": {
72
+ "source": "onnxruntime/test/contrib_ops/bias_add_op_test.cc",
73
+ "test": "BiasAddTest.BiasAddTest_HiddenSize_1280",
74
+ "notes": "Small deterministic projection with hidden size divisible by 4, matching the ORT channel-bias pattern and exercising the vec4 path."
75
+ },
76
+ "inputs": {
77
+ "X": {
78
+ "dtype": "float32",
79
+ "shape": [1, 2, 4],
80
+ "data": { "kind": "values", "values": [-4.0, -1.0, 0.0, 2.0, 5.0, 6.0, -7.0, 8.0] }
81
+ },
82
+ "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, -0.25, 1.0, -2.0] } },
83
+ "skip": {
84
+ "dtype": "float32",
85
+ "shape": [1, 2, 4],
86
+ "data": { "kind": "values", "values": [10.0, 0.0, -1.0, 1.0, -5.0, 2.0, 3.0, -4.0] }
87
+ }
88
+ },
89
+ "outputs": {
90
+ "Y": {
91
+ "dtype": "float32",
92
+ "shape": [1, 2, 4],
93
+ "tolerance": 0,
94
+ "data": { "kind": "values", "values": [6.5, -1.25, 0.0, 1.0, 0.5, 7.75, -3.0, 2.0] }
95
+ }
96
+ }
97
+ },
98
+ {
99
+ "name": "f32_hidden320_vec4_compact",
100
+ "inputs": {
101
+ "X": {
102
+ "dtype": "float32",
103
+ "shape": [2, 16, 320],
104
+ "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.5 }
105
+ },
106
+ "bias": {
107
+ "dtype": "float32",
108
+ "shape": [320],
109
+ "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.17, "scale": 0.25 }
110
+ },
111
+ "skip": {
112
+ "dtype": "float32",
113
+ "shape": [2, 16, 320],
114
+ "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 0.25 }
115
+ }
116
+ },
117
+ "outputs": { "Y": { "dtype": "float32", "shape": [2, 16, 320], "tolerance": 0.000001 } }
118
+ },
119
+ {
120
+ "name": "ort_fused_skip_bias_scalar_f16",
121
+ "provenance": {
122
+ "source": "onnxruntime/test/contrib_ops/bias_add_op_test.cc",
123
+ "test": "BiasAddTest.BiasAddTest_HiddenSize_768",
124
+ "notes": "Small deterministic float16 projection with hidden size not divisible by 4, so the scalar path covers ORT's fused skip+bias pattern."
125
+ },
126
+ "inputs": {
127
+ "X": {
128
+ "dtype": "float16",
129
+ "shape": [1, 3, 3],
130
+ "data": { "kind": "values", "values": [0.5, -0.5, 1.0, 2.0, -2.0, 0.25, -0.25, 4.0, -4.0] }
131
+ },
132
+ "bias": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.25, -1.5, 2.0] } },
133
+ "skip": {
134
+ "dtype": "float16",
135
+ "shape": [1, 3, 3],
136
+ "data": { "kind": "values", "values": [1.0, 2.0, 3.0, -1.0, -2.0, -3.0, 0.5, -0.5, 1.0] }
137
+ }
138
+ },
139
+ "outputs": {
140
+ "Y": {
141
+ "dtype": "float16",
142
+ "shape": [1, 3, 3],
143
+ "tolerance": 0,
144
+ "data": { "kind": "values", "values": [1.75, 0.0, 6.0, 1.25, -5.5, -0.75, 0.5, 2.0, -1.0] }
145
+ }
146
+ }
147
+ },
148
+ {
149
+ "name": "ort_fused_skip_bias_image3_vec4_f32",
150
+ "provenance": {
151
+ "source": "onnxruntime/test/contrib_ops/bias_add_op_test.cc",
152
+ "test": "BiasAddTest.BiasAddTest_HiddenSize_1536",
153
+ "notes": "Small deterministic projection of ORT's batch=1, image_size=3, channel-bias path with hidden size divisible by 4."
154
+ },
155
+ "inputs": {
156
+ "X": {
157
+ "dtype": "float32",
158
+ "shape": [1, 3, 4],
159
+ "data": { "kind": "values", "values": [-2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, 4.0, -4.0, 0.5, -0.5] }
160
+ },
161
+ "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, -0.5, 1.5, -1.25] } },
162
+ "skip": {
163
+ "dtype": "float32",
164
+ "shape": [1, 3, 4],
165
+ "data": { "kind": "values", "values": [10.0, 0.0, -10.0, 2.0, -1.0, 1.0, 3.0, -3.0, 0.5, -0.5, 2.0, -2.0] }
166
+ }
167
+ },
168
+ "outputs": {
169
+ "Y": {
170
+ "dtype": "float32",
171
+ "shape": [1, 3, 4],
172
+ "tolerance": 0,
173
+ "data": {
174
+ "kind": "values",
175
+ "values": [8.25, -1.5, -8.5, 1.75, 1.25, 3.5, 1.5, -6.25, 4.75, -5.0, 4.0, -3.75]
176
+ }
177
+ }
178
+ }
179
+ },
180
+ {
181
+ "name": "empty_zero_dim",
182
+ "inputs": {
183
+ "X": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } },
184
+ "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4] } },
185
+ "skip": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } }
186
+ },
187
+ "outputs": { "Y": { "dtype": "float32", "shape": [1, 0, 4], "tolerance": 0 } }
188
+ },
189
+ {
190
+ "name": "empty_zero_dim_f16",
191
+ "inputs": {
192
+ "X": { "dtype": "float16", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } },
193
+ "bias": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4] } },
194
+ "skip": { "dtype": "float16", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } }
195
+ },
196
+ "outputs": { "Y": { "dtype": "float16", "shape": [1, 0, 4], "tolerance": 0 } }
197
+ },
198
+ {
199
+ "name": "empty_zero_hidden",
200
+ "inputs": {
201
+ "X": { "dtype": "float32", "shape": [1, 2, 0], "data": { "kind": "values", "values": [] } },
202
+ "bias": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } },
203
+ "skip": { "dtype": "float32", "shape": [1, 2, 0], "data": { "kind": "values", "values": [] } }
204
+ },
205
+ "outputs": { "Y": { "dtype": "float32", "shape": [1, 2, 0], "tolerance": 0 } }
206
+ },
207
+ {
208
+ "name": "scalar_hidden5_bias_broadcast_row_wrap",
209
+ "inputs": {
210
+ "X": { "dtype": "float32", "shape": [1, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 2.0 } },
211
+ "bias": {
212
+ "dtype": "float32",
213
+ "shape": [5],
214
+ "data": { "kind": "values", "values": [0.5, -1.0, 2.0, -0.25, 1.5] }
215
+ },
216
+ "skip": { "dtype": "float32", "shape": [1, 3, 5], "data": { "kind": "linspace", "start": 1.0, "end": -1.0 } }
217
+ },
218
+ "outputs": { "Y": { "dtype": "float32", "shape": [1, 3, 5], "tolerance": 0.000001 } }
219
+ }
220
+ ]
221
+ }