Hugging Face's logo

Kernels:
webgpu-kernels
/
com.microsoft.BiasGelu

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

sync 2e7068faf55e

Browse files
Files changed (6) hide show
  1. README.md +57 -0
  2. build/webgpu/bench.json +80 -0
  3. build/webgpu/elementwise-bias-gelu.wgsl.jinja +90 -0
  4. build/webgpu/manifest.json +160 -0
  5. build/webgpu/metadata.json +18 -0
  6. build/webgpu/test.json +243 -0
README.md CHANGED
@@ -1,3 +1,60 @@
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.BiasGelu
10
+
11
+ `com.microsoft` · ONNX Runtime contrib operator · contrib since_version 1
12
+
13
+ ## Description
14
+
15
+ Applies GELU to `A + B`, where the 1-D bias `B` is broadcast along the last dimension of `A`. This implementation supports float16 and float32; the schema's double and bfloat16 types are not implemented.
16
+
17
+ See the [ONNX Runtime `BiasGelu` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.BiasGelu) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `A` | `A` | `T` | — | — | The main input tensor of any shape. | required |
24
+ | `B` | `B` | `T` | `1` | — | 1-D bias tensor whose length equals the last dimension of `A`. | required |
25
+
26
+ ## Outputs
27
+
28
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
29
+ | --- | --- | --- | --- | --- | --- | --- |
30
+ | `C` | `C` | `T` | same as `A` | same as `A` | Output tensor after applying GELU to `A + B`; same shape as `A`. | required |
31
+
32
+ ## Type constraints
33
+
34
+ | Variable | Allowed dtypes |
35
+ | --- | --- |
36
+ | `T` | `float32`, `float16` |
37
+
38
+ ## Files
39
+
40
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
41
+ - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
42
+ - [`test.json`](build/webgpu/test.json) — correctness cases
43
+ - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
44
+ - [`elementwise-bias-gelu.wgsl.jinja`](build/webgpu/elementwise-bias-gelu.wgsl.jinja)
45
+
46
+ ## Use with `@huggingface/kernels`
47
+
48
+ The loader derives every required output's shape and logical dtype from the manifest contract and this call.
49
+ It then allocates the result tensors automatically.
50
+
51
+ The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
52
+
53
+ Replace each `*Data` placeholder with a typed array containing the corresponding input data.
54
+
55
+ ```js
56
+ import { getKernel } from "@huggingface/kernels";
57
+
58
+ const kernel = await getKernel("webgpu-kernels/com.microsoft.BiasGelu", { version: 1 });
59
+ const { C } = await kernel({ A: { data: AData, shape: [3] }, B: { data: BData, shape: [3] } });
60
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,80 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.BiasGelu",
3
+ "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
+ "cases": [
5
+ {
6
+ "name": "biasgelu-f32-4096x3072",
7
+ "preset": "smoke",
8
+ "vars": { "dtype": "float32" },
9
+ "inputs": {
10
+ "A": { "shape": [4096, 3072], "dtype": "float32", "dist": "normal", "seed": 420, "scale": 2 },
11
+ "B": { "shape": [3072], "dtype": "float32", "dist": "normal", "seed": 421, "scale": 1 }
12
+ },
13
+ "outputs": { "C": { "shape": [4096, 3072], "dtype": "float32" } },
14
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "4096 * 3072 * 2 * dtypeBytes(args.dtype)" }] }
15
+ },
16
+ {
17
+ "name": "biasgelu-f32-scalar-large-odd-hidden",
18
+ "preset": "stress",
19
+ "vars": { "dtype": "float32" },
20
+ "inputs": {
21
+ "A": { "shape": [2097152, 3], "dtype": "float32", "dist": "normal", "seed": 430, "scale": 2 },
22
+ "B": { "shape": [3], "dtype": "float32", "dist": "normal", "seed": 431, "scale": 1 }
23
+ },
24
+ "outputs": { "C": { "shape": [2097152, 3], "dtype": "float32" } },
25
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "2097152 * 3 * 2 * dtypeBytes(args.dtype)" }] }
26
+ },
27
+ {
28
+ "name": "biasgelu-bert-base-b8-s512-i3072",
29
+ "preset": "model",
30
+ "provenance": {
31
+ "notes": "BERT class defaults (hidden_size 768, intermediate_size 3072): the FFN bias+GELU for a batch of 8 full-length sequences."
32
+ },
33
+ "vars": { "dtype": "float32", "rows": 4096, "inter": 3072 },
34
+ "inputs": {
35
+ "A": { "shape": [4096, 3072], "dtype": "float32", "dist": "normal", "seed": 5500, "scale": 0.5 },
36
+ "B": { "shape": [3072], "dtype": "float32", "dist": "normal", "seed": 5501, "scale": 0.1 }
37
+ },
38
+ "outputs": { "C": { "shape": [4096, 3072], "dtype": "float32" } },
39
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.inter * 2 * dtypeBytes(args.dtype)" }] }
40
+ },
41
+ {
42
+ "name": "biasgelu-vit-base-b64-p197-i3072",
43
+ "preset": "model",
44
+ "provenance": {
45
+ "notes": "ViT class defaults (hidden_size 768, intermediate_size 3072, image_size 224, patch_size 16 -> 196 patches + CLS) at a 64-image batch."
46
+ },
47
+ "vars": { "dtype": "float32", "rows": 12608, "inter": 3072 },
48
+ "inputs": {
49
+ "A": { "shape": [12608, 3072], "dtype": "float32", "dist": "normal", "seed": 5600, "scale": 0.5 },
50
+ "B": { "shape": [3072], "dtype": "float32", "dist": "normal", "seed": 5601, "scale": 0.1 }
51
+ },
52
+ "outputs": { "C": { "shape": [12608, 3072], "dtype": "float32" } },
53
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.inter * 2 * dtypeBytes(args.dtype)" }] }
54
+ },
55
+ {
56
+ "name": "biasgelu-t5-base-b8-s512-i2048",
57
+ "preset": "model",
58
+ "provenance": { "notes": "T5 class defaults (d_model 512, d_ff 2048) at a batch of 8 full-length sequences." },
59
+ "vars": { "dtype": "float32", "rows": 4096, "inter": 2048 },
60
+ "inputs": {
61
+ "A": { "shape": [4096, 2048], "dtype": "float32", "dist": "normal", "seed": 5700, "scale": 0.5 },
62
+ "B": { "shape": [2048], "dtype": "float32", "dist": "normal", "seed": 5701, "scale": 0.1 }
63
+ },
64
+ "outputs": { "C": { "shape": [4096, 2048], "dtype": "float32" } },
65
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.inter * 2 * dtypeBytes(args.dtype)" }] }
66
+ },
67
+ {
68
+ "name": "biasgelu-deberta-v2-b8-s512-i6144",
69
+ "preset": "model",
70
+ "provenance": { "notes": "DeBERTa-v2 class defaults (hidden_size 1536, intermediate_size 6144)." },
71
+ "vars": { "dtype": "float32", "rows": 4096, "inter": 6144 },
72
+ "inputs": {
73
+ "A": { "shape": [4096, 6144], "dtype": "float32", "dist": "normal", "seed": 5800, "scale": 0.5 },
74
+ "B": { "shape": [6144], "dtype": "float32", "dist": "normal", "seed": 5801, "scale": 0.1 }
75
+ },
76
+ "outputs": { "C": { "shape": [4096, 6144], "dtype": "float32" } },
77
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.inter * 2 * dtypeBytes(args.dtype)" }] }
78
+ }
79
+ ]
80
+ }
build/webgpu/elementwise-bias-gelu.wgsl.jinja ADDED
@@ -0,0 +1,90 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {% if usesF16 %}
2
+ enable f16;
3
+ {% endif %}
4
+ {{ env.wgsl.resourceDeclarations }}
5
+ {% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
6
+
7
+ // Bias plus GELU, with a specialization-selected tanh or erf approximation.
8
+ // The optional bias is a rank-1 vector broadcast over the innermost (hidden)
9
+ // axis: bias index = element_index % HIDDEN. The `vec4` path requires
10
+ // HIDDEN % 4 == 0 and numel % 4 == 0 so a vec4 group never crosses the hidden
11
+ // axis (the bias slice is then contiguous). `vec4Tail` keeps scalar bindings but
12
+ // evaluates four guarded lanes per invocation, so odd hidden sizes retain the
13
+ // same parallel efficiency without crossing row/bias boundaries. Gelu math and overflow guards match
14
+ // the vectorized unary implementation: tanh saturates to +/-1 by |x|~9, and the
15
+ // erf path uses the same rational approximation as Gelu.
16
+ {% if approximate == "erf" %}
17
+ fn erf_approx(x: f32) -> f32 {
18
+ let ax = abs(x);
19
+ // The polynomial has a small nonzero floor near zero. Use erf(x) ~=
20
+ // 2/sqrt(pi)*x below 2^-20 to preserve erf(0) == 0, odd symmetry, and the
21
+ // correctly rounded f32 result. The exactly representable threshold keeps
22
+ // scalar and vector branching identical. NaN falls through to the polynomial
23
+ // and propagates.
24
+ if (ax < 9.5367431640625e-7) {
25
+ return 1.1283791670955126 * x;
26
+ }
27
+ let sign = select(-1.0, 1.0, x >= 0.0);
28
+ let t = 1.0 / (1.0 + 0.3275911 * ax);
29
+ let y = 1.0 - (((((1.061405429 * t - 1.453152027) * t) + 1.421413741) * t - 0.284496736) * t + 0.254829592) * t * exp(-(ax * ax));
30
+ return sign * y;
31
+ }
32
+ {% else %}
33
+ fn tanh_safe(x: f32) -> f32 {
34
+ if (x > 10.0) { return 1.0; }
35
+ if (x < -10.0) { return -1.0; }
36
+ return tanh(x);
37
+ }
38
+ {% endif %}
39
+ fn gelu_value(v: f32) -> f32 {
40
+ {% if approximate == "erf" %}
41
+ return 0.5 * v * (1.0 + erf_approx(v * 0.7071067811865476));
42
+ {% else %}
43
+ return 0.5 * v * (1.0 + tanh_safe(0.7978845608028654 * (v + 0.044715 * v * v * v)));
44
+ {% endif %}
45
+ }
46
+ {% if hasBias %}
47
+
48
+ const HIDDEN: u32 = {{ hidden }}u;
49
+
50
+ {% endif %}
51
+ @compute @workgroup_size({{ wg }})
52
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
53
+ // 2D-folded flat index: gid.y carries the high bits past the
54
+ // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
55
+ let i = gid.x + gid.y * nwg.x * {{ wg }}u;
56
+ if (i >= params.count) {
57
+ return;
58
+ }
59
+ {% if vec4Tail %}
60
+ let base = i * 4u;
61
+ {% for lane in range(4) %}
62
+ if (base + {{ lane }}u < params.count) {
63
+ let xv{{ lane }} = f32(x[base + {{ lane }}u]);
64
+ {% if hasBias %}
65
+ let v{{ lane }} = xv{{ lane }} + f32(bias[(base + {{ lane }}u) % HIDDEN]);
66
+ {% else %}
67
+ let v{{ lane }} = xv{{ lane }};
68
+ {% endif %}
69
+ y[base + {{ lane }}u] = {{ scalar }}(gelu_value(v{{ lane }}));
70
+ }
71
+ {% endfor %}
72
+ {% elif vec4 %}
73
+ let xv = vec4<f32>(x[i]);
74
+ {% if hasBias %}
75
+ let bcol = (i * 4u) % HIDDEN;
76
+ let v = xv + vec4<f32>(f32(bias[bcol]), f32(bias[bcol + 1u]), f32(bias[bcol + 2u]), f32(bias[bcol + 3u]));
77
+ {% else %}
78
+ let v = xv;
79
+ {% endif %}
80
+ y[i] = vec4<{{ scalar }}>(vec4<f32>(gelu_value(v.x), gelu_value(v.y), gelu_value(v.z), gelu_value(v.w)));
81
+ {% else %}
82
+ let xv = f32(x[i]);
83
+ {% if hasBias %}
84
+ let v = xv + f32(bias[i % HIDDEN]);
85
+ {% else %}
86
+ let v = xv;
87
+ {% endif %}
88
+ y[i] = {{ scalar }}(gelu_value(v));
89
+ {% endif %}
90
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,160 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "com.microsoft",
3
+ "name": "BiasGelu",
4
+ "sinceVersion": 1,
5
+ "description": "Applies GELU to `A + B`, where the 1-D bias `B` is broadcast along the last dimension of `A`. This implementation supports float16 and float32; the schema's double and bfloat16 types are not implemented.",
6
+ "inputs": [
7
+ { "role": "A", "dtype": "T", "description": "The main input tensor of any shape." },
8
+ {
9
+ "role": "B",
10
+ "dtype": "T",
11
+ "rank": 1,
12
+ "description": "1-D bias tensor whose length equals the last dimension of `A`."
13
+ }
14
+ ],
15
+ "outputs": [
16
+ {
17
+ "role": "C",
18
+ "dtype": "T",
19
+ "rank": "ranks.A",
20
+ "shape": "shapes.A",
21
+ "description": "Output tensor after applying GELU to `A + B`; same shape as `A`."
22
+ }
23
+ ],
24
+ "typeConstraints": { "T": ["float32", "float16"] },
25
+ "args": {
26
+ "A": { "kind": "tensor", "semantic": "A", "role": "input" },
27
+ "B": { "kind": "tensor", "semantic": "B", "role": "input" },
28
+ "C": { "kind": "tensor", "semantic": "C", "role": "output" }
29
+ },
30
+ "tunables": { "WORKGROUP_SIZE": 64 },
31
+ "derive": {
32
+ "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
33
+ "workgroupOk": "tunables.WORKGROUP_SIZE > 0 and tunables.WORKGROUP_SIZE <= deviceWorkgroupCap",
34
+ "baseOk": "workgroupOk and ranks.A >= 1 and ranks.C == ranks.A and sameShape(shapes.A, shapes.C) and f16Ok(dtypes.T) and ranks.B == 1 and dim(shapes.B, 0) == dim(shapes.A, ranks.A - 1)",
35
+ "vec4Ok": "numel(shapes.A) > 0 and numel(shapes.A) % 4 == 0 and dim(shapes.A, ranks.A - 1) % 4 == 0"
36
+ },
37
+ "bindingSets": {
38
+ "scalarTail": [
39
+ { "name": "x", "arg": "A", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
40
+ {
41
+ "name": "bias",
42
+ "arg": "B",
43
+ "semantic": "B",
44
+ "buffer": { "type": "read-only-storage" },
45
+ "elementType": "$scalar"
46
+ },
47
+ { "name": "y", "arg": "C", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$scalar" },
48
+ {
49
+ "name": "params",
50
+ "semantic": "kernel.params",
51
+ "buffer": { "type": "uniform" },
52
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.A)" }] }
53
+ }
54
+ ]
55
+ },
56
+ "variants": [
57
+ {
58
+ "id": "vec4",
59
+ "priority": 20,
60
+ "when": ["baseOk", "vec4Ok"],
61
+ "constants": {
62
+ "scalar": "dtypes.T",
63
+ "usesF16": "dtypes.T == \"f16\"",
64
+ "approximate": "\"erf\"",
65
+ "vec4": true,
66
+ "vec4Tail": false,
67
+ "hasBias": true,
68
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
69
+ "hidden": "dim(shapes.A, ranks.A - 1) if dim(shapes.A, ranks.A - 1) > 0 else 1"
70
+ },
71
+ "passes": [
72
+ {
73
+ "id": "main",
74
+ "name": "BiasGelu.vec4",
75
+ "shader": "elementwise-bias-gelu.wgsl.jinja",
76
+ "bindings": [
77
+ {
78
+ "name": "x",
79
+ "arg": "A",
80
+ "semantic": "A",
81
+ "buffer": { "type": "read-only-storage" },
82
+ "elementType": "$vectorScalar"
83
+ },
84
+ {
85
+ "name": "bias",
86
+ "arg": "B",
87
+ "semantic": "B",
88
+ "buffer": { "type": "read-only-storage" },
89
+ "elementType": "$scalar",
90
+ "length": "$hidden"
91
+ },
92
+ {
93
+ "name": "y",
94
+ "arg": "C",
95
+ "semantic": "C",
96
+ "buffer": { "type": "storage" },
97
+ "elementType": "$vectorScalar"
98
+ },
99
+ {
100
+ "name": "params",
101
+ "semantic": "kernel.params",
102
+ "buffer": { "type": "uniform" },
103
+ "struct": {
104
+ "name": "Params",
105
+ "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.A) / 4" }]
106
+ }
107
+ }
108
+ ],
109
+ "dispatch": { "threads": "numel(shapes.A) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
110
+ }
111
+ ]
112
+ },
113
+ {
114
+ "id": "vec4_tail",
115
+ "priority": 10,
116
+ "when": ["baseOk", "numel(shapes.A) > 0"],
117
+ "constants": {
118
+ "scalar": "dtypes.T",
119
+ "usesF16": "dtypes.T == \"f16\"",
120
+ "approximate": "\"erf\"",
121
+ "vec4": false,
122
+ "vec4Tail": true,
123
+ "hasBias": true,
124
+ "hidden": "dim(shapes.A, ranks.A - 1) if dim(shapes.A, ranks.A - 1) > 0 else 1"
125
+ },
126
+ "passes": [
127
+ {
128
+ "id": "main",
129
+ "name": "BiasGelu.vec4Tail",
130
+ "shader": "elementwise-bias-gelu.wgsl.jinja",
131
+ "bindings": "scalarTail",
132
+ "dispatch": { "threads": "ceilDiv(numel(shapes.A), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
133
+ }
134
+ ]
135
+ },
136
+ {
137
+ "id": "scalar",
138
+ "priority": 0,
139
+ "when": ["baseOk", "true"],
140
+ "constants": {
141
+ "scalar": "dtypes.T",
142
+ "usesF16": "dtypes.T == \"f16\"",
143
+ "approximate": "\"erf\"",
144
+ "vec4": false,
145
+ "vec4Tail": false,
146
+ "hasBias": true,
147
+ "hidden": "dim(shapes.A, ranks.A - 1) if dim(shapes.A, ranks.A - 1) > 0 else 1"
148
+ },
149
+ "passes": [
150
+ {
151
+ "id": "main",
152
+ "name": "BiasGelu.scalar",
153
+ "shader": "elementwise-bias-gelu.wgsl.jinja",
154
+ "bindings": "scalarTail",
155
+ "dispatch": { "threads": "numel(shapes.A)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
156
+ }
157
+ ]
158
+ }
159
+ ]
160
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,18 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "com.microsoft.BiasGelu",
3
+ "id": "_com_microsoft_biasgelu_webgpu_0db9e22",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "UiPLr8ZUfUCDcEfplbEhxBCd01I/EdY03qb4n/Qtzyc=",
11
+ "elementwise-bias-gelu.wgsl.jinja": "Eg8N2jJCMce+IsYNcCzuxvs8CSv7yTQXdorvv+c0L58=",
12
+ "manifest.json": "YkUZuF88k5utKaSJd858ft/Msp3otF/iijavqWmOK3E=",
13
+ "test.json": "2XoIgkpiRq0YEMEp23vi/cmZk8pZFcSt7x7vVzsBRyk="
14
+ }
15
+ },
16
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
17
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.BiasGelu" }
18
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,243 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.BiasGelu",
3
+ "cases": [
4
+ {
5
+ "name": "dispatch_cliff_vec4_f32",
6
+ "requires": { "limits": { "maxBufferSize": 268435488, "maxStorageBufferBindingSize": 268435488 } },
7
+ "provenance": {
8
+ "notes": "Its largest tensor is 268435488 bytes, so the case needs an adapter whose maxBufferSize and maxStorageBufferBindingSize both reach it — declared, because the WebGPU guaranteed minimums (256 MiB / 128 MiB) do not, and a device at them must report the case inapplicable rather than fail allocating it."
9
+ },
10
+ "inputs": {
11
+ "A": { "dtype": "float32", "shape": [8388609, 8], "data": { "kind": "linspace", "start": -2.0, "end": 2.0 } },
12
+ "B": {
13
+ "dtype": "float32",
14
+ "shape": [8],
15
+ "data": { "kind": "values", "values": [0.25, -0.25, 0.5, -0.5, 0.1, -0.1, 0.0, 1.0] }
16
+ }
17
+ },
18
+ "outputs": { "C": { "dtype": "float32", "shape": [8388609, 8], "tolerance": 0.0001 } }
19
+ },
20
+ {
21
+ "name": "ort_float32_rank2_bias_vector",
22
+ "provenance": {
23
+ "source": "onnxruntime/test/contrib_ops/element_wise_ops_test.cc",
24
+ "test": "BiasGeluTest.Float",
25
+ "notes": "Small deterministic rank-2 instance of ORT's bias-vector GELU coverage."
26
+ },
27
+ "inputs": {
28
+ "A": {
29
+ "dtype": "float32",
30
+ "shape": [2, 4],
31
+ "data": { "kind": "values", "values": [-1.5, -0.5, 0.5, 1.5, 2.0, -2.0, 0.25, -0.25] }
32
+ },
33
+ "B": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, -0.25, 0.5, -0.5] } }
34
+ },
35
+ "outputs": {
36
+ "C": {
37
+ "dtype": "float32",
38
+ "shape": [2, 4],
39
+ "tolerance": 0.000001,
40
+ "data": {
41
+ "kind": "values",
42
+ "values": [-0.13206221, -0.16997051, 0.8413447, 0.8413447, 2.222495, -0.02750505, 0.5800295, -0.16997051]
43
+ }
44
+ }
45
+ }
46
+ },
47
+ {
48
+ "name": "ort_float16_rank2_bias_vector",
49
+ "provenance": {
50
+ "source": "onnxruntime/test/contrib_ops/element_wise_ops_test.cc",
51
+ "test": "BiasGeluTest.MLFloat16",
52
+ "notes": "Compact deterministic f16 projection of ORT's bias-vector GELU coverage."
53
+ },
54
+ "inputs": {
55
+ "A": {
56
+ "dtype": "float16",
57
+ "shape": [2, 4],
58
+ "data": { "kind": "values", "values": [-1.5, -0.5, 0.5, 1.5, 2.0, -2.0, 0.25, -0.25] }
59
+ },
60
+ "B": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.25, -0.25, 0.5, -0.5] } }
61
+ },
62
+ "outputs": {
63
+ "C": {
64
+ "dtype": "float16",
65
+ "shape": [2, 4],
66
+ "tolerance": 0.001,
67
+ "data": {
68
+ "kind": "values",
69
+ "values": [-0.1320623, -0.16997046, 0.8413447, 0.8413447, 2.222495, -0.02750498, 0.58002954, -0.16997046]
70
+ }
71
+ }
72
+ }
73
+ },
74
+ {
75
+ "name": "ort_float32_rank3_bias_vector",
76
+ "provenance": {
77
+ "source": "onnxruntime/test/contrib_ops/element_wise_ops_test.cc",
78
+ "test": "BiasGeluTest.Float",
79
+ "notes": "Small deterministic rank-3 instance of ORT's bias-vector GELU coverage."
80
+ },
81
+ "inputs": {
82
+ "A": {
83
+ "dtype": "float32",
84
+ "shape": [2, 2, 3],
85
+ "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, 2.0, -2.0, 0.5, -0.5, 0.25, -0.25, 3.0, -3.0, 0.0] }
86
+ },
87
+ "B": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, -0.5, 0.25] } }
88
+ },
89
+ "outputs": {
90
+ "C": {
91
+ "dtype": "float32",
92
+ "shape": [2, 2, 3],
93
+ "tolerance": 0.000001,
94
+ "data": {
95
+ "kind": "values",
96
+ "values": [-0.15426877, -0.15426877, 1.1179378, 2.4844761, -0.015524104, 0.58002949, 0.0, -0.10032342, 0.0, 3.4991858, -0.00081422925, 0.14967658]
97
+ }
98
+ }
99
+ }
100
+ },
101
+ {
102
+ "name": "ort_float32_rank4_bias_vector",
103
+ "provenance": {
104
+ "source": "onnxruntime/test/contrib_ops/element_wise_ops_test.cc",
105
+ "test": "BiasGeluTest.Float",
106
+ "notes": "Small deterministic rank-4 instance of ORT's bias-vector GELU coverage."
107
+ },
108
+ "inputs": {
109
+ "A": {
110
+ "dtype": "float32",
111
+ "shape": [1, 2, 2, 3],
112
+ "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, 2.0, -2.0, 0.5, -0.5, 0.25, -0.25, 3.0, -3.0, 0.0] }
113
+ },
114
+ "B": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.25, -0.25, 0.5] } }
115
+ },
116
+ "outputs": {
117
+ "C": {
118
+ "dtype": "float32",
119
+ "shape": [1, 2, 2, 3],
120
+ "tolerance": 0.000001,
121
+ "data": {
122
+ "kind": "values",
123
+ "values": [-0.16997051, -0.10032342, 1.3997892, 2.2224948, -0.027505063, 0.8413447, -0.10032342, 0.0, 0.14967658, 3.2481246, -0.0018753314, 0.34573123]
124
+ }
125
+ }
126
+ }
127
+ },
128
+ {
129
+ "name": "ort_float32_empty_middle_dim",
130
+ "provenance": {
131
+ "source": "onnxruntime/test/contrib_ops/element_wise_ops_test.cc",
132
+ "test": "BiasGeluTest.Float",
133
+ "notes": "Additional edge: empty tensors should preserve shape and produce no values."
134
+ },
135
+ "inputs": {
136
+ "A": { "dtype": "float32", "shape": [2, 0, 4], "data": { "kind": "values", "values": [] } },
137
+ "B": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, -0.25, 0.5, -0.5] } }
138
+ },
139
+ "outputs": { "C": { "dtype": "float32", "shape": [2, 0, 4], "data": { "kind": "values", "values": [] } } }
140
+ },
141
+ {
142
+ "name": "empty_zero_hidden",
143
+ "provenance": {
144
+ "notes": "Zero-length last (bias-broadcast) axis: the bias vector itself is empty, so no kernel may size a binding from the hidden extent."
145
+ },
146
+ "inputs": {
147
+ "A": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
148
+ "B": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
149
+ },
150
+ "outputs": { "C": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } } }
151
+ },
152
+ {
153
+ "name": "f32_subnormal_linear_region_zero_bias_vec4_gpu_gap",
154
+ "skipGpu": {
155
+ "category": "permanent",
156
+ "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
157
+ },
158
+ "provenance": {
159
+ "source": "onnxruntime/test/contrib_ops/element_wise_ops_test.cc",
160
+ "test": "BiasGeluTest.Float",
161
+ "notes": "BiasGelu with zero bias reduces to exact erf-form GELU; finite subnormal inputs should produce the x/2 linear-region tail through the vec4 bias path."
162
+ },
163
+ "inputs": {
164
+ "A": {
165
+ "dtype": "float32",
166
+ "shape": [1, 4],
167
+ "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38, -1e-38] }
168
+ },
169
+ "B": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
170
+ },
171
+ "outputs": {
172
+ "C": {
173
+ "dtype": "float32",
174
+ "shape": [1, 4],
175
+ "tolerance": 2e-45,
176
+ "data": {
177
+ "kind": "values",
178
+ "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39, -4.999999675228202e-39]
179
+ }
180
+ }
181
+ }
182
+ },
183
+ {
184
+ "name": "f32_subnormal_linear_region_zero_bias_scalar_gpu_gap",
185
+ "skipGpu": {
186
+ "category": "permanent",
187
+ "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
188
+ },
189
+ "provenance": {
190
+ "source": "onnxruntime/test/contrib_ops/element_wise_ops_test.cc",
191
+ "test": "BiasGeluTest.Float",
192
+ "notes": "Scalar-path companion for BiasGelu subnormal linear-region behavior with zero bias."
193
+ },
194
+ "inputs": {
195
+ "A": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } },
196
+ "B": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
197
+ },
198
+ "outputs": {
199
+ "C": {
200
+ "dtype": "float32",
201
+ "shape": [1, 3],
202
+ "tolerance": 2e-45,
203
+ "data": { "kind": "values", "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39] }
204
+ }
205
+ }
206
+ },
207
+ {
208
+ "name": "scalar_dispatch_cliff_2d_fold",
209
+ "inputs": {
210
+ "A": { "dtype": "float32", "shape": [8388609, 3], "data": { "kind": "linspace", "start": -2.0, "end": 2.0 } },
211
+ "B": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.1, -0.1, 0.3] } }
212
+ },
213
+ "outputs": { "C": { "dtype": "float32", "shape": [8388609, 3], "tolerance": 0.0001 } }
214
+ },
215
+ {
216
+ "name": "scalar_rank1_odd_hidden_bias_broadcast",
217
+ "inputs": {
218
+ "A": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.5, -1.5, 0.7] } },
219
+ "B": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.2, -0.2, 0.3] } }
220
+ },
221
+ "outputs": {
222
+ "C": {
223
+ "dtype": "float32",
224
+ "shape": [3],
225
+ "tolerance": 0.000001,
226
+ "data": { "kind": "values", "values": [1.62423876, -0.07576124, 0.84134474] }
227
+ }
228
+ }
229
+ },
230
+ {
231
+ "name": "scalar_numel_div4_but_last_dim_not",
232
+ "inputs": {
233
+ "A": {
234
+ "dtype": "float32",
235
+ "shape": [4, 3],
236
+ "data": { "kind": "values", "values": [-1.0, 0.5, 2.0, 1.5, -0.5, -1.5, 0.0, 1.0, -2.0, 0.25, -0.25, 0.75] }
237
+ },
238
+ "B": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.2, -0.3, 0.1] } }
239
+ },
240
+ "outputs": { "C": { "dtype": "float32", "shape": [4, 3], "tolerance": 0.000001 } }
241
+ }
242
+ ]
243
+ }