Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Atan

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

sync 2e7068faf55e

Browse files
Files changed (7) hide show
  1. README.md +57 -0
  2. build/webgpu/bench.json +49 -0
  3. build/webgpu/manifest.json +78 -0
  4. build/webgpu/metadata.json +19 -0
  5. build/webgpu/test.json +210 -0
  6. build/webgpu/unary-scalar.wgsl.jinja +84 -0
  7. build/webgpu/unary-vec4.wgsl.jinja +77 -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
+ # ai.onnx.Atan
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 7
12
+
13
+ ## Description
14
+
15
+ Computes the arctangent (`atan`) of each element of the input tensor. The output has the same shape as the input.
16
+
17
+ See the [ONNX `Atan` spec](https://onnx.ai/onnx/operators/onnx__Atan.html) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `input` | `x` | `T` | — | — | Values to map elementwise to arctangents in radians. | required |
24
+
25
+ ## Outputs
26
+
27
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
28
+ | --- | --- | --- | --- | --- | --- | --- |
29
+ | `output` | `y` | `T` | same as `input` | same as `input` | The arctangent of each element of `input`, computed elementwise. | 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
+ - [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
44
+ - [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.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/ai.onnx.Atan", { version: 1 });
59
+ const { y } = await kernel({ x: { data: xData, shape: [3] } });
60
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,49 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.Atan",
3
+ "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
+ "cases": [
5
+ {
6
+ "name": "atan-f32-1m-vec4",
7
+ "preset": "smoke",
8
+ "vars": { "dtype": "float32", "count": 1048576 },
9
+ "inputs": { "x": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 820, "scale": 4 } },
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": "atan-f32-scalar-non-mult4",
18
+ "preset": "smoke",
19
+ "vars": { "dtype": "float32", "count": 1048575 },
20
+ "inputs": { "x": { "shape": [1048575], "dtype": "float32", "dist": "normal", "seed": 821, "scale": 4 } },
21
+ "outputs": { "y": { "shape": [1048575], "dtype": "float32" } },
22
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
23
+ },
24
+ {
25
+ "name": "atan-f16-8m-vec4",
26
+ "preset": "smoke",
27
+ "vars": { "dtype": "float16", "count": 8388608 },
28
+ "inputs": { "x": { "shape": [8388608], "dtype": "float16", "dist": "normal", "seed": 822, "scale": 4 } },
29
+ "outputs": { "y": { "shape": [8388608], "dtype": "float16" } },
30
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
31
+ },
32
+ {
33
+ "name": "atan-f16-scalar-non-mult4",
34
+ "preset": "smoke",
35
+ "vars": { "dtype": "float16", "count": 8388607 },
36
+ "inputs": { "x": { "shape": [8388607], "dtype": "float16", "dist": "normal", "seed": 823, "scale": 4 } },
37
+ "outputs": { "y": { "shape": [8388607], "dtype": "float16" } },
38
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
39
+ },
40
+ {
41
+ "name": "atan-f32-8m-dispatch",
42
+ "preset": "smoke",
43
+ "vars": { "dtype": "float32", "count": 8388608 },
44
+ "inputs": { "x": { "shape": [8388608], "dtype": "float32", "dist": "normal", "seed": 824, "scale": 4 } },
45
+ "outputs": { "y": { "shape": [8388608], "dtype": "float32" } },
46
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
47
+ }
48
+ ]
49
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,78 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "ai.onnx",
3
+ "name": "Atan",
4
+ "sinceVersion": 7,
5
+ "description": "Computes the arctangent (`atan`) of each element of the input tensor. The output has the same shape as the input.",
6
+ "inputs": [{ "role": "input", "dtype": "T", "description": "Values to map elementwise to arctangents in radians." }],
7
+ "outputs": [
8
+ {
9
+ "role": "output",
10
+ "dtype": "T",
11
+ "rank": "ranks.input",
12
+ "description": "The arctangent of each element of `input`, computed elementwise.",
13
+ "shape": "shapes.input"
14
+ }
15
+ ],
16
+ "typeConstraints": { "T": ["float32", "float16"] },
17
+ "args": {
18
+ "x": { "kind": "tensor", "semantic": "input", "role": "input" },
19
+ "y": { "kind": "tensor", "semantic": "output", "role": "output" }
20
+ },
21
+ "tunables": { "WORKGROUP_SIZE": 256 },
22
+ "variants": [
23
+ {
24
+ "id": "same_layout_vec4",
25
+ "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
26
+ "constants": {
27
+ "scalar": "dtypes.T",
28
+ "usesF16": "dtypes.T == \"f16\"",
29
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
30
+ },
31
+ "passes": [
32
+ {
33
+ "id": "main",
34
+ "name": "Atan.vec4",
35
+ "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"atan\"" } },
36
+ "bindings": [
37
+ { "name": "x", "arg": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$vectorScalar" },
38
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
39
+ {
40
+ "name": "params",
41
+ "semantic": "kernel.params",
42
+ "buffer": { "type": "uniform" },
43
+ "struct": {
44
+ "name": "Params",
45
+ "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }]
46
+ }
47
+ }
48
+ ],
49
+ "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
50
+ }
51
+ ],
52
+ "priority": 20
53
+ },
54
+ {
55
+ "id": "elementwise",
56
+ "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
57
+ "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
58
+ "passes": [
59
+ {
60
+ "id": "main",
61
+ "name": "Atan",
62
+ "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"atan\"", "itemsPerInvocation": 4 } },
63
+ "bindings": [
64
+ { "name": "x", "arg": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
65
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
66
+ {
67
+ "name": "params",
68
+ "semantic": "kernel.params",
69
+ "buffer": { "type": "uniform" },
70
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
71
+ }
72
+ ],
73
+ "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
74
+ }
75
+ ]
76
+ }
77
+ ]
78
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,19 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "ai.onnx.Atan",
3
+ "id": "_ai_onnx_atan_webgpu_4f57e81",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "TItmPVJYgjxju1u//NYkrMhw8HFNoUkrft0yYWmY7fk=",
11
+ "manifest.json": "cL4O0izwr8OqCuDTzUuCZn6MCEEwvEZ+AxKLoGuT9Io=",
12
+ "test.json": "sJLtryPr/wpc6udFIJVVoS8w/NBZUXL7PYicdInn/0s=",
13
+ "unary-scalar.wgsl.jinja": "MXAwWOBnc0HvmEQ+bWF3AaYkaXDARTYiMU9urQZNb/w=",
14
+ "unary-vec4.wgsl.jinja": "LWSY/fzvxogBMUS0MYUut2CpXJpivEP/OlZNcrxr8jQ="
15
+ }
16
+ },
17
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
18
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Atan" }
19
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,210 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.Atan",
3
+ "cases": [
4
+ {
5
+ "name": "f32_values",
6
+ "inputs": {
7
+ "x": {
8
+ "dtype": "float32",
9
+ "shape": [6],
10
+ "data": { "kind": "values", "values": [-4.0, -1.0, 0.0, 0.5, 1.0, 4.0] }
11
+ }
12
+ },
13
+ "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.000001 } }
14
+ },
15
+ {
16
+ "name": "f32_subnormal_identity_tail_gpu_gap",
17
+ "skipGpu": {
18
+ "category": "permanent",
19
+ "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."
20
+ },
21
+ "provenance": {
22
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
23
+ "test": "MathOpTest.Atan",
24
+ "notes": "For tiny finite inputs atan(x) rounds back to x in float32; zero-flushing erases the signed tail."
25
+ },
26
+ "inputs": {
27
+ "x": {
28
+ "dtype": "float32",
29
+ "shape": [4],
30
+ "data": { "kind": "values", "values": [-1e-39, -1e-40, 1e-40, 1e-39] }
31
+ }
32
+ },
33
+ "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
34
+ },
35
+ {
36
+ "name": "f32_subnormal_identity_tail_scalar_gpu_gap",
37
+ "skipGpu": {
38
+ "category": "permanent",
39
+ "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."
40
+ },
41
+ "provenance": {
42
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
43
+ "test": "MathOpTest.Atan",
44
+ "notes": "Scalar-path companion: subnormal inputs are valid finite Atan outputs."
45
+ },
46
+ "inputs": {
47
+ "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } }
48
+ },
49
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
50
+ },
51
+ {
52
+ "name": "f16_values",
53
+ "inputs": {
54
+ "x": {
55
+ "dtype": "float16",
56
+ "shape": [2, 3],
57
+ "data": { "kind": "values", "values": [-4.0, -1.0, 0.0, 0.5, 1.0, 4.0] }
58
+ }
59
+ },
60
+ "outputs": { "y": { "dtype": "float16", "shape": [2, 3] } },
61
+ "tolerance": 0.002
62
+ },
63
+ {
64
+ "name": "special_values_infinity_nan",
65
+ "inputs": {
66
+ "x": {
67
+ "dtype": "float32",
68
+ "shape": [5],
69
+ "data": { "kind": "values", "values": ["-Infinity", -1.0, 0.0, "Infinity", "NaN"] }
70
+ }
71
+ },
72
+ "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0.000001, "allowNaN": true } }
73
+ },
74
+ {
75
+ "name": "ort_float_wide_range",
76
+ "provenance": {
77
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
78
+ "test": "MathOpTest.Atan"
79
+ },
80
+ "inputs": {
81
+ "x": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [-10.0, -5.0, 0.0, 5.0, 10.0] } }
82
+ },
83
+ "outputs": {
84
+ "y": {
85
+ "dtype": "float32",
86
+ "shape": [5],
87
+ "tolerance": 0.00001,
88
+ "data": {
89
+ "kind": "values",
90
+ "values": [-1.4711276292800903, -1.3734008073806763, 0.0, 1.3734008073806763, 1.4711276292800903]
91
+ }
92
+ }
93
+ }
94
+ },
95
+ {
96
+ "name": "onnx_backend_example",
97
+ "provenance": {
98
+ "source": "cmake/external/onnx/onnx/backend/test/data/node/test_atan_example",
99
+ "test": "test_atan_example"
100
+ },
101
+ "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } } },
102
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
103
+ },
104
+ {
105
+ "name": "onnx_backend_atan",
106
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_atan" },
107
+ "inputs": {
108
+ "x": {
109
+ "dtype": "float32",
110
+ "shape": [3, 4, 5],
111
+ "data": {
112
+ "kind": "values",
113
+ "values": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
114
+ }
115
+ }
116
+ },
117
+ "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0.00001 } }
118
+ },
119
+ {
120
+ "name": "onnx_backend_atan_example",
121
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_atan_example" },
122
+ "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } } },
123
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.00001 } }
124
+ },
125
+ {
126
+ "name": "vec4_f16_lanes",
127
+ "inputs": {
128
+ "x": {
129
+ "dtype": "float16",
130
+ "shape": [16],
131
+ "data": {
132
+ "kind": "values",
133
+ "values": [-8.0, -4.0, -2.0, -1.0, -0.5, -0.25, 0.0, 0.25, 0.5, 1.0, 2.0, 4.0, 8.0, 16.0, 50.0, 100.0]
134
+ }
135
+ }
136
+ },
137
+ "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0.001, "relTolerance": 0.002 } }
138
+ },
139
+ {
140
+ "name": "empty_input_zero_dim",
141
+ "inputs": { "x": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
142
+ "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
143
+ },
144
+ {
145
+ "name": "f16_vec4_realistic_finite",
146
+ "provenance": {
147
+ "notes": "MANDATE A clean coverage: f16 at a multiple-of-4 numel exercises the same_layout_vec4 packed path (atan(vec4<f16>)) over an ordinary finite range a real model would feed (e.g. coordinate/angle decomposition). Existing f16 vec4 case is only 16 elements; this is a more realistic feature-map row and a denser finite sweep."
148
+ },
149
+ "inputs": {
150
+ "x": {
151
+ "dtype": "float16",
152
+ "shape": [64],
153
+ "data": {
154
+ "kind": "values",
155
+ "values": [-6.0, -5.0, -4.0, -3.5, -3.0, -2.5, -2.0, -1.75, -1.5, -1.25, -1.0, -0.875, -0.75, -0.625, -0.5, -0.375, -0.25, -0.125, -0.0625, -0.03125, 0.0, 0.03125, 0.0625, 0.125, 0.25, 0.375, 0.5, 0.625, 0.75, 0.875, 1.0, 1.25, 1.5, 1.75, 2.0, 2.5, 3.0, 3.5, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 12.0, 16.0, 20.0, 24.0, 32.0, -7.0, -8.0, -9.0, -10.0, -12.0, -16.0, -20.0, -24.0, -32.0, 0.1875, -0.1875, 0.4375, -0.4375, 1.125]
156
+ }
157
+ }
158
+ },
159
+ "outputs": { "y": { "dtype": "float16", "shape": [64], "tolerance": 0.001, "relTolerance": 0.002 } }
160
+ },
161
+ {
162
+ "name": "f16_scalar_path_non_multiple_of_4",
163
+ "provenance": {
164
+ "notes": "MANDATE A clean coverage: f16 with numel%4 != 0 (15) fails the same_layout_vec4 when-gate (numel(X) % 4 == 0) and must route through the scalar elementwise variant. Verifies f16 atan on the scalar fallback over a normal finite range; the only other non-%4 f16 case is shape [2,3]=6 with a tiny 6-value sweep."
165
+ },
166
+ "inputs": {
167
+ "x": {
168
+ "dtype": "float16",
169
+ "shape": [15],
170
+ "data": {
171
+ "kind": "values",
172
+ "values": [-3.0, -2.0, -1.5, -1.0, -0.5, -0.25, -0.125, 0.0, 0.125, 0.25, 0.5, 1.0, 1.5, 2.0, 3.0]
173
+ }
174
+ }
175
+ },
176
+ "outputs": { "y": { "dtype": "float16", "shape": [15], "tolerance": 0.001, "relTolerance": 0.002 } }
177
+ },
178
+ {
179
+ "name": "f32_scalar_path_non_multiple_of_4",
180
+ "provenance": {
181
+ "notes": "MANDATE A clean coverage: f32 numel%4 != 0 (7) routes through the scalar elementwise variant (same_layout_vec4 requires numel%4==0). No existing f32 case with non-%4 numel exercises the scalar path on the float32 dtype; ordinary finite range."
182
+ },
183
+ "inputs": {
184
+ "x": {
185
+ "dtype": "float32",
186
+ "shape": [7],
187
+ "data": { "kind": "values", "values": [-2.5, -1.0, -0.25, 0.0, 0.25, 1.0, 2.5] }
188
+ }
189
+ },
190
+ "outputs": { "y": { "dtype": "float32", "shape": [7], "tolerance": 0.000001 } }
191
+ },
192
+ {
193
+ "name": "f16_2d_featuremap_vec4",
194
+ "provenance": {
195
+ "notes": "MANDATE A clean coverage: rank-2 f16 with a flat numel (4x8=32) that is a multiple of 4 -> same_layout_vec4. Realistic small feature-map tile; existing f16 rank-2 case ([2,3]) is non-%4 (scalar). Confirms vec4 packing is correct when the leading dim is >1."
196
+ },
197
+ "inputs": {
198
+ "x": {
199
+ "dtype": "float16",
200
+ "shape": [4, 8],
201
+ "data": {
202
+ "kind": "values",
203
+ "values": [-4.0, -3.0, -2.0, -1.0, -0.5, 0.0, 0.5, 1.0, 2.0, 3.0, 4.0, 5.0, -5.0, -1.5, 1.5, -2.5, 2.5, -0.75, 0.75, -0.125, 0.125, -8.0, 8.0, -10.0, 10.0, -0.0625, 0.0625, 6.0, -6.0, 7.0, -7.0, 0.25]
204
+ }
205
+ }
206
+ },
207
+ "outputs": { "y": { "dtype": "float16", "shape": [4, 8], "tolerance": 0.001, "relTolerance": 0.002 } }
208
+ }
209
+ ]
210
+ }
build/webgpu/unary-scalar.wgsl.jinja ADDED
@@ -0,0 +1,84 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {% macro flat_tail_open() %}
2
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
3
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
4
+ // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
5
+ // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
6
+ let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
7
+ // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
8
+ // the logical tensor length (or its storage binding) to be vec4 aligned.
9
+ let begin = invocation * {{ source.itemsPerInvocation }}u;
10
+ let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
11
+ for (var i = begin; i < end; i = i + 1u) {
12
+ {%- endmacro %}
13
+ {% macro flat_tail_close() %}
14
+ }
15
+ {% endmacro %}
16
+
17
+ // Scalar unary fallback. Each branch retains the operation's numeric hardening,
18
+ // including Payne-Hanek trigonometric range reduction and NaN/overflow guards.
19
+ {% if usesF16 %}
20
+ enable f16;
21
+ {% endif %}
22
+ {{ env.wgsl.resourceDeclarations }}
23
+ // Backend-stable f32 atan/asin/acos helpers. The atan core range-reduces to
24
+ // [-tan(pi/8), tan(pi/8)] before a degree-19 odd series. Asin/acos use stable
25
+ // half-angle identities, avoiding cancellation in 1-x*x near the domain ends.
26
+ // Maximum absolute error over finite f32 inputs is below 3e-7.
27
+ const PRECISE_INV_TRIG_HALF_PI: f32 = 1.5707963267948966;
28
+ const PRECISE_INV_TRIG_QUARTER_PI: f32 = 0.7853981633974483;
29
+ const PRECISE_INV_TRIG_HALF_PI_HI: f32 = 1.570796251296997;
30
+ const PRECISE_INV_TRIG_HALF_PI_LO: f32 = 7.549789415861596e-8;
31
+
32
+ fn precise_atan_unit(z: f32) -> f32 {
33
+ let z2 = z * z;
34
+ var polynomial = -0.05263157894736842;
35
+ polynomial = fma(polynomial, z2, 0.058823529411764705);
36
+ polynomial = fma(polynomial, z2, -0.06666666666666667);
37
+ polynomial = fma(polynomial, z2, 0.07692307692307693);
38
+ polynomial = fma(polynomial, z2, -0.09090909090909091);
39
+ polynomial = fma(polynomial, z2, 0.1111111111111111);
40
+ polynomial = fma(polynomial, z2, -0.14285714285714285);
41
+ polynomial = fma(polynomial, z2, 0.2);
42
+ polynomial = fma(polynomial, z2, -0.3333333333333333);
43
+ return z * fma(polynomial, z2, 1.0);
44
+ }
45
+
46
+ fn precise_atan_positive(value: f32) -> f32 {
47
+ if (value > 2.414213562373095) {
48
+ let tail = fma(-1.0, precise_atan_unit(1.0 / value),
49
+ PRECISE_INV_TRIG_HALF_PI_LO);
50
+ // The bitcast round-trip is a rounding barrier. It prevents reassociation
51
+ // of HI + (LO - x) into PI - x, which loses the low word near the domain edge.
52
+ let stableTail = bitcast<f32>(bitcast<u32>(tail));
53
+ // Spell the final high/low combination as one fused operation to prevent
54
+ // reassociation into rounded(pi/2) - x and select the result with f32 FMA rounding.
55
+ return fma(1.0, stableTail, PRECISE_INV_TRIG_HALF_PI_HI);
56
+ }
57
+ if (value > 0.41421356237309503) {
58
+ return PRECISE_INV_TRIG_QUARTER_PI
59
+ + precise_atan_unit((value - 1.0) / (value + 1.0));
60
+ }
61
+ return precise_atan_unit(value);
62
+ }
63
+ fn atan_accurate(value: f32) -> f32 {
64
+ let bits = bitcast<u32>(value);
65
+ let magnitudeBits = bits & 0x7fffffffu;
66
+ if (magnitudeBits > 0x7f800000u) {
67
+ return value;
68
+ }
69
+ if (magnitudeBits == 0x7f800000u) {
70
+ return select(PRECISE_INV_TRIG_HALF_PI, -PRECISE_INV_TRIG_HALF_PI,
71
+ (bits & 0x80000000u) != 0u);
72
+ }
73
+ if (magnitudeBits == 0u) {
74
+ return value;
75
+ }
76
+ let result = precise_atan_positive(abs(value));
77
+ return select(result, -result, (bits & 0x80000000u) != 0u);
78
+ }
79
+
80
+ {{ flat_tail_open() }}
81
+ let v = f32(x[i]);
82
+ y[i] = {{ scalar }}(atan_accurate(v));
83
+ {{ flat_tail_close() -}}
84
+ }
build/webgpu/unary-vec4.wgsl.jinja ADDED
@@ -0,0 +1,77 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ // Loads and stores vec4<T> (128 bits) while retaining scalar per-component
2
+ // arithmetic, including per-component helper calls for guard-heavy operations.
3
+ {% if usesF16 %}
4
+ enable f16;
5
+ {% endif %}
6
+ {{ env.wgsl.resourceDeclarations }}
7
+
8
+ // Backend-stable f32 atan/asin/acos helpers. The atan core range-reduces to
9
+ // [-tan(pi/8), tan(pi/8)] before a degree-19 odd series. Asin/acos use stable
10
+ // half-angle identities, avoiding cancellation in 1-x*x near the domain ends.
11
+ // Maximum absolute error over finite f32 inputs is below 3e-7.
12
+ const PRECISE_INV_TRIG_HALF_PI: f32 = 1.5707963267948966;
13
+ const PRECISE_INV_TRIG_QUARTER_PI: f32 = 0.7853981633974483;
14
+ const PRECISE_INV_TRIG_HALF_PI_HI: f32 = 1.570796251296997;
15
+ const PRECISE_INV_TRIG_HALF_PI_LO: f32 = 7.549789415861596e-8;
16
+
17
+ fn precise_atan_unit(z: f32) -> f32 {
18
+ let z2 = z * z;
19
+ var polynomial = -0.05263157894736842;
20
+ polynomial = fma(polynomial, z2, 0.058823529411764705);
21
+ polynomial = fma(polynomial, z2, -0.06666666666666667);
22
+ polynomial = fma(polynomial, z2, 0.07692307692307693);
23
+ polynomial = fma(polynomial, z2, -0.09090909090909091);
24
+ polynomial = fma(polynomial, z2, 0.1111111111111111);
25
+ polynomial = fma(polynomial, z2, -0.14285714285714285);
26
+ polynomial = fma(polynomial, z2, 0.2);
27
+ polynomial = fma(polynomial, z2, -0.3333333333333333);
28
+ return z * fma(polynomial, z2, 1.0);
29
+ }
30
+
31
+ fn precise_atan_positive(value: f32) -> f32 {
32
+ if (value > 2.414213562373095) {
33
+ let tail = fma(-1.0, precise_atan_unit(1.0 / value),
34
+ PRECISE_INV_TRIG_HALF_PI_LO);
35
+ // The bitcast round-trip is a rounding barrier. It prevents reassociation
36
+ // of HI + (LO - x) into PI - x, which loses the low word near the domain edge.
37
+ let stableTail = bitcast<f32>(bitcast<u32>(tail));
38
+ // Spell the final high/low combination as one fused operation to prevent
39
+ // reassociation into rounded(pi/2) - x and select the result with f32 FMA rounding.
40
+ return fma(1.0, stableTail, PRECISE_INV_TRIG_HALF_PI_HI);
41
+ }
42
+ if (value > 0.41421356237309503) {
43
+ return PRECISE_INV_TRIG_QUARTER_PI
44
+ + precise_atan_unit((value - 1.0) / (value + 1.0));
45
+ }
46
+ return precise_atan_unit(value);
47
+ }
48
+ fn atan_accurate(value: f32) -> f32 {
49
+ let bits = bitcast<u32>(value);
50
+ let magnitudeBits = bits & 0x7fffffffu;
51
+ if (magnitudeBits > 0x7f800000u) {
52
+ return value;
53
+ }
54
+ if (magnitudeBits == 0x7f800000u) {
55
+ return select(PRECISE_INV_TRIG_HALF_PI, -PRECISE_INV_TRIG_HALF_PI,
56
+ (bits & 0x80000000u) != 0u);
57
+ }
58
+ if (magnitudeBits == 0u) {
59
+ return value;
60
+ }
61
+ let result = precise_atan_positive(abs(value));
62
+ return select(result, -result, (bits & 0x80000000u) != 0u);
63
+ }
64
+
65
+
66
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
67
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
68
+ // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
69
+ // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
70
+ let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
71
+ if (i >= params.count) {
72
+ return;
73
+ }
74
+ let xv = x[i];
75
+ let fv = vec4<f32>(xv);
76
+ y[i] = {{ vectorScalar }}(vec4<f32>(atan_accurate(fv.x), atan_accurate(fv.y), atan_accurate(fv.z), atan_accurate(fv.w)));
77
+ }