Hugging Face's logo

Kernels:
webgpu-kernels
/
com.microsoft.Gelu

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
9ca4c42
·
verified ·
1 Parent(s): bbbd4e6

sync 2e7068faf55e

Browse files
Files changed (6) hide show
  1. README.md +56 -0
  2. build/webgpu/bench.json +66 -0
  3. build/webgpu/elementwise-bias-gelu.wgsl.jinja +90 -0
  4. build/webgpu/manifest.json +102 -0
  5. build/webgpu/metadata.json +18 -0
  6. build/webgpu/test.json +204 -0
README.md CHANGED
@@ -1,3 +1,59 @@
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.Gelu
10
+
11
+ `com.microsoft` · ONNX Runtime contrib operator · contrib since_version 1
12
+
13
+ ## Description
14
+
15
+ Applies the Gaussian Error Linear Unit (GELU) activation elementwise: `Y = 0.5 * X * (1 + erf(X / sqrt(2)))`. The output has the same shape as the input. Float16 and float32 are supported; the schema's double and bfloat16 types are not.
16
+
17
+ See the [ONNX Runtime `Gelu` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.Gelu) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `X` | `X` | `T` | — | — | Values transformed elementwise by the exact GELU activation. | required |
24
+
25
+ ## Outputs
26
+
27
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
28
+ | --- | --- | --- | --- | --- | --- | --- |
29
+ | `Y` | `Y` | `T` | same as `X` | same as `X` | Output tensor after applying GELU; same shape as the input. | required |
30
+
31
+ ## Type constraints
32
+
33
+ | Variable | Allowed dtypes |
34
+ | --- | --- |
35
+ | `T` | `float32`, `float16` |
36
+
37
+ ## Files
38
+
39
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
40
+ - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
41
+ - [`test.json`](build/webgpu/test.json) — correctness cases
42
+ - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
43
+ - [`elementwise-bias-gelu.wgsl.jinja`](build/webgpu/elementwise-bias-gelu.wgsl.jinja)
44
+
45
+ ## Use with `@huggingface/kernels`
46
+
47
+ The loader derives every required output's shape and logical dtype from the manifest contract and this call.
48
+ It then allocates the result tensors automatically.
49
+
50
+ The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
51
+
52
+ Replace each `*Data` placeholder with a typed array containing the corresponding input data.
53
+
54
+ ```js
55
+ import { getKernel } from "@huggingface/kernels";
56
+
57
+ const kernel = await getKernel("webgpu-kernels/com.microsoft.Gelu", { version: 1 });
58
+ const { Y } = await kernel({ X: { data: XData, shape: [] } });
59
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,66 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.Gelu",
3
+ "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
+ "cases": [
5
+ {
6
+ "name": "gelu-f32-1m",
7
+ "preset": "smoke",
8
+ "vars": { "dtype": "float32", "count": 1048576 },
9
+ "inputs": { "X": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 650, "scale": 0.5 } },
10
+ "outputs": { "Y": { "shape": [1048576], "dtype": "float32" } },
11
+ "bench": {
12
+ "primary": true,
13
+ "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }]
14
+ }
15
+ },
16
+ {
17
+ "name": "gelu-f32-8m",
18
+ "preset": "smoke",
19
+ "vars": { "dtype": "float32", "count": 8388608 },
20
+ "inputs": { "X": { "shape": [8388608], "dtype": "float32", "seed": 7001, "dist": "normal", "scale": 0.5 } },
21
+ "outputs": { "Y": { "shape": [8388608], "dtype": "float32" } },
22
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
23
+ },
24
+ {
25
+ "name": "gelu-f16-8m",
26
+ "preset": "smoke",
27
+ "vars": { "dtype": "float16", "count": 8388608 },
28
+ "inputs": { "X": { "shape": [8388608], "dtype": "float16", "seed": 7002, "dist": "normal", "scale": 0.5 } },
29
+ "outputs": { "Y": { "shape": [8388608], "dtype": "float16" } },
30
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
31
+ },
32
+ {
33
+ "name": "gelu-bert-base-b8-s512-i3072",
34
+ "preset": "model",
35
+ "provenance": {
36
+ "notes": "BERT class defaults (intermediate_size 3072): one FFN activation tensor for a batch of 8 full-length sequences."
37
+ },
38
+ "vars": { "dtype": "float32", "count": 12582912 },
39
+ "inputs": { "X": { "shape": [12582912], "dtype": "float32", "dist": "normal", "seed": 5900, "scale": 0.5 } },
40
+ "outputs": { "Y": { "shape": [12582912], "dtype": "float32" } },
41
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
42
+ },
43
+ {
44
+ "name": "gelu-whisper-encoder-s1500-i1536",
45
+ "preset": "model",
46
+ "provenance": {
47
+ "notes": "Whisper class defaults (encoder_ffn_dim 1536, max_source_positions 1500): one encoder FFN activation for a full 30-second window."
48
+ },
49
+ "vars": { "dtype": "float32", "count": 2304000 },
50
+ "inputs": { "X": { "shape": [2304000], "dtype": "float32", "dist": "normal", "seed": 6000, "scale": 0.5 } },
51
+ "outputs": { "Y": { "shape": [2304000], "dtype": "float32" } },
52
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
53
+ },
54
+ {
55
+ "name": "gelu-clip-text-b64-s77-i2048",
56
+ "preset": "model",
57
+ "provenance": {
58
+ "notes": "CLIP text class defaults (intermediate_size 2048, max_position_embeddings 77) at a 64-caption batch."
59
+ },
60
+ "vars": { "dtype": "float32", "count": 10092544 },
61
+ "inputs": { "X": { "shape": [10092544], "dtype": "float32", "dist": "normal", "seed": 6100, "scale": 0.5 } },
62
+ "outputs": { "Y": { "shape": [10092544], "dtype": "float32" } },
63
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
64
+ }
65
+ ]
66
+ }
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,102 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "com.microsoft",
3
+ "name": "Gelu",
4
+ "sinceVersion": 1,
5
+ "description": "Applies the Gaussian Error Linear Unit (GELU) activation elementwise: `Y = 0.5 * X * (1 + erf(X / sqrt(2)))`. The output has the same shape as the input. Float16 and float32 are supported; the schema's double and bfloat16 types are not.",
6
+ "inputs": [
7
+ { "role": "X", "dtype": "T", "description": "Values transformed elementwise by the exact GELU activation." }
8
+ ],
9
+ "outputs": [
10
+ {
11
+ "role": "Y",
12
+ "dtype": "T",
13
+ "rank": "ranks.X",
14
+ "shape": "shapes.X",
15
+ "description": "Output tensor after applying GELU; same shape as the input."
16
+ }
17
+ ],
18
+ "typeConstraints": { "T": ["float32", "float16"] },
19
+ "args": {
20
+ "X": { "kind": "tensor", "semantic": "X", "role": "input" },
21
+ "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
22
+ },
23
+ "tunables": { "WORKGROUP_SIZE": 256 },
24
+ "constants": {
25
+ "scalar": "dtypes.T",
26
+ "usesF16": "dtypes.T == \"f16\"",
27
+ "approximate": "\"erf\"",
28
+ "vec4Tail": false,
29
+ "hasBias": false
30
+ },
31
+ "variants": [
32
+ {
33
+ "id": "vec4",
34
+ "priority": 20,
35
+ "when": ["numel(shapes.X) == numel(shapes.Y)", "f16Ok(dtypes.T)", "numel(shapes.X) > 0", "numel(shapes.X) % 4 == 0"],
36
+ "constants": { "vec4": true, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
37
+ "passes": [
38
+ {
39
+ "id": "main",
40
+ "name": "Gelu.vec4",
41
+ "shader": "elementwise-bias-gelu.wgsl.jinja",
42
+ "bindings": [
43
+ {
44
+ "name": "x",
45
+ "arg": "X",
46
+ "semantic": "X",
47
+ "buffer": { "type": "read-only-storage" },
48
+ "elementType": "$vectorScalar"
49
+ },
50
+ {
51
+ "name": "y",
52
+ "arg": "Y",
53
+ "semantic": "Y",
54
+ "buffer": { "type": "storage" },
55
+ "elementType": "$vectorScalar"
56
+ },
57
+ {
58
+ "name": "params",
59
+ "semantic": "kernel.params",
60
+ "buffer": { "type": "uniform" },
61
+ "struct": {
62
+ "name": "Params",
63
+ "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }]
64
+ }
65
+ }
66
+ ],
67
+ "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
68
+ }
69
+ ]
70
+ },
71
+ {
72
+ "id": "scalar",
73
+ "priority": 0,
74
+ "when": ["numel(shapes.X) == numel(shapes.Y)", "f16Ok(dtypes.T)"],
75
+ "constants": { "vec4": false },
76
+ "passes": [
77
+ {
78
+ "id": "main",
79
+ "name": "Gelu.scalar",
80
+ "shader": "elementwise-bias-gelu.wgsl.jinja",
81
+ "bindings": [
82
+ {
83
+ "name": "x",
84
+ "arg": "X",
85
+ "semantic": "X",
86
+ "buffer": { "type": "read-only-storage" },
87
+ "elementType": "$scalar"
88
+ },
89
+ { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
90
+ {
91
+ "name": "params",
92
+ "semantic": "kernel.params",
93
+ "buffer": { "type": "uniform" },
94
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
95
+ }
96
+ ],
97
+ "dispatch": { "threads": "numel(shapes.X)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
98
+ }
99
+ ]
100
+ }
101
+ ]
102
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,18 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "com.microsoft.Gelu",
3
+ "id": "_com_microsoft_gelu_webgpu_6511522",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "cFEvXUyEhaDZGPv9ppB2H/jtirbnrkEOd8ngzRycX+k=",
11
+ "elementwise-bias-gelu.wgsl.jinja": "Eg8N2jJCMce+IsYNcCzuxvs8CSv7yTQXdorvv+c0L58=",
12
+ "manifest.json": "89wktvMtLewo412huOWKAmpUsL1Q9k9RYkeWTxrhepE=",
13
+ "test.json": "CmfJcfaIcytJIPtgyCKKLh72PzJvUlrpycr/FB0EZr8="
14
+ }
15
+ },
16
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
17
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.Gelu" }
18
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,204 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "com.microsoft.Gelu",
3
+ "cases": [
4
+ {
5
+ "name": "dispatch_cliff_scalar_f32",
6
+ "inputs": {
7
+ "X": { "dtype": "float32", "shape": [16776961], "data": { "kind": "linspace", "start": -2.0, "end": 2.0 } }
8
+ },
9
+ "outputs": { "Y": { "dtype": "float32", "shape": [16776961], "tolerance": 0.0001 } }
10
+ },
11
+ {
12
+ "name": "ort_float32_erf_extreme_edges",
13
+ "provenance": {
14
+ "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
15
+ "test": "ActivationOpTest.Gelu",
16
+ "notes": "Exact erf-form GELU with large finite values that should saturate to zero/pass-through without overflowing."
17
+ },
18
+ "inputs": {
19
+ "X": {
20
+ "dtype": "float32",
21
+ "shape": [1, 1, 9],
22
+ "data": { "kind": "values", "values": [-1000.0, -100.0, -10.0, -1.0, 0.0, 1.0, 10.0, 100.0, 1000.0] }
23
+ }
24
+ },
25
+ "outputs": {
26
+ "Y": {
27
+ "dtype": "float32",
28
+ "shape": [1, 1, 9],
29
+ "tolerance": 0.000001,
30
+ "data": {
31
+ "kind": "values",
32
+ "values": [0.0, 0.0, 0.0, -0.15865525603294373, 0.0, 0.8413447141647339, 10.0, 100.0, 1000.0]
33
+ }
34
+ }
35
+ }
36
+ },
37
+ {
38
+ "name": "ort_float32_nonfinite_edges",
39
+ "provenance": {
40
+ "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
41
+ "test": "ActivationOpTest.Gelu_bfloat16",
42
+ "notes": "Float32 counterpart of ORT's nonfinite activation coverage; -Infinity produces NaN under the exact erf expression."
43
+ },
44
+ "inputs": {
45
+ "X": {
46
+ "dtype": "float32",
47
+ "shape": [5],
48
+ "data": { "kind": "values", "values": ["-Infinity", "Infinity", "NaN", 0.0, 0.0] }
49
+ }
50
+ },
51
+ "outputs": {
52
+ "Y": {
53
+ "dtype": "float32",
54
+ "shape": [5],
55
+ "allowNaN": true,
56
+ "data": { "kind": "values", "values": ["NaN", "Infinity", "NaN", 0.0, 0.0] }
57
+ }
58
+ }
59
+ },
60
+ {
61
+ "name": "ort_float32_empty_rank3",
62
+ "provenance": {
63
+ "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
64
+ "notes": "Empty tensors should preserve shape and produce no values."
65
+ },
66
+ "inputs": { "X": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } },
67
+ "outputs": { "Y": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } }
68
+ },
69
+ {
70
+ "name": "f32_subnormal_linear_region_vec4_gpu_gap",
71
+ "skipGpu": {
72
+ "category": "permanent",
73
+ "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."
74
+ },
75
+ "provenance": {
76
+ "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
77
+ "test": "ActivationOpTest.Gelu",
78
+ "notes": "Near zero, exact erf-form GELU is approximately x/2; the com.microsoft path should preserve finite subnormal outputs."
79
+ },
80
+ "inputs": {
81
+ "X": {
82
+ "dtype": "float32",
83
+ "shape": [4],
84
+ "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38, -1e-38] }
85
+ }
86
+ },
87
+ "outputs": {
88
+ "Y": {
89
+ "dtype": "float32",
90
+ "shape": [4],
91
+ "tolerance": 2e-45,
92
+ "data": {
93
+ "kind": "values",
94
+ "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39, -4.999999675228202e-39]
95
+ }
96
+ }
97
+ }
98
+ },
99
+ {
100
+ "name": "f32_subnormal_linear_region_scalar_gpu_gap",
101
+ "skipGpu": {
102
+ "category": "permanent",
103
+ "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."
104
+ },
105
+ "provenance": {
106
+ "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
107
+ "test": "ActivationOpTest.Gelu",
108
+ "notes": "Scalar-path companion for com.microsoft.Gelu subnormal linear-region behavior."
109
+ },
110
+ "inputs": {
111
+ "X": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } }
112
+ },
113
+ "outputs": {
114
+ "Y": {
115
+ "dtype": "float32",
116
+ "shape": [3],
117
+ "tolerance": 2e-45,
118
+ "data": { "kind": "values", "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39] }
119
+ }
120
+ }
121
+ },
122
+ {
123
+ "name": "rank0_negative_scalar",
124
+ "provenance": {
125
+ "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
126
+ "test": "ActivationOpTest.Gelu",
127
+ "notes": "Additional edge: scalar tensors use the same exact erf-form GELU path."
128
+ },
129
+ "inputs": { "X": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-0.5] } } },
130
+ "outputs": {
131
+ "Y": {
132
+ "dtype": "float32",
133
+ "shape": [],
134
+ "tolerance": 0.000001,
135
+ "data": { "kind": "values", "values": [-0.15426877] }
136
+ }
137
+ }
138
+ },
139
+ {
140
+ "name": "vec4_f32_4x8",
141
+ "inputs": {
142
+ "X": {
143
+ "dtype": "float32",
144
+ "shape": [4, 8],
145
+ "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 2.0 }
146
+ }
147
+ },
148
+ "outputs": { "Y": { "dtype": "float32", "shape": [4, 8], "tolerance": 0.00001 } }
149
+ },
150
+ {
151
+ "name": "vec4_f16_4x8",
152
+ "inputs": {
153
+ "X": {
154
+ "dtype": "float16",
155
+ "shape": [4, 8],
156
+ "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 2.0 }
157
+ }
158
+ },
159
+ "outputs": { "Y": { "dtype": "float16", "shape": [4, 8], "tolerance": 0.01 } }
160
+ },
161
+ {
162
+ "name": "scalar_numel_not_div4_inline",
163
+ "inputs": {
164
+ "X": {
165
+ "dtype": "float32",
166
+ "shape": [6],
167
+ "data": { "kind": "values", "values": [1.5, -1.5, 0.7, -0.7, 0.3, -0.3] }
168
+ }
169
+ },
170
+ "outputs": {
171
+ "Y": {
172
+ "dtype": "float32",
173
+ "shape": [6],
174
+ "tolerance": 0.000001,
175
+ "data": {
176
+ "kind": "values",
177
+ "values": [1.39978915, -0.10021085, 0.53062549, -0.16937451, 0.18537341, -0.11462659]
178
+ }
179
+ }
180
+ }
181
+ },
182
+ {
183
+ "name": "vec4_odd_last_dim_still_routes_vec4",
184
+ "inputs": {
185
+ "X": {
186
+ "dtype": "float32",
187
+ "shape": [6, 2],
188
+ "data": { "kind": "values", "values": [-1.0, 0.5, 1.0, -0.5, 2.0, -2.0, 0.0, 1.5, -1.5, 0.25, -0.25, 0.75] }
189
+ }
190
+ },
191
+ "outputs": {
192
+ "Y": {
193
+ "dtype": "float32",
194
+ "shape": [6, 2],
195
+ "tolerance": 0.000001,
196
+ "data": {
197
+ "kind": "values",
198
+ "values": [-0.15865527, 0.34573123, 0.84134471, -0.15426877, 1.95449984, -0.04550013, 0.0, 1.39978909, -0.10021085, 0.14967656, -0.10032343, 0.58002955]
199
+ }
200
+ }
201
+ }
202
+ }
203
+ ]
204
+ }