Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Tan

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

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 +80 -0
  4. build/webgpu/metadata.json +19 -0
  5. build/webgpu/test.json +229 -0
  6. build/webgpu/unary-scalar.wgsl.jinja +229 -0
  7. build/webgpu/unary-vec4.wgsl.jinja +222 -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.Tan
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 7
12
+
13
+ ## Description
14
+
15
+ Computes the tangent of each element of the input tensor. Output has the same shape and type as the input.
16
+
17
+ See the [ONNX `Tan` spec](https://onnx.ai/onnx/operators/onnx__Tan.html) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `input` | `x` | `T` | — | — | Angles in radians whose tangent is computed elementwise. | 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` | Elementwise tangent of the input tensor. | 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.Tan", { 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.Tan",
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": "tan-f32-1m-vec4",
11
+ "preset": "smoke",
12
+ "vars": { "dtype": "float32", "count": 1048576 },
13
+ "inputs": {
14
+ "x": { "shape": [1048576], "dtype": "float32", "dist": "uniform", "seed": 7101, "scale": 2.5, "offset": -1.25 }
15
+ },
16
+ "outputs": { "y": { "shape": [1048576], "dtype": "float32" } },
17
+ "bench": {
18
+ "primary": true,
19
+ "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }]
20
+ }
21
+ },
22
+ {
23
+ "name": "tan-f16-1m-vec4",
24
+ "preset": "smoke",
25
+ "vars": { "dtype": "float16", "count": 1048576 },
26
+ "inputs": {
27
+ "x": { "shape": [1048576], "dtype": "float16", "dist": "uniform", "seed": 7102, "scale": 2.5, "offset": -1.25 }
28
+ },
29
+ "outputs": { "y": { "shape": [1048576], "dtype": "float16" } },
30
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
31
+ },
32
+ {
33
+ "name": "tan-f32-1m-plus-1-scalar-fallback",
34
+ "preset": "smoke",
35
+ "vars": { "dtype": "float32", "count": 1048577 },
36
+ "inputs": {
37
+ "x": { "shape": [1048577], "dtype": "float32", "dist": "uniform", "seed": 7103, "scale": 2.5, "offset": -1.25 }
38
+ },
39
+ "outputs": { "y": { "shape": [1048577], "dtype": "float32" } },
40
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
41
+ },
42
+ {
43
+ "name": "tan-f16-1m-plus-1-scalar-fallback",
44
+ "preset": "smoke",
45
+ "vars": { "dtype": "float16", "count": 1048577 },
46
+ "inputs": {
47
+ "x": { "shape": [1048577], "dtype": "float16", "dist": "uniform", "seed": 7104, "scale": 2.5, "offset": -1.25 }
48
+ },
49
+ "outputs": { "y": { "shape": [1048577], "dtype": "float16" } },
50
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
51
+ }
52
+ ]
53
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,80 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "ai.onnx",
3
+ "name": "Tan",
4
+ "sinceVersion": 7,
5
+ "description": "Computes the tangent of each element of the input tensor. Output has the same shape and type as the input.",
6
+ "inputs": [
7
+ { "role": "input", "dtype": "T", "description": "Angles in radians whose tangent is computed elementwise." }
8
+ ],
9
+ "outputs": [
10
+ {
11
+ "role": "output",
12
+ "dtype": "T",
13
+ "rank": "ranks.input",
14
+ "description": "Elementwise tangent of the input tensor.",
15
+ "shape": "shapes.input"
16
+ }
17
+ ],
18
+ "typeConstraints": { "T": ["float32", "float16"] },
19
+ "args": {
20
+ "x": { "kind": "tensor", "semantic": "input", "role": "input" },
21
+ "y": { "kind": "tensor", "semantic": "output", "role": "output" }
22
+ },
23
+ "tunables": { "WORKGROUP_SIZE": 256 },
24
+ "variants": [
25
+ {
26
+ "id": "same_layout_vec4",
27
+ "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
28
+ "constants": {
29
+ "scalar": "dtypes.T",
30
+ "usesF16": "dtypes.T == \"f16\"",
31
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
32
+ },
33
+ "passes": [
34
+ {
35
+ "id": "main",
36
+ "name": "Tan.vec4",
37
+ "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"tan\"" } },
38
+ "bindings": [
39
+ { "name": "x", "arg": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$vectorScalar" },
40
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
41
+ {
42
+ "name": "params",
43
+ "semantic": "kernel.params",
44
+ "buffer": { "type": "uniform" },
45
+ "struct": {
46
+ "name": "Params",
47
+ "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }]
48
+ }
49
+ }
50
+ ],
51
+ "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
52
+ }
53
+ ],
54
+ "priority": 20
55
+ },
56
+ {
57
+ "id": "elementwise",
58
+ "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
59
+ "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
60
+ "passes": [
61
+ {
62
+ "id": "main",
63
+ "name": "Tan",
64
+ "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"tan\"", "itemsPerInvocation": 4 } },
65
+ "bindings": [
66
+ { "name": "x", "arg": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
67
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
68
+ {
69
+ "name": "params",
70
+ "semantic": "kernel.params",
71
+ "buffer": { "type": "uniform" },
72
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
73
+ }
74
+ ],
75
+ "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
76
+ }
77
+ ]
78
+ }
79
+ ]
80
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,19 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "ai.onnx.Tan",
3
+ "id": "_ai_onnx_tan_webgpu_3a9b5fc",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "2CXV9SpZgdr6/7N0OR5rZ1wmrXkfUgB3zDtaD5qOy1A=",
11
+ "manifest.json": "5ypWz6oe87DcGIBgESYuQRKGCzPq28CxLJD+M8blfd0=",
12
+ "test.json": "fVVVlfiWuVBfJUiOGWsi07VGA3Ofu6oegzBWn0pj6Lk=",
13
+ "unary-scalar.wgsl.jinja": "WzaWbiyC7b96tk5M2dpFdOTRzgTaIkF7eyBPFjTwkUE=",
14
+ "unary-vec4.wgsl.jinja": "2TOUEu3TR4elnjHDM/TxNE2kzUUtm6NlX2iC/Eov/gg="
15
+ }
16
+ },
17
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
18
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Tan" }
19
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,229 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.Tan",
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, 1.25] }
11
+ }
12
+ },
13
+ "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.00001 } }
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.Tan",
24
+ "notes": "For tiny finite inputs tan(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.Tan",
44
+ "notes": "Scalar-path companion: subnormal inputs are valid finite Tan 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": [-1.0, -0.5, 0.0, 0.5, 1.0, 1.25] }
58
+ }
59
+ },
60
+ "outputs": { "y": { "dtype": "float16", "shape": [2, 3] } },
61
+ "tolerance": 0.002
62
+ },
63
+ {
64
+ "name": "large_argument_range_reduction",
65
+ "inputs": {
66
+ "x": {
67
+ "dtype": "float32",
68
+ "shape": [8],
69
+ "data": {
70
+ "kind": "values",
71
+ "values": [1000000.0, 10000000.0, 10000000000000.0, 100000000000000000000.0, -1000000000000000.0, -123456.78, 314159.265, 2500000000.0]
72
+ }
73
+ }
74
+ },
75
+ "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.001 } }
76
+ },
77
+ {
78
+ "name": "f32_large_argument_range_reduction_accuracy_gpu_gap",
79
+ "skipGpu": {
80
+ "category": "todo",
81
+ "reason": "The current f32 large-argument range reduction differs from the CPU reference by about one ULP. A more accurate range reduction or software-extended precision could close this implementable gap."
82
+ },
83
+ "provenance": {
84
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
85
+ "test": "MathOpTest.Tan",
86
+ "notes": "Finite large arguments stress Tan's range reduction; ORT and the CPU reference agree more tightly than the current WGSL builtin path on these float32 inputs."
87
+ },
88
+ "inputs": {
89
+ "x": {
90
+ "dtype": "float32",
91
+ "shape": [8],
92
+ "data": {
93
+ "kind": "values",
94
+ "values": [1000000.0, 10000000.0, 10000000000000.0, 100000000000000000000.0, -1000000000000000.0, -123456.78, 314159.265, 2500000000.0]
95
+ }
96
+ }
97
+ },
98
+ "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.000001, "relTolerance": 0 } }
99
+ },
100
+ {
101
+ "name": "ort_float_large_arguments",
102
+ "provenance": {
103
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
104
+ "test": "MathOpTest.Tan"
105
+ },
106
+ "inputs": {
107
+ "x": {
108
+ "dtype": "float32",
109
+ "shape": [5],
110
+ "data": { "kind": "values", "values": [-100.0, -50.0, 0.0, 50.0, 100.0] }
111
+ }
112
+ },
113
+ "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0.0001 } }
114
+ },
115
+ {
116
+ "name": "ort_nonfinite_values",
117
+ "provenance": {
118
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
119
+ "test": "MathOpTest.Tan",
120
+ "notes": "Extends ORT's large-argument case with signed infinities and NaN."
121
+ },
122
+ "inputs": {
123
+ "x": {
124
+ "dtype": "float32",
125
+ "shape": [7],
126
+ "data": { "kind": "values", "values": ["-Infinity", -1.0, 0.0, 0.0, 1.0, "Infinity", "NaN"] }
127
+ }
128
+ },
129
+ "outputs": { "y": { "dtype": "float32", "shape": [7], "tolerance": 0.0001, "allowNaN": true } }
130
+ },
131
+ {
132
+ "name": "onnx_backend_example",
133
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_tan_example" },
134
+ "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } } },
135
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
136
+ },
137
+ {
138
+ "name": "onnx_backend_tan",
139
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_tan" },
140
+ "inputs": {
141
+ "x": {
142
+ "dtype": "float32",
143
+ "shape": [3, 4, 5],
144
+ "data": {
145
+ "kind": "values",
146
+ "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]
147
+ }
148
+ }
149
+ },
150
+ "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0.00001 } }
151
+ },
152
+ {
153
+ "name": "onnx_backend_tan_example",
154
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_tan_example" },
155
+ "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } } },
156
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.00001 } }
157
+ },
158
+ {
159
+ "name": "vec4_f16_lanes",
160
+ "inputs": {
161
+ "x": {
162
+ "dtype": "float16",
163
+ "shape": [16],
164
+ "data": {
165
+ "kind": "values",
166
+ "values": [-1.2, -1.0, -0.75, -0.5, -0.25, -0.125, 0.0, 0.125, 0.25, 0.5, 0.75, 1.0, 1.2, 2.0, 3.0, 5.0]
167
+ }
168
+ }
169
+ },
170
+ "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0.001, "relTolerance": 0.005 } }
171
+ },
172
+ {
173
+ "name": "vec4_f32_nonfinite",
174
+ "inputs": {
175
+ "x": {
176
+ "dtype": "float32",
177
+ "shape": [8],
178
+ "data": { "kind": "values", "values": ["-Infinity", -1.0, 0.0, 0.5, 1.0, "Infinity", "NaN", 2.0] }
179
+ }
180
+ },
181
+ "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.000001, "allowNaN": true } }
182
+ },
183
+ {
184
+ "name": "empty_input_zero_dim",
185
+ "inputs": { "x": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
186
+ "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
187
+ },
188
+ {
189
+ "name": "near_pole_argument_accuracy",
190
+ "provenance": {
191
+ "source": "ai.onnx.Tan reference (Math.tan)",
192
+ "notes": "Tan blows up near odd multiples of pi/2; tan(x) ~ 1/(pi/2 - x) so a tiny argument error is amplified into a large relative output error. These arguments are below the kernel's 1e4 range-reduction threshold and so route through the hardware tan() directly. Validates that the Metal hardware tan() matches the CPU Math.tan reference at moderate near-pole arguments (the riskier untested regime; pi/4 and pi anchors are well-conditioned)."
193
+ },
194
+ "inputs": {
195
+ "x": {
196
+ "dtype": "float32",
197
+ "shape": [8],
198
+ "data": {
199
+ "kind": "values",
200
+ "values": [1.5706963267948966, 1.5697963267948967, -1.5706963267948966, 4.711388980384689, 0.7853981633974483, -0.7853981633974483, 3.1405926535897932, 0.0]
201
+ }
202
+ }
203
+ },
204
+ "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.01, "relTolerance": 0.0005 } }
205
+ },
206
+ {
207
+ "name": "f16_scalar_near_f16_max",
208
+ "inputs": {
209
+ "x": {
210
+ "dtype": "float16",
211
+ "shape": [7],
212
+ "data": { "kind": "values", "values": [65504.0, -65504.0, 20000.0, -20000.0, 50000.0, -50000.0, 0.0] }
213
+ }
214
+ },
215
+ "outputs": { "y": { "dtype": "float16", "shape": [7], "tolerance": 0.5, "relTolerance": 0.01 } }
216
+ },
217
+ {
218
+ "name": "f16_vec4_large_arg_above_reduce_threshold",
219
+ "inputs": {
220
+ "x": {
221
+ "dtype": "float16",
222
+ "shape": [8],
223
+ "data": { "kind": "values", "values": [20000.0, -20000.0, 40000.0, -40000.0, 0.5, -0.5, 1.0, -1.0] }
224
+ }
225
+ },
226
+ "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.5, "relTolerance": 0.01 } }
227
+ }
228
+ ]
229
+ }
build/webgpu/unary-scalar.wgsl.jinja ADDED
@@ -0,0 +1,229 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
+ {% macro emit_is_nan_f32() %}
24
+ fn is_nan_f32(value: f32) -> bool {
25
+ let bits = bitcast<u32>(value);
26
+ return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
27
+ }{% endmacro %}
28
+ {% macro emit_is_inf_f32() %}
29
+ fn is_inf_f32(value: f32) -> bool {
30
+ return (bitcast<u32>(value) & 0x7fffffffu) == 0x7f800000u;
31
+ }{% endmacro %}
32
+ {% macro emit_nan_from_inf_f32() %}
33
+ fn nan_from_inf_f32(value: f32) -> f32 {
34
+ return bitcast<f32>(bitcast<u32>(value) | 0x00400000u);
35
+ }{% endmacro %}
36
+ {% macro emit_reduce_pio2() %}
37
+ fn reduce_pio2(ax: f32) -> Pio2 {
38
+ let ix = bitcast<u32>(ax);
39
+ let e = (ix >> 23u) & 0xffu;
40
+ let mant = (ix & 0x7fffffu) | 0x800000u; // 24-bit mantissa
41
+ let p = i32(e) - 150; // ax = mant * 2^p
42
+ // 96-bit accumulator a2:a1:a0 holding (ax * 2/pi mod 4) * 2^94.
43
+ var a0 = 0u; var a1 = 0u; var a2 = 0u;
44
+ for (var k = 0u; k < 14u; k = k + 1u) {
45
+ let shift = p - 24 * (i32(k) + 1) + 94;
46
+ if (shift <= -48 || shift >= 96) { continue; }
47
+ // prod = mant * TWO_OVER_PI[k] (<= 48 bits) -> phi32:plo
48
+ let t = TWO_OVER_PI[k];
49
+ let mlo = mant & 0xffffu; let mhi = mant >> 16u;
50
+ let tlo = t & 0xffffu; let thi = t >> 16u;
51
+ let ll = mlo * tlo;
52
+ let mid = mlo * thi + mhi * tlo; // <= 25 bits, no overflow
53
+ let hh = mhi * thi;
54
+ let plo = ll + ((mid & 0xffffu) << 16u);
55
+ let carry = select(0u, 1u, plo < ll);
56
+ let phi32 = hh + (mid >> 16u) + carry; // <= 16 bits
57
+ // Place (phi32:plo) << shift as a 64-bit value (v1:v0) at offset off >= 0.
58
+ var v0 = 0u; var v1 = 0u; var off = 0;
59
+ if (shift < 0) {
60
+ let s = u32(-shift); // 1..47
61
+ if (s < 32u) {
62
+ v0 = (plo >> s) | (phi32 << (32u - s));
63
+ v1 = phi32 >> s;
64
+ } else {
65
+ v0 = phi32 >> (s - 32u);
66
+ v1 = 0u;
67
+ }
68
+ off = 0;
69
+ } else {
70
+ v0 = plo; v1 = phi32; off = shift; // 0..95
71
+ }
72
+ // Spread (v1:v0) across 96-bit limbs starting at bit `off`.
73
+ let li = off / 32; let b = u32(off % 32);
74
+ var w0 = 0u; var w1 = 0u; var w2 = 0u;
75
+ if (b == 0u) { w0 = v0; w1 = v1; w2 = 0u; }
76
+ else {
77
+ w0 = v0 << b;
78
+ w1 = (v1 << b) | (v0 >> (32u - b));
79
+ w2 = v1 >> (32u - b);
80
+ }
81
+ var t0 = 0u; var t1 = 0u; var t2 = 0u;
82
+ if (li == 0) { t0 = w0; t1 = w1; t2 = w2; }
83
+ else if (li == 1) { t1 = w0; t2 = w1; }
84
+ else { t2 = w0; }
85
+ // 96-bit add acc += t (bits above 2^96 are multiples of 4, so discarded).
86
+ let n0 = a0 + t0; let c0 = select(0u, 1u, n0 < a0);
87
+ let n1a = a1 + t1; let c1a = select(0u, 1u, n1a < a1);
88
+ let n1 = n1a + c0; let c1b = select(0u, 1u, n1 < n1a);
89
+ let n2 = a2 + t2 + c1a + c1b;
90
+ a0 = n0; a1 = n1; a2 = n2;
91
+ }
92
+ var n = a2 >> 30u; // integer part mod 4
93
+ var phi = f32(a2 & 0x3fffffffu) * (1.0 / 1073741824.0)
94
+ + f32(a1) * (1.0 / 4611686018427387904.0); // fraction in [0,1)
95
+ if (phi >= 0.5) { phi = phi - 1.0; n = n + 1u; }
96
+ var out: Pio2;
97
+ out.octant = n & 3u;
98
+ out.r = phi * PIO2_F;
99
+ return out;
100
+ }{% endmacro %}
101
+ {% macro emit_reduce_pio2_fast() %}
102
+ fn reduce_pio2_fast(ax: f32) -> Pio2 {
103
+ // Cody-Waite reduction for ordinary magnitudes. Splitting pi/2 keeps the
104
+ // residual accurate even when n is thousands; larger inputs use Payne-Hanek.
105
+ let nFloat = floor(fma(ax, INV_PIO2_F, 0.5));
106
+ var residual = fma(-nFloat, PIO2_HI_F, ax);
107
+ residual = fma(-nFloat, PIO2_LO_F, residual);
108
+ var out: Pio2;
109
+ out.octant = u32(nFloat) & 3u;
110
+ out.r = residual;
111
+ return out;
112
+ }{% endmacro %}
113
+ {% macro emit_trig_reduction_support() %}
114
+ {% set needPreciseTrigCentered = false %}
115
+ {% set needPreciseTrigTwoPi = false %}
116
+ // Backend-stable f32 sine/cosine core.
117
+ //
118
+ // Shader transcendental accuracy is implementation-defined, and some portable
119
+ // backends are only accurate to roughly 1e-4. Each path retains the most accurate
120
+ // available phase representation, reduces it to [-pi, pi], then uses these
121
+ // polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
122
+ // truncation error is well below one f32 ULP over its documented interval.
123
+ {% if needPreciseTrigCentered %}
124
+ const PRECISE_TRIG_PI: f32 = 3.141592653589793;
125
+ {% endif %}
126
+ {% if needPreciseTrigTwoPi %}
127
+ const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
128
+ {% endif %}
129
+ {% if needPreciseTrigCentered %}
130
+ const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
131
+ {% endif %}
132
+
133
+ // Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
134
+ fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
135
+ let x2 = x * x;
136
+
137
+ var sinPolynomial = 1.6059043836821613e-10;
138
+ sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
139
+ sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
140
+ sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
141
+ sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
142
+ sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
143
+ let sine = x * fma(sinPolynomial, x2, 1.0);
144
+
145
+ var cosPolynomial = 2.08767569878681e-9;
146
+ cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
147
+ cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
148
+ cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
149
+ cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
150
+ cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
151
+ let cosine = fma(cosPolynomial, x2, 1.0);
152
+
153
+ return vec2<f32>(cosine, sine);
154
+ }
155
+ {% if needPreciseTrigCentered %}
156
+
157
+ // Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
158
+ fn precise_sincos_centered(x: f32) -> vec2<f32> {
159
+ var folded = x;
160
+ var cosineSign = 1.0;
161
+ if (folded > PRECISE_TRIG_HALF_PI) {
162
+ folded = PRECISE_TRIG_PI - folded;
163
+ cosineSign = -1.0;
164
+ } else if (folded < -PRECISE_TRIG_HALF_PI) {
165
+ folded = -PRECISE_TRIG_PI - folded;
166
+ cosineSign = -1.0;
167
+ }
168
+ let value = precise_sincos_half_pi(folded);
169
+ return vec2<f32>(cosineSign * value.x, value.y);
170
+ }
171
+ {% endif %}
172
+
173
+
174
+ // 24-bit words of 2/pi used by the exact Payne-Hanek large-input reducer.
175
+ const TWO_OVER_PI: array<u32, 14> = array<u32, 14>(
176
+ 0xa2f983u, 0x6e4e44u, 0x1529fcu, 0x2757d1u, 0xf534ddu, 0xc0db62u,
177
+ 0x95993cu, 0x439041u, 0xfe5163u, 0xabdebbu, 0xc561b7u, 0x246e3au,
178
+ 0x424dd2u, 0xe00649u
179
+ );
180
+ const PIO2_F: f32 = 1.5707963267948966;
181
+ const INV_PIO2_F: f32 = 0.6366197723675814;
182
+ const PIO2_HI_F: f32 = 1.570796251296997;
183
+ const PIO2_LO_F: f32 = 7.549789415861596e-8;
184
+ const REDUCE_THRESHOLD: f32 = 1.0e4;
185
+
186
+ struct Pio2 { octant: u32, r: f32 };
187
+
188
+ {{ emit_reduce_pio2_fast() }}
189
+
190
+ {{ emit_reduce_pio2() }}
191
+ {%- endmacro %}
192
+ {% macro emit_tan_accurate() %}
193
+ fn tan_accurate(x: f32) -> f32 {
194
+ if (is_nan_f32(x)) {
195
+ return x;
196
+ }
197
+ if (is_inf_f32(x)) {
198
+ return nan_from_inf_f32(x);
199
+ }
200
+ let ax = abs(x);
201
+ var red: Pio2;
202
+ if (ax < REDUCE_THRESHOLD) { red = reduce_pio2_fast(ax); }
203
+ else { red = reduce_pio2(ax); }
204
+ let reduced = precise_sincos_half_pi(red.r);
205
+ let sr = reduced.y;
206
+ let cr = reduced.x;
207
+ var sn: f32; var cs: f32;
208
+ switch (red.octant) {
209
+ case 0u: { sn = sr; cs = cr; }
210
+ case 1u: { sn = cr; cs = -sr; }
211
+ case 2u: { sn = -sr; cs = -cr; }
212
+ default: { sn = -cr; cs = sr; }
213
+ }
214
+ let tn = sn / cs;
215
+ return select(tn, -tn, x < 0.0); // tan is odd
216
+ }{% endmacro %}
217
+ {{ emit_trig_reduction_support() }}
218
+
219
+ {{ emit_is_nan_f32() }}
220
+
221
+ {{ emit_is_inf_f32() }}
222
+
223
+ {{ emit_nan_from_inf_f32() }}
224
+
225
+ {{ emit_tan_accurate() }}
226
+ {{ flat_tail_open() }}
227
+ y[i] = {{ scalar }}(tan_accurate(f32(x[i])));
228
+ {{ flat_tail_close() -}}
229
+ }
build/webgpu/unary-vec4.wgsl.jinja ADDED
@@ -0,0 +1,222 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
+ {% macro emit_is_nan_f32() %}
9
+ fn is_nan_f32(value: f32) -> bool {
10
+ let bits = bitcast<u32>(value);
11
+ return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
12
+ }{% endmacro %}
13
+ {% macro emit_is_inf_f32() %}
14
+ fn is_inf_f32(value: f32) -> bool {
15
+ return (bitcast<u32>(value) & 0x7fffffffu) == 0x7f800000u;
16
+ }{% endmacro %}
17
+ {% macro emit_nan_from_inf_f32() %}
18
+ fn nan_from_inf_f32(value: f32) -> f32 {
19
+ return bitcast<f32>(bitcast<u32>(value) | 0x00400000u);
20
+ }{% endmacro %}
21
+ {% macro emit_reduce_pio2() %}
22
+ fn reduce_pio2(ax: f32) -> Pio2 {
23
+ let ix = bitcast<u32>(ax);
24
+ let e = (ix >> 23u) & 0xffu;
25
+ let mant = (ix & 0x7fffffu) | 0x800000u; // 24-bit mantissa
26
+ let p = i32(e) - 150; // ax = mant * 2^p
27
+ // 96-bit accumulator a2:a1:a0 holding (ax * 2/pi mod 4) * 2^94.
28
+ var a0 = 0u; var a1 = 0u; var a2 = 0u;
29
+ for (var k = 0u; k < 14u; k = k + 1u) {
30
+ let shift = p - 24 * (i32(k) + 1) + 94;
31
+ if (shift <= -48 || shift >= 96) { continue; }
32
+ // prod = mant * TWO_OVER_PI[k] (<= 48 bits) -> phi32:plo
33
+ let t = TWO_OVER_PI[k];
34
+ let mlo = mant & 0xffffu; let mhi = mant >> 16u;
35
+ let tlo = t & 0xffffu; let thi = t >> 16u;
36
+ let ll = mlo * tlo;
37
+ let mid = mlo * thi + mhi * tlo; // <= 25 bits, no overflow
38
+ let hh = mhi * thi;
39
+ let plo = ll + ((mid & 0xffffu) << 16u);
40
+ let carry = select(0u, 1u, plo < ll);
41
+ let phi32 = hh + (mid >> 16u) + carry; // <= 16 bits
42
+ // Place (phi32:plo) << shift as a 64-bit value (v1:v0) at offset off >= 0.
43
+ var v0 = 0u; var v1 = 0u; var off = 0;
44
+ if (shift < 0) {
45
+ let s = u32(-shift); // 1..47
46
+ if (s < 32u) {
47
+ v0 = (plo >> s) | (phi32 << (32u - s));
48
+ v1 = phi32 >> s;
49
+ } else {
50
+ v0 = phi32 >> (s - 32u);
51
+ v1 = 0u;
52
+ }
53
+ off = 0;
54
+ } else {
55
+ v0 = plo; v1 = phi32; off = shift; // 0..95
56
+ }
57
+ // Spread (v1:v0) across 96-bit limbs starting at bit `off`.
58
+ let li = off / 32; let b = u32(off % 32);
59
+ var w0 = 0u; var w1 = 0u; var w2 = 0u;
60
+ if (b == 0u) { w0 = v0; w1 = v1; w2 = 0u; }
61
+ else {
62
+ w0 = v0 << b;
63
+ w1 = (v1 << b) | (v0 >> (32u - b));
64
+ w2 = v1 >> (32u - b);
65
+ }
66
+ var t0 = 0u; var t1 = 0u; var t2 = 0u;
67
+ if (li == 0) { t0 = w0; t1 = w1; t2 = w2; }
68
+ else if (li == 1) { t1 = w0; t2 = w1; }
69
+ else { t2 = w0; }
70
+ // 96-bit add acc += t (bits above 2^96 are multiples of 4, so discarded).
71
+ let n0 = a0 + t0; let c0 = select(0u, 1u, n0 < a0);
72
+ let n1a = a1 + t1; let c1a = select(0u, 1u, n1a < a1);
73
+ let n1 = n1a + c0; let c1b = select(0u, 1u, n1 < n1a);
74
+ let n2 = a2 + t2 + c1a + c1b;
75
+ a0 = n0; a1 = n1; a2 = n2;
76
+ }
77
+ var n = a2 >> 30u; // integer part mod 4
78
+ var phi = f32(a2 & 0x3fffffffu) * (1.0 / 1073741824.0)
79
+ + f32(a1) * (1.0 / 4611686018427387904.0); // fraction in [0,1)
80
+ if (phi >= 0.5) { phi = phi - 1.0; n = n + 1u; }
81
+ var out: Pio2;
82
+ out.octant = n & 3u;
83
+ out.r = phi * PIO2_F;
84
+ return out;
85
+ }{% endmacro %}
86
+ {% macro emit_reduce_pio2_fast() %}
87
+ fn reduce_pio2_fast(ax: f32) -> Pio2 {
88
+ // Cody-Waite reduction for ordinary magnitudes. Splitting pi/2 keeps the
89
+ // residual accurate even when n is thousands; larger inputs use Payne-Hanek.
90
+ let nFloat = floor(fma(ax, INV_PIO2_F, 0.5));
91
+ var residual = fma(-nFloat, PIO2_HI_F, ax);
92
+ residual = fma(-nFloat, PIO2_LO_F, residual);
93
+ var out: Pio2;
94
+ out.octant = u32(nFloat) & 3u;
95
+ out.r = residual;
96
+ return out;
97
+ }{% endmacro %}
98
+ {% macro emit_trig_reduction_support() %}
99
+ {% set needPreciseTrigCentered = false %}
100
+ {% set needPreciseTrigTwoPi = false %}
101
+ // Backend-stable f32 sine/cosine core.
102
+ //
103
+ // Shader transcendental accuracy is implementation-defined, and some portable
104
+ // backends are only accurate to roughly 1e-4. Each path retains the most accurate
105
+ // available phase representation, reduces it to [-pi, pi], then uses these
106
+ // polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
107
+ // truncation error is well below one f32 ULP over its documented interval.
108
+ {% if needPreciseTrigCentered %}
109
+ const PRECISE_TRIG_PI: f32 = 3.141592653589793;
110
+ {% endif %}
111
+ {% if needPreciseTrigTwoPi %}
112
+ const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
113
+ {% endif %}
114
+ {% if needPreciseTrigCentered %}
115
+ const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
116
+ {% endif %}
117
+
118
+ // Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
119
+ fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
120
+ let x2 = x * x;
121
+
122
+ var sinPolynomial = 1.6059043836821613e-10;
123
+ sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
124
+ sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
125
+ sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
126
+ sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
127
+ sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
128
+ let sine = x * fma(sinPolynomial, x2, 1.0);
129
+
130
+ var cosPolynomial = 2.08767569878681e-9;
131
+ cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
132
+ cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
133
+ cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
134
+ cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
135
+ cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
136
+ let cosine = fma(cosPolynomial, x2, 1.0);
137
+
138
+ return vec2<f32>(cosine, sine);
139
+ }
140
+ {% if needPreciseTrigCentered %}
141
+
142
+ // Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
143
+ fn precise_sincos_centered(x: f32) -> vec2<f32> {
144
+ var folded = x;
145
+ var cosineSign = 1.0;
146
+ if (folded > PRECISE_TRIG_HALF_PI) {
147
+ folded = PRECISE_TRIG_PI - folded;
148
+ cosineSign = -1.0;
149
+ } else if (folded < -PRECISE_TRIG_HALF_PI) {
150
+ folded = -PRECISE_TRIG_PI - folded;
151
+ cosineSign = -1.0;
152
+ }
153
+ let value = precise_sincos_half_pi(folded);
154
+ return vec2<f32>(cosineSign * value.x, value.y);
155
+ }
156
+ {% endif %}
157
+
158
+
159
+ // 24-bit words of 2/pi used by the exact Payne-Hanek large-input reducer.
160
+ const TWO_OVER_PI: array<u32, 14> = array<u32, 14>(
161
+ 0xa2f983u, 0x6e4e44u, 0x1529fcu, 0x2757d1u, 0xf534ddu, 0xc0db62u,
162
+ 0x95993cu, 0x439041u, 0xfe5163u, 0xabdebbu, 0xc561b7u, 0x246e3au,
163
+ 0x424dd2u, 0xe00649u
164
+ );
165
+ const PIO2_F: f32 = 1.5707963267948966;
166
+ const INV_PIO2_F: f32 = 0.6366197723675814;
167
+ const PIO2_HI_F: f32 = 1.570796251296997;
168
+ const PIO2_LO_F: f32 = 7.549789415861596e-8;
169
+ const REDUCE_THRESHOLD: f32 = 1.0e4;
170
+
171
+ struct Pio2 { octant: u32, r: f32 };
172
+
173
+ {{ emit_reduce_pio2_fast() }}
174
+
175
+ {{ emit_reduce_pio2() }}
176
+ {%- endmacro %}
177
+ {% macro emit_tan_accurate() %}
178
+ fn tan_accurate(x: f32) -> f32 {
179
+ if (is_nan_f32(x)) {
180
+ return x;
181
+ }
182
+ if (is_inf_f32(x)) {
183
+ return nan_from_inf_f32(x);
184
+ }
185
+ let ax = abs(x);
186
+ var red: Pio2;
187
+ if (ax < REDUCE_THRESHOLD) { red = reduce_pio2_fast(ax); }
188
+ else { red = reduce_pio2(ax); }
189
+ let reduced = precise_sincos_half_pi(red.r);
190
+ let sr = reduced.y;
191
+ let cr = reduced.x;
192
+ var sn: f32; var cs: f32;
193
+ switch (red.octant) {
194
+ case 0u: { sn = sr; cs = cr; }
195
+ case 1u: { sn = cr; cs = -sr; }
196
+ case 2u: { sn = -sr; cs = -cr; }
197
+ default: { sn = -cr; cs = sr; }
198
+ }
199
+ let tn = sn / cs;
200
+ return select(tn, -tn, x < 0.0); // tan is odd
201
+ }{% endmacro %}
202
+ {{ emit_is_nan_f32() }}
203
+ {{ emit_trig_reduction_support() }}
204
+
205
+ {{ emit_is_inf_f32() }}
206
+
207
+ {{ emit_nan_from_inf_f32() }}
208
+
209
+ {{ emit_tan_accurate() }}
210
+
211
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
212
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
213
+ // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
214
+ // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
215
+ let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
216
+ if (i >= params.count) {
217
+ return;
218
+ }
219
+ let xv = x[i];
220
+ let fv = vec4<f32>(xv);
221
+ y[i] = {{ vectorScalar }}(vec4<f32>(tan_accurate(fv.x), tan_accurate(fv.y), tan_accurate(fv.z), tan_accurate(fv.w)));
222
+ }