Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Acos

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
181d18c
·
verified ·
1 Parent(s): 8d1ebb5

sync 2e7068faf55e

Browse files
Files changed (7) hide show
  1. README.md +57 -0
  2. build/webgpu/bench.json +53 -0
  3. build/webgpu/manifest.json +78 -0
  4. build/webgpu/metadata.json +19 -0
  5. build/webgpu/test.json +164 -0
  6. build/webgpu/unary-scalar.wgsl.jinja +96 -0
  7. build/webgpu/unary-vec4.wgsl.jinja +89 -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.Acos
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 7
12
+
13
+ ## Description
14
+
15
+ Computes the arccosine (inverse cosine) of each element of the input tensor. The output has the same shape as the input.
16
+
17
+ See the [ONNX `Acos` spec](https://onnx.ai/onnx/operators/onnx__Acos.html) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `input` | `x` | `T` | — | — | Cosine values to map elementwise to angles 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 arccosine of the input tensor, 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.Acos", { version: 1 });
59
+ const { y } = await kernel({ x: { data: xData, shape: [3] } });
60
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,53 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.Acos",
3
+ "cases": [
4
+ {
5
+ "name": "1m_f32",
6
+ "inputs": { "x": { "dtype": "float32", "shape": [1048576] } },
7
+ "outputs": { "y": { "dtype": "float32", "shape": [1048576] } }
8
+ },
9
+ {
10
+ "name": "acos-f32-1m-vec4",
11
+ "preset": "smoke",
12
+ "vars": { "dtype": "float32", "count": 1048576 },
13
+ "inputs": {
14
+ "x": { "shape": [1048576], "dtype": "float32", "dist": "uniform", "seed": 9101, "scale": 2, "offset": -1 }
15
+ },
16
+ "outputs": { "y": { "shape": [1048576], "dtype": "float32" } },
17
+ "bench": {
18
+ "primary": true,
19
+ "metrics": [{ "type": "bandwidth", "value": "numel(shapes.x) * dtypeBytes(args.dtype) * 2" }]
20
+ }
21
+ },
22
+ {
23
+ "name": "acos-f32-1m-minus1-scalar-fallback",
24
+ "preset": "smoke",
25
+ "vars": { "dtype": "float32", "count": 1048575 },
26
+ "inputs": {
27
+ "x": { "shape": [1048575], "dtype": "float32", "dist": "uniform", "seed": 9102, "scale": 2, "offset": -1 }
28
+ },
29
+ "outputs": { "y": { "shape": [1048575], "dtype": "float32" } },
30
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.x) * dtypeBytes(args.dtype) * 2" }] }
31
+ },
32
+ {
33
+ "name": "acos-f16-1m-vec4",
34
+ "preset": "smoke",
35
+ "vars": { "dtype": "float16", "count": 1048576 },
36
+ "inputs": {
37
+ "x": { "shape": [1048576], "dtype": "float16", "dist": "uniform", "seed": 9103, "scale": 2, "offset": -1 }
38
+ },
39
+ "outputs": { "y": { "shape": [1048576], "dtype": "float16" } },
40
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.x) * dtypeBytes(args.dtype) * 2" }] }
41
+ },
42
+ {
43
+ "name": "acos-f16-1m-minus1-scalar-fallback",
44
+ "preset": "smoke",
45
+ "vars": { "dtype": "float16", "count": 1048575 },
46
+ "inputs": {
47
+ "x": { "shape": [1048575], "dtype": "float16", "dist": "uniform", "seed": 9104, "scale": 2, "offset": -1 }
48
+ },
49
+ "outputs": { "y": { "shape": [1048575], "dtype": "float16" } },
50
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.x) * dtypeBytes(args.dtype) * 2" }] }
51
+ }
52
+ ]
53
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,78 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "ai.onnx",
3
+ "name": "Acos",
4
+ "sinceVersion": 7,
5
+ "description": "Computes the arccosine (inverse cosine) of each element of the input tensor. The output has the same shape as the input.",
6
+ "inputs": [{ "role": "input", "dtype": "T", "description": "Cosine values to map elementwise to angles in radians." }],
7
+ "outputs": [
8
+ {
9
+ "role": "output",
10
+ "dtype": "T",
11
+ "rank": "ranks.input",
12
+ "description": "The arccosine of the input tensor, 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": "Acos.vec4",
35
+ "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"acos\"" } },
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": "Acos",
62
+ "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"acos\"", "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.Acos",
3
+ "id": "_ai_onnx_acos_webgpu_1054d90",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "33CNgk6mE3efmcd/KRM4IQdIQN57/Tbq/ra3Jyd5ft8=",
11
+ "manifest.json": "6LWa8VwpwgCVNBl+Y5m0/aKk+JbgmQ6ZkYyEqo2PD4c=",
12
+ "test.json": "od4CUyrfNDuubxZ0V6kOoLRSD6T1iVY5JhRIGAnGsB8=",
13
+ "unary-scalar.wgsl.jinja": "QS+cUNrZympm3XnYQz1pY4TNUwwK2JgQf08CHg3EauM=",
14
+ "unary-vec4.wgsl.jinja": "0DU51TVj67Sc14WWDqxJ37vHcmQ23wBCegC07Ydy4Ug="
15
+ }
16
+ },
17
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
18
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Acos" }
19
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,164 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.Acos",
3
+ "cases": [
4
+ {
5
+ "name": "f32_values",
6
+ "inputs": {
7
+ "x": {
8
+ "dtype": "float32",
9
+ "shape": [6],
10
+ "data": { "kind": "values", "values": [-1.0, -0.5, 0.0, 0.5, 1.0, 0.25] }
11
+ }
12
+ },
13
+ "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.000001 } }
14
+ },
15
+ {
16
+ "name": "f16_values",
17
+ "inputs": {
18
+ "x": {
19
+ "dtype": "float16",
20
+ "shape": [2, 3],
21
+ "data": { "kind": "values", "values": [-1.0, -0.5, 0.0, 0.5, 1.0, 0.25] }
22
+ }
23
+ },
24
+ "outputs": { "y": { "dtype": "float16", "shape": [2, 3] } },
25
+ "tolerance": 0.002
26
+ },
27
+ {
28
+ "name": "out_of_domain_returns_nan",
29
+ "inputs": {
30
+ "x": {
31
+ "dtype": "float32",
32
+ "shape": [6],
33
+ "data": { "kind": "values", "values": [1.5, 2.0, -1.5, -2.0, 5.0, -5.0] }
34
+ }
35
+ },
36
+ "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.000001, "allowNaN": true } }
37
+ },
38
+ {
39
+ "name": "f32_near_domain_edges_accuracy_scalar",
40
+ "provenance": {
41
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
42
+ "test": "MathOpTest.Acos",
43
+ "notes": "Scalar-path companion for Acos values immediately inside the domain boundary; exercises the split-pi complement path."
44
+ },
45
+ "inputs": {
46
+ "x": {
47
+ "dtype": "float32",
48
+ "shape": [3],
49
+ "data": { "kind": "values", "values": [0.9999999403953552, -0.9999999403953552, 0.0] }
50
+ }
51
+ },
52
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 1e-8, "relTolerance": 2.4e-7 } }
53
+ },
54
+ {
55
+ "name": "f16_domain_and_nonfinite_values",
56
+ "inputs": {
57
+ "x": {
58
+ "dtype": "float16",
59
+ "shape": [8],
60
+ "data": { "kind": "values", "values": [-1.0, -1.001, 0.0, 1.0, 1.001, "Infinity", "-Infinity", "NaN"] }
61
+ }
62
+ },
63
+ "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.002, "allowNaN": true } }
64
+ },
65
+ {
66
+ "name": "ort_float_unit_interval",
67
+ "provenance": {
68
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
69
+ "test": "MathOpTest.Acos"
70
+ },
71
+ "inputs": {
72
+ "x": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [-1.0, -0.5, 0.0, 0.5, 1.0] } }
73
+ },
74
+ "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0.00001 } }
75
+ },
76
+ {
77
+ "name": "ort_domain_and_nonfinite_values",
78
+ "provenance": {
79
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
80
+ "test": "MathOpTest.Acos",
81
+ "notes": "Extends ORT's unit-interval case with ONNX-valid out-of-domain and non-finite inputs."
82
+ },
83
+ "inputs": {
84
+ "x": {
85
+ "dtype": "float32",
86
+ "shape": [9],
87
+ "data": {
88
+ "kind": "values",
89
+ "values": [-1.0, 0.0, 0.0, 1.0, 1.0000001, -1.0000001, "Infinity", "-Infinity", "NaN"]
90
+ }
91
+ }
92
+ },
93
+ "outputs": { "y": { "dtype": "float32", "shape": [9], "tolerance": 0.00001, "allowNaN": true } }
94
+ },
95
+ {
96
+ "name": "onnx_backend_example",
97
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_acos_example" },
98
+ "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-0.5, 0.0, 0.5] } } },
99
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
100
+ },
101
+ {
102
+ "name": "onnx_backend_acos",
103
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_acos" },
104
+ "inputs": {
105
+ "x": {
106
+ "dtype": "float32",
107
+ "shape": [3, 4, 5],
108
+ "data": {
109
+ "kind": "values",
110
+ "values": [0.54881352186203, 0.7151893377304077, 0.6027633547782898, 0.5448831915855408, 0.42365479469299316, 0.6458941102027893, 0.4375872015953064, 0.891772985458374, 0.9636627435684204, 0.3834415078163147, 0.7917250394821167, 0.5288949012756348, 0.5680445432662964, 0.9255966544151306, 0.07103605568408966, 0.08712930232286453, 0.020218396559357643, 0.832619845867157, 0.7781567573547363, 0.8700121641159058, 0.978618323802948, 0.7991585731506348, 0.4614793658256531, 0.7805292010307312, 0.11827442795038223, 0.6399210095405579, 0.14335328340530396, 0.9446688890457153, 0.5218483209609985, 0.4146619439125061, 0.26455560326576233, 0.7742336988449097, 0.4561503231525421, 0.568433940410614, 0.018789799883961678, 0.6176354885101318, 0.6120957136154175, 0.6169340014457703, 0.9437480568885803, 0.681820273399353, 0.35950788855552673, 0.43703195452690125, 0.6976311802864075, 0.0602254718542099, 0.6667667031288147, 0.670637845993042, 0.21038256585597992, 0.12892629206180573, 0.31542834639549255, 0.36371076107025146, 0.5701967477798462, 0.4386015236377716, 0.9883738160133362, 0.10204481333494186, 0.20887675881385803, 0.16130951046943665, 0.6531082987785339, 0.25329160690307617, 0.4663107693195343, 0.24442559480667114]
111
+ }
112
+ }
113
+ },
114
+ "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0.00001, "allowNaN": true } }
115
+ },
116
+ {
117
+ "name": "onnx_backend_acos_example",
118
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_acos_example" },
119
+ "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-0.5, 0.0, 0.5] } } },
120
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.00001 } }
121
+ },
122
+ {
123
+ "name": "empty_input_zero_dim",
124
+ "inputs": { "x": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
125
+ "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
126
+ },
127
+ {
128
+ "name": "f16_just_inside_domain_edge_high_slope",
129
+ "inputs": {
130
+ "x": {
131
+ "dtype": "float16",
132
+ "shape": [8],
133
+ "data": { "kind": "values", "values": [0.999, -0.999, 0.99, -0.99, 1.0, -1.0, 0.95, -0.95] }
134
+ }
135
+ },
136
+ "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.01 } }
137
+ },
138
+ {
139
+ "name": "f32_vec4_mixed_domain_nan",
140
+ "inputs": {
141
+ "x": {
142
+ "dtype": "float32",
143
+ "shape": [8],
144
+ "data": { "kind": "values", "values": [-1.0, 0.0, 0.5, 1.0, 1.5, -1.5, 2.0, -2.0] }
145
+ }
146
+ },
147
+ "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.000001, "allowNaN": true } }
148
+ },
149
+ {
150
+ "name": "f32_scalar_tail_4095_domain",
151
+ "provenance": {
152
+ "notes": "Compact sibling for Acos scalar-tail benchmarks; odd numel forces the elementwise path while bounded generated values stay inside [-1, 1]."
153
+ },
154
+ "inputs": {
155
+ "x": {
156
+ "dtype": "float32",
157
+ "shape": [4095],
158
+ "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.45 }
159
+ }
160
+ },
161
+ "outputs": { "y": { "dtype": "float32", "shape": [4095], "tolerance": 0.00001 } }
162
+ }
163
+ ]
164
+ }
build/webgpu/unary-scalar.wgsl.jinja ADDED
@@ -0,0 +1,96 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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_PI: f32 = 3.141592653589793;
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
+
64
+ fn precise_domain_nan(value: f32) -> f32 {
65
+ return bitcast<f32>((bitcast<u32>(value) & 0x80000000u) | 0x7fc00000u);
66
+ }
67
+
68
+ fn acos_accurate(value: f32) -> f32 {
69
+ let bits = bitcast<u32>(value);
70
+ let magnitudeBits = bits & 0x7fffffffu;
71
+ if (magnitudeBits > 0x7f800000u) {
72
+ return value;
73
+ }
74
+ if (abs(value) > 1.0) {
75
+ return precise_domain_nan(value);
76
+ }
77
+ if (value == 1.0) {
78
+ return 0.0;
79
+ }
80
+ if (value == -1.0) {
81
+ return PRECISE_INV_TRIG_PI;
82
+ }
83
+ if (value >= 0.0) {
84
+ return 2.0 * precise_atan_positive(sqrt((1.0 - value) / (1.0 + value)));
85
+ }
86
+ // The reciprocal half-angle form produces the near-pi result directly and
87
+ // avoids a reassociable final pi-small subtraction.
88
+ return 2.0 * precise_atan_positive(
89
+ sqrt((1.0 - value) / (1.0 + value)));
90
+ }
91
+
92
+ {{ flat_tail_open() }}
93
+ let v = f32(x[i]);
94
+ y[i] = {{ scalar }}(acos_accurate(v));
95
+ {{ flat_tail_close() -}}
96
+ }
build/webgpu/unary-vec4.wgsl.jinja ADDED
@@ -0,0 +1,89 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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_PI: f32 = 3.141592653589793;
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
+
49
+ fn precise_domain_nan(value: f32) -> f32 {
50
+ return bitcast<f32>((bitcast<u32>(value) & 0x80000000u) | 0x7fc00000u);
51
+ }
52
+
53
+ fn acos_accurate(value: f32) -> f32 {
54
+ let bits = bitcast<u32>(value);
55
+ let magnitudeBits = bits & 0x7fffffffu;
56
+ if (magnitudeBits > 0x7f800000u) {
57
+ return value;
58
+ }
59
+ if (abs(value) > 1.0) {
60
+ return precise_domain_nan(value);
61
+ }
62
+ if (value == 1.0) {
63
+ return 0.0;
64
+ }
65
+ if (value == -1.0) {
66
+ return PRECISE_INV_TRIG_PI;
67
+ }
68
+ if (value >= 0.0) {
69
+ return 2.0 * precise_atan_positive(sqrt((1.0 - value) / (1.0 + value)));
70
+ }
71
+ // The reciprocal half-angle form produces the near-pi result directly and
72
+ // avoids a reassociable final pi-small subtraction.
73
+ return 2.0 * precise_atan_positive(
74
+ sqrt((1.0 - value) / (1.0 + value)));
75
+ }
76
+
77
+
78
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
79
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
80
+ // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
81
+ // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
82
+ let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
83
+ if (i >= params.count) {
84
+ return;
85
+ }
86
+ let xv = x[i];
87
+ let fv = vec4<f32>(xv);
88
+ y[i] = {{ vectorScalar }}(vec4<f32>(acos_accurate(fv.x), acos_accurate(fv.y), acos_accurate(fv.z), acos_accurate(fv.w)));
89
+ }