Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Clip

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
b7f95e2
·
verified ·
1 Parent(s): 1eadc80

sync 6fdf6301e2bb

Browse files
Files changed (6) hide show
  1. README.md +2 -2
  2. build/webgpu/bench.json +1 -1
  3. build/webgpu/manifest.json +57 -65
  4. build/webgpu/metadata.json +9 -9
  5. build/webgpu/unary-scalar.wgsl.jinja +8 -14
  6. build/webgpu/unary-vec4.wgsl.jinja +24 -16
README.md CHANGED
@@ -41,14 +41,14 @@ See the [ONNX `Clip` spec](https://onnx.ai/onnx/operators/onnx__Clip.html) for t
41
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
42
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
43
  - [`test.json`](build/webgpu/test.json) — correctness cases
44
- - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
45
  - [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
46
  - [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
47
 
48
  ## Use with `@huggingface/kernels`
49
 
50
  ```sh
51
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
52
  ```
53
 
54
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 
41
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
42
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
43
  - [`test.json`](build/webgpu/test.json) — correctness cases
44
+ - [`bench.json`](build/webgpu/bench.json) — benchmark cases
45
  - [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
46
  - [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
47
 
48
  ## Use with `@huggingface/kernels`
49
 
50
  ```sh
51
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
52
  ```
53
 
54
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
build/webgpu/bench.json CHANGED
@@ -38,7 +38,7 @@
38
  "name": "clip-f32-input-bounds-large-scalar-path",
39
  "provenance": {
40
  "source": "local fixture",
41
- "notes": "The 8,388,608-element float32 input is divisible by four and selects the vectorized input-bounds route."
42
  },
43
  "preset": "stress",
44
  "vars": { "dtype": "float32", "count": 8388608 },
 
38
  "name": "clip-f32-input-bounds-large-scalar-path",
39
  "provenance": {
40
  "source": "local fixture",
41
+ "notes": "An 8,388,608-element float32 input divisible by four with both input bounds present."
42
  },
43
  "preset": "stress",
44
  "vars": { "dtype": "float32", "count": 8388608 },
build/webgpu/manifest.json CHANGED
@@ -19,23 +19,15 @@
19
  "defaultMin": "-65504 if tensorDtypes.input == \"float16\" else (-340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (-128 if tensorDtypes.input == \"int8\" else (-2147483648 if tensorDtypes.input == \"int32\" else 0)))",
20
  "defaultMax": "65504 if tensorDtypes.input == \"float16\" else (340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (127 if tensorDtypes.input == \"int8\" else (2147483647 if tensorDtypes.input == \"int32\" else (255 if tensorDtypes.input == \"uint8\" else 4294967295))))",
21
  "wideVec4StorageOk": "dtypes.T != \"f32\" or device.features.has(\"subgroups\") or not (has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 8 and device.adapterInfo.subgroupMaxSize <= 32)",
22
- "scalar": "dtypes.T",
23
- "usesF16": "dtypes.T == \"f16\""
24
  },
25
  "bindings": {
26
- "x": { "arg": "input", "buffer": "read-only-storage", "elementType": "$vectorScalar" },
27
- "y": { "arg": "output", "buffer": "storage", "elementType": "$vectorScalar" },
28
- "params": {
29
- "buffer": "uniform",
30
- "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
31
- },
32
- "params_2": {
33
- "name": "params",
34
- "buffer": "uniform",
35
- "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
36
- },
37
- "clipMin": { "arg": "min", "buffer": "read-only-storage", "elementType": "$scalar" },
38
- "clipMax": { "arg": "max", "buffer": "read-only-storage", "elementType": "$scalar" }
39
  },
40
  "variants": [
41
  {
@@ -46,7 +38,7 @@
46
  "passes": [
47
  {
48
  "id": "main",
49
- "name": "Clip.defaultBoundsVec4",
50
  "shader": "unary-vec4.wgsl.jinja",
51
  "derive": {
52
  "op": "\"clip\"",
@@ -70,7 +62,7 @@
70
  "passes": [
71
  {
72
  "id": "main",
73
- "name": "Clip.defaultBounds",
74
  "shader": "unary-scalar.wgsl.jinja",
75
  "derive": {
76
  "op": "\"clip\"",
@@ -79,7 +71,7 @@
79
  "maxValue": "defaultMax",
80
  "itemsPerInvocation": 4
81
  },
82
- "bindings": ["input", "output", "params_2"],
83
  "dispatch": {
84
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
85
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -96,14 +88,14 @@
96
  "passes": [
97
  {
98
  "id": "main",
99
- "name": "Clip.inputBoundsVec4",
100
  "shader": "unary-vec4.wgsl.jinja",
101
  "derive": {
102
  "op": "\"clip\"",
103
  "boundsFromInput": true,
 
104
  "clipMinFromInput": true,
105
- "clipMaxFromInput": true,
106
- "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1"
107
  },
108
  "bindings": ["x", "y", "clipMin", "clipMax", "params"],
109
  "dispatch": {
@@ -115,49 +107,48 @@
115
  ]
116
  },
117
  {
118
- "id": "min_input_vec4",
119
- "priority": 40,
120
- "when": ["minInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
121
- "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
122
  "passes": [
123
  {
124
  "id": "main",
125
- "name": "Clip.minInputVec4",
126
- "shader": "unary-vec4.wgsl.jinja",
127
  "derive": {
128
  "op": "\"clip\"",
129
  "boundsFromInput": true,
 
130
  "clipMinFromInput": true,
131
- "clipMaxFromInput": false,
132
- "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1"
133
  },
134
- "bindings": ["x", "y", "clipMin", "params"],
135
  "dispatch": {
136
- "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
137
- "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
138
  "z": 1
139
  }
140
  }
141
- ]
 
142
  },
143
  {
144
- "id": "max_input_vec4",
145
  "priority": 40,
146
- "when": ["maxInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
147
  "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
148
  "passes": [
149
  {
150
  "id": "main",
151
- "name": "Clip.maxInputVec4",
152
  "shader": "unary-vec4.wgsl.jinja",
153
  "derive": {
154
  "op": "\"clip\"",
155
  "boundsFromInput": true,
156
- "clipMinFromInput": false,
157
- "clipMaxFromInput": true,
158
- "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1"
159
  },
160
- "bindings": ["x", "y", "clipMax", "params"],
161
  "dispatch": {
162
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
163
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -167,50 +158,51 @@
167
  ]
168
  },
169
  {
170
- "id": "input_bounds",
171
- "priority": 30,
172
- "when": ["inputBoundsContract"],
173
  "passes": [
174
  {
175
  "id": "main",
176
- "name": "Clip.inputBounds",
177
  "shader": "unary-scalar.wgsl.jinja",
178
  "derive": {
179
  "op": "\"clip\"",
180
  "boundsFromInput": true,
 
181
  "clipMinFromInput": true,
182
- "clipMaxFromInput": true,
183
- "itemsPerInvocation": 4
184
  },
185
- "bindings": ["input", "output", "clipMin", "clipMax", "params_2"],
186
  "dispatch": {
187
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
188
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
189
  "z": 1
190
  }
191
  }
192
- ]
 
193
  },
194
  {
195
- "id": "min_input",
196
- "priority": 30,
197
- "when": ["minInputContract"],
 
198
  "passes": [
199
  {
200
  "id": "main",
201
- "name": "Clip.minInput",
202
- "shader": "unary-scalar.wgsl.jinja",
203
  "derive": {
204
  "op": "\"clip\"",
205
  "boundsFromInput": true,
206
- "clipMinFromInput": true,
207
- "clipMaxFromInput": false,
208
- "itemsPerInvocation": 4
209
  },
210
- "bindings": ["input", "output", "clipMin", "params_2"],
211
  "dispatch": {
212
- "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
213
- "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
214
  "z": 1
215
  }
216
  }
@@ -218,28 +210,28 @@
218
  },
219
  {
220
  "id": "max_input",
221
- "priority": 30,
222
  "when": ["maxInputContract"],
223
  "passes": [
224
  {
225
  "id": "main",
226
- "name": "Clip.maxInput",
227
  "shader": "unary-scalar.wgsl.jinja",
228
  "derive": {
229
  "op": "\"clip\"",
230
  "boundsFromInput": true,
 
231
  "clipMinFromInput": false,
232
- "clipMaxFromInput": true,
233
- "itemsPerInvocation": 4
234
  },
235
- "bindings": ["input", "output", "clipMax", "params_2"],
236
  "dispatch": {
237
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
238
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
239
  "z": 1
240
  }
241
  }
242
- ]
 
243
  }
244
  ]
245
  }
 
19
  "defaultMin": "-65504 if tensorDtypes.input == \"float16\" else (-340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (-128 if tensorDtypes.input == \"int8\" else (-2147483648 if tensorDtypes.input == \"int32\" else 0)))",
20
  "defaultMax": "65504 if tensorDtypes.input == \"float16\" else (340282346638528859811704183484516925440 if tensorDtypes.input == \"float32\" else (127 if tensorDtypes.input == \"int8\" else (2147483647 if tensorDtypes.input == \"int32\" else (255 if tensorDtypes.input == \"uint8\" else 4294967295))))",
21
  "wideVec4StorageOk": "dtypes.T != \"f32\" or device.features.has(\"subgroups\") or not (has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 8 and device.adapterInfo.subgroupMaxSize <= 32)",
22
+ "scalar": "dtypes.T"
 
23
  },
24
  "bindings": {
25
+ "x": { "arg": "input", "elementType": "$vectorScalar" },
26
+ "y": { "arg": "output", "elementType": "$vectorScalar" },
27
+ "params": { "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }] },
28
+ "params_main": { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] },
29
+ "clipMin": { "arg": "min", "elementType": "$scalar" },
30
+ "clipMax": { "arg": "max", "elementType": "$scalar" }
 
 
 
 
 
 
 
31
  },
32
  "variants": [
33
  {
 
38
  "passes": [
39
  {
40
  "id": "main",
41
+ "name": "Clip.DefaultBoundsVec4",
42
  "shader": "unary-vec4.wgsl.jinja",
43
  "derive": {
44
  "op": "\"clip\"",
 
62
  "passes": [
63
  {
64
  "id": "main",
65
+ "name": "Clip.DefaultBounds",
66
  "shader": "unary-scalar.wgsl.jinja",
67
  "derive": {
68
  "op": "\"clip\"",
 
71
  "maxValue": "defaultMax",
72
  "itemsPerInvocation": 4
73
  },
74
+ "bindings": ["input", "output", "params_main"],
75
  "dispatch": {
76
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
77
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
 
88
  "passes": [
89
  {
90
  "id": "main",
91
+ "name": "Clip.InputBoundsVec4",
92
  "shader": "unary-vec4.wgsl.jinja",
93
  "derive": {
94
  "op": "\"clip\"",
95
  "boundsFromInput": true,
96
+ "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1",
97
  "clipMinFromInput": true,
98
+ "clipMaxFromInput": true
 
99
  },
100
  "bindings": ["x", "y", "clipMin", "clipMax", "params"],
101
  "dispatch": {
 
107
  ]
108
  },
109
  {
110
+ "id": "input_bounds",
111
+ "when": ["inputBoundsContract"],
 
 
112
  "passes": [
113
  {
114
  "id": "main",
115
+ "name": "Clip.InputBounds",
116
+ "shader": "unary-scalar.wgsl.jinja",
117
  "derive": {
118
  "op": "\"clip\"",
119
  "boundsFromInput": true,
120
+ "itemsPerInvocation": 4,
121
  "clipMinFromInput": true,
122
+ "clipMaxFromInput": true
 
123
  },
124
+ "bindings": ["input", "output", "clipMin", "clipMax", "params_main"],
125
  "dispatch": {
126
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
127
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
128
  "z": 1
129
  }
130
  }
131
+ ],
132
+ "priority": 30
133
  },
134
  {
135
+ "id": "min_input_vec4",
136
  "priority": 40,
137
+ "when": ["minInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
138
  "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
139
  "passes": [
140
  {
141
  "id": "main",
142
+ "name": "Clip.MinInputVec4",
143
  "shader": "unary-vec4.wgsl.jinja",
144
  "derive": {
145
  "op": "\"clip\"",
146
  "boundsFromInput": true,
147
+ "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1",
148
+ "clipMinFromInput": true,
149
+ "clipMaxFromInput": false
150
  },
151
+ "bindings": ["x", "y", "clipMin", "params"],
152
  "dispatch": {
153
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
154
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
 
158
  ]
159
  },
160
  {
161
+ "id": "min_input",
162
+ "when": ["minInputContract"],
 
163
  "passes": [
164
  {
165
  "id": "main",
166
+ "name": "Clip.MinInput",
167
  "shader": "unary-scalar.wgsl.jinja",
168
  "derive": {
169
  "op": "\"clip\"",
170
  "boundsFromInput": true,
171
+ "itemsPerInvocation": 4,
172
  "clipMinFromInput": true,
173
+ "clipMaxFromInput": false
 
174
  },
175
+ "bindings": ["input", "output", "clipMin", "params_main"],
176
  "dispatch": {
177
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
178
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
179
  "z": 1
180
  }
181
  }
182
+ ],
183
+ "priority": 30
184
  },
185
  {
186
+ "id": "max_input_vec4",
187
+ "priority": 40,
188
+ "when": ["maxInputContract", "numel(shapes.input) > 0", "numel(shapes.input) % 4 == 0", "wideVec4StorageOk"],
189
+ "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
190
  "passes": [
191
  {
192
  "id": "main",
193
+ "name": "Clip.MaxInputVec4",
194
+ "shader": "unary-vec4.wgsl.jinja",
195
  "derive": {
196
  "op": "\"clip\"",
197
  "boundsFromInput": true,
198
+ "vec4PerThread": "4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1",
199
+ "clipMinFromInput": false,
200
+ "clipMaxFromInput": true
201
  },
202
+ "bindings": ["x", "y", "clipMax", "params"],
203
  "dispatch": {
204
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
205
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output) / 4, 4 if numel(shapes.output) * dtypeBytes(tensorDtypes.output) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
206
  "z": 1
207
  }
208
  }
 
210
  },
211
  {
212
  "id": "max_input",
 
213
  "when": ["maxInputContract"],
214
  "passes": [
215
  {
216
  "id": "main",
217
+ "name": "Clip.MaxInput",
218
  "shader": "unary-scalar.wgsl.jinja",
219
  "derive": {
220
  "op": "\"clip\"",
221
  "boundsFromInput": true,
222
+ "itemsPerInvocation": 4,
223
  "clipMinFromInput": false,
224
+ "clipMaxFromInput": true
 
225
  },
226
+ "bindings": ["input", "output", "clipMax", "params_main"],
227
  "dispatch": {
228
  "x": "min(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
229
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.output), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
230
  "z": 1
231
  }
232
  }
233
+ ],
234
+ "priority": 30
235
  }
236
  ]
237
  }
build/webgpu/metadata.json CHANGED
@@ -1,30 +1,30 @@
1
  {
2
  "name": "ai.onnx.Clip",
3
- "id": "_ai_onnx_clip_webgpu_de72431",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "F12x3+s84ktuo167e5515+F2hnUMtcowAYHmEgoDQOg=",
11
- "manifest.json": "kwVE+OfpodbDYkXNlYpHzflNPmAefXobA+ugNpwKfME=",
12
  "test.json": "a604ZK+4B6A13Q4yccOjW3dVe84WzyJ57YjvMttgjuQ=",
13
- "unary-scalar.wgsl.jinja": "rsQQXT2OWXvamkOuvViLDBj3K2kZY/CpmMOvsDNXU7Y=",
14
- "unary-vec4.wgsl.jinja": "/UXGg3mvbuFOUdyjrpM1npD9tJ5fovKFZ6dt0LFgRRc="
15
  }
16
  },
17
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
18
  "webgpu": {
19
- "manifestSpec": "2.0",
20
  "variants": {
21
  "default_bounds_vec4": ["unary-vec4.wgsl.jinja"],
22
  "default_bounds": ["unary-scalar.wgsl.jinja"],
23
  "input_bounds_vec4": ["unary-vec4.wgsl.jinja"],
24
- "min_input_vec4": ["unary-vec4.wgsl.jinja"],
25
- "max_input_vec4": ["unary-vec4.wgsl.jinja"],
26
  "input_bounds": ["unary-scalar.wgsl.jinja"],
 
27
  "min_input": ["unary-scalar.wgsl.jinja"],
 
28
  "max_input": ["unary-scalar.wgsl.jinja"]
29
  }
30
  }
 
1
  {
2
  "name": "ai.onnx.Clip",
3
+ "id": "_ai_onnx_clip_webgpu_027ee2f",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "s28IzawynTG4b4k5Kzojf1kYj4dKCaD3XKquiaTpiEM=",
11
+ "manifest.json": "290wjNBwOPzSyr6Gg5tvm4v+jBicMtbaCGepPyXVEno=",
12
  "test.json": "a604ZK+4B6A13Q4yccOjW3dVe84WzyJ57YjvMttgjuQ=",
13
+ "unary-scalar.wgsl.jinja": "9UgQeCwuevyuggj5HYhjW9ex8KKbD8QIQAFtxJMzAs4=",
14
+ "unary-vec4.wgsl.jinja": "Su4keeFV0z0Yf/X8GHcgvTatrTEc6fqLS2R2/rQCkaA="
15
  }
16
  },
17
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
18
  "webgpu": {
19
+ "manifestSpec": "2.1",
20
  "variants": {
21
  "default_bounds_vec4": ["unary-vec4.wgsl.jinja"],
22
  "default_bounds": ["unary-scalar.wgsl.jinja"],
23
  "input_bounds_vec4": ["unary-vec4.wgsl.jinja"],
 
 
24
  "input_bounds": ["unary-scalar.wgsl.jinja"],
25
+ "min_input_vec4": ["unary-vec4.wgsl.jinja"],
26
  "min_input": ["unary-scalar.wgsl.jinja"],
27
+ "max_input_vec4": ["unary-vec4.wgsl.jinja"],
28
  "max_input": ["unary-scalar.wgsl.jinja"]
29
  }
30
  }
build/webgpu/unary-scalar.wgsl.jinja CHANGED
@@ -12,23 +12,17 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
12
  for (var i = begin; i < end; i = i + 1u) {
13
  {%- endmacro %}
14
  {% macro flat_tail_close() %}
15
- }
16
- {% endmacro %}
17
-
18
- {% set boundsFromInput = boundsFromInput is defined and boundsFromInput %}
19
- // Scalar unary elementwise implementation. Specialization emits only the
20
- // selected operation and any numerical helper it requires.
21
- {% if usesF16 %}
22
- enable f16;
23
- {% endif %}
24
  {{ env.wgsl.resourceDeclarations }}
25
- {% macro emit_is_nan_f32() %}
 
26
  fn is_nan_f32(value: f32) -> bool {
27
  let bits = bitcast<u32>(value);
28
  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
29
- }{% endmacro %}
30
- {% if op == "clip" and scalar != "i32" and scalar != "u32" %}
31
- {{ emit_is_nan_f32() }}
32
  {% endif %}
33
  {{ flat_tail_open() }}
34
  {% set clipMinFromInput = clipMinFromInput | default(false) %}
@@ -67,5 +61,5 @@ fn is_nan_f32(value: f32) -> bool {
67
  }
68
  output[i] = {{ clipScalar }}(out);
69
  {% endif %}
70
- {{ flat_tail_close() -}}
71
  }
 
12
  for (var i = begin; i < end; i = i + 1u) {
13
  {%- endmacro %}
14
  {% macro flat_tail_close() %}
15
+ }{% endmacro %}
16
+ // Scalar unary elementwise kernel: applies the selected operation, with the
17
+ // numerical helpers it uses, to each element.
 
 
 
 
 
 
18
  {{ env.wgsl.resourceDeclarations }}
19
+
20
+ {% if (op == "clip" or op == "relu") and scalar != "i32" and scalar != "u32" %}
21
  fn is_nan_f32(value: f32) -> bool {
22
  let bits = bitcast<u32>(value);
23
  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
24
+ }
25
+
 
26
  {% endif %}
27
  {{ flat_tail_open() }}
28
  {% set clipMinFromInput = clipMinFromInput | default(false) %}
 
61
  }
62
  output[i] = {{ clipScalar }}(out);
63
  {% endif %}
64
+ {{ flat_tail_close() }}
65
  }
build/webgpu/unary-vec4.wgsl.jinja CHANGED
@@ -2,15 +2,14 @@
2
  // component.
3
  {{ env.wgsl.resourceDeclarations }}
4
 
5
- {% macro emit_is_nan_f32() %}
6
  fn is_nan_f32(value: f32) -> bool {
7
  let bits = bitcast<u32>(value);
8
  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
9
- }{% endmacro %}
10
- {% if (op == "clip" and scalar != "i32" and scalar != "u32") or (op == "relu" and scalar != "i32") or op == "sin" or op == "tan" %}
11
- {{ emit_is_nan_f32() }}
12
  {% endif %}
13
- {% if not (boundsFromInput | default(false)) and scalar != "i32" and scalar != "u32" %}
14
  {% set clipMinValue = minValue | default(0) %}
15
  {% set clipMaxValue = maxValue | default(0) %}
16
  fn clip_value(v: f32) -> f32 {
@@ -20,35 +19,44 @@ fn clip_value(v: f32) -> f32 {
20
  }
21
  return out;
22
  }
23
- {% endif %}
24
 
25
- {% set vec4PerThread = vec4PerThread %}
 
 
 
 
 
 
 
 
 
 
26
  {% if vec4PerThread > 1 %}
27
  const ITEMS: u32 = {{ vec4PerThread }}u;
28
- {% endif %}
29
 
 
30
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
31
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
32
- // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
33
- // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
34
  {% if vec4PerThread > 1 %}
35
  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
36
  // access consecutive words on every step, while each lane can keep several
37
  // independent loads in flight.
38
- let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
39
  let span = (params.count + ITEMS - 1u) / ITEMS;
 
 
 
 
 
 
40
  for (var j = 0u; j < ITEMS; j = j + 1u) {
41
  let i = tid + j * span;
42
  if (i >= params.count) {
43
  break;
44
  }
45
  {% else %}
46
- let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
47
- if (i >= params.count) {
48
- return;
49
- }
50
  {% endif %}
51
-
52
  let xv = x[i];
53
  {% set clipMinValue = minValue | default(0) %}
54
  {% set clipMaxValue = maxValue | default(0) %}
 
2
  // component.
3
  {{ env.wgsl.resourceDeclarations }}
4
 
5
+ {% if ((op == "clip" or op == "relu") and scalar != "i32" and scalar != "u32") or op == "sin" or op == "tan" %}
6
  fn is_nan_f32(value: f32) -> bool {
7
  let bits = bitcast<u32>(value);
8
  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
9
+ }
10
+
 
11
  {% endif %}
12
+ {% if op == "clip" and not (boundsFromInput | default(false)) and scalar != "i32" and scalar != "u32" %}
13
  {% set clipMinValue = minValue | default(0) %}
14
  {% set clipMaxValue = maxValue | default(0) %}
15
  fn clip_value(v: f32) -> f32 {
 
19
  }
20
  return out;
21
  }
 
22
 
23
+ {% endif %}
24
+ {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
25
+ {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
26
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
27
+ // per-axis workgroup fold width.
28
+ {% if bound == "" %}
29
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% else %}
30
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
31
+ if ({{ name }} >= {{ bound }}) {
32
+ return;
33
+ }{% endif %}{% endmacro %}
34
  {% if vec4PerThread > 1 %}
35
  const ITEMS: u32 = {{ vec4PerThread }}u;
 
36
 
37
+ {% endif %}
38
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
39
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 
 
40
  {% if vec4PerThread > 1 %}
41
  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
42
  // access consecutive words on every step, while each lane can keep several
43
  // independent loads in flight.
44
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE, "tid", "") }}
45
  let span = (params.count + ITEMS - 1u) / ITEMS;
46
+ // Lanes at or past span exist only because the dispatch rounds up to whole
47
+ // workgroups. Lane span + k would start on lane k's second group and rewrite
48
+ // up to ITEMS - 1 groups another lane already stored.
49
+ if (tid >= span) {
50
+ return;
51
+ }
52
  for (var j = 0u; j < ITEMS; j = j + 1u) {
53
  let i = tid + j * span;
54
  if (i >= params.count) {
55
  break;
56
  }
57
  {% else %}
58
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE) }}
 
 
 
59
  {% endif %}
 
60
  let xv = x[i];
61
  {% set clipMinValue = minValue | default(0) %}
62
  {% set clipMaxValue = maxValue | default(0) %}