Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.HammingWindow

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
d68c445
·
verified ·
1 Parent(s): 76cd045

sync 6fdf6301e2bb

Browse files
Files changed (6) hide show
  1. README.md +2 -2
  2. build/webgpu/bench.json +33 -0
  3. build/webgpu/manifest.json +6 -10
  4. build/webgpu/metadata.json +7 -7
  5. build/webgpu/test.json +11 -2
  6. build/webgpu/window-cosine-sum.wgsl.jinja +3 -3
README.md CHANGED
@@ -49,13 +49,13 @@ Default values (overridable per request):
49
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
50
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
51
  - [`test.json`](build/webgpu/test.json) — correctness cases
52
- - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
53
  - [`window-cosine-sum.wgsl.jinja`](build/webgpu/window-cosine-sum.wgsl.jinja)
54
 
55
  ## Use with `@huggingface/kernels`
56
 
57
  ```sh
58
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
59
  ```
60
 
61
  Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
 
49
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
50
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
51
  - [`test.json`](build/webgpu/test.json) — correctness cases
52
+ - [`bench.json`](build/webgpu/bench.json) — benchmark cases
53
  - [`window-cosine-sum.wgsl.jinja`](build/webgpu/window-cosine-sum.wgsl.jinja)
54
 
55
  ## Use with `@huggingface/kernels`
56
 
57
  ```sh
58
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
59
  ```
60
 
61
  Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
build/webgpu/bench.json CHANGED
@@ -19,6 +19,39 @@
19
  "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1048576] } } },
20
  "outputs": { "y": { "dtype": "float32", "shape": [1048576], "dist": "empty" } },
21
  "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.y) * 4", "name": "write-BW" }] }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
22
  }
23
  ]
24
  }
 
19
  "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1048576] } } },
20
  "outputs": { "y": { "dtype": "float32", "shape": [1048576], "dist": "empty" } },
21
  "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.y) * 4", "name": "write-BW" }] }
22
+ },
23
+ {
24
+ "name": "hammingwindow-float32-size1048583-periodic1",
25
+ "preset": "edge",
26
+ "attrs": { "periodic": 1, "output_datatype": 1 },
27
+ "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1048583] } } },
28
+ "outputs": { "y": { "dtype": "float32", "shape": [1048583] } },
29
+ "bench": { "metrics": [{ "type": "bandwidth", "name": "write-BW", "value": "numel(shapes.y) * 4" }] },
30
+ "provenance": {
31
+ "notes": "Rational-period window with a non-power-of-two length, an output tail, or a period beyond exact float32 integer representation."
32
+ }
33
+ },
34
+ {
35
+ "name": "hammingwindow-float32-size1048577-periodic0",
36
+ "preset": "edge",
37
+ "attrs": { "periodic": 0, "output_datatype": 1 },
38
+ "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1048577] } } },
39
+ "outputs": { "y": { "dtype": "float32", "shape": [1048577] } },
40
+ "bench": { "metrics": [{ "type": "bandwidth", "name": "write-BW", "value": "numel(shapes.y) * 4" }] },
41
+ "provenance": {
42
+ "notes": "Rational-period window with a non-power-of-two length, an output tail, or a period beyond exact float32 integer representation."
43
+ }
44
+ },
45
+ {
46
+ "name": "hammingwindow-float16-size65537-periodic0",
47
+ "preset": "edge",
48
+ "attrs": { "periodic": 0, "output_datatype": 10 },
49
+ "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [65537] } } },
50
+ "outputs": { "y": { "dtype": "float16", "shape": [65537] } },
51
+ "bench": { "metrics": [{ "type": "bandwidth", "name": "write-BW", "value": "numel(shapes.y) * 2" }] },
52
+ "provenance": {
53
+ "notes": "Rational-period window with a non-power-of-two length, an output tail, or a period beyond exact float32 integer representation."
54
+ }
55
  }
56
  ]
57
  }
build/webgpu/manifest.json CHANGED
@@ -14,7 +14,8 @@
14
  "when": ["ranks.y == 1 and attrs.output_datatype == onnxDtypeCode(logicalDtypes.T2) and f16Ok(tensorDtypes.y)"],
15
  "derive": {
16
  "integerOutput": "tensorDtypes.y != \"float32\" and tensorDtypes.y != \"float16\"",
17
- "outScalar": "dtypes.T2"
 
18
  },
19
  "passes": [
20
  {
@@ -28,21 +29,16 @@
28
  "name": "params",
29
  "struct": [
30
  { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
31
- {
32
- "name": "denom",
33
- "type": "f32",
34
- "value": "dim(shapes.y, 0) if attrs.periodic else dim(shapes.y, 0) - 1"
35
- }
36
  ]
37
  }
38
  ],
39
  "dispatch": {
40
- "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
41
- "y": 1,
42
- "z": 1
43
  }
44
  }
45
  ]
46
  }
47
- ]
 
48
  }
 
14
  "when": ["ranks.y == 1 and attrs.output_datatype == onnxDtypeCode(logicalDtypes.T2) and f16Ok(tensorDtypes.y)"],
15
  "derive": {
16
  "integerOutput": "tensorDtypes.y != \"float32\" and tensorDtypes.y != \"float16\"",
17
+ "outScalar": "dtypes.T2",
18
+ "denom": "dim(shapes.y, 0) if attrs.periodic else dim(shapes.y, 0) - 1"
19
  },
20
  "passes": [
21
  {
 
29
  "name": "params",
30
  "struct": [
31
  { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
32
+ { "name": "denom", "type": "f32", "value": "denom" }
 
 
 
 
33
  ]
34
  }
35
  ],
36
  "dispatch": {
37
+ "x": "min(ceilDiv(numel(shapes.y), tunables.WORKGROUP_SIZE), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
 
 
38
  }
39
  }
40
  ]
41
  }
42
+ ],
43
+ "bindings": {}
44
  }
build/webgpu/metadata.json CHANGED
@@ -1,18 +1,18 @@
1
  {
2
  "name": "ai.onnx.HammingWindow",
3
- "id": "_ai_onnx_hammingwindow_webgpu_0636ea9",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "Cc46HEb4s6IKKoK8SD2mKgLqX2TlsMRg2ZMdv2BsYfo=",
11
- "manifest.json": "T+EIWkHDvr76V02BuqLKZOdO+vggDLEjVOH7vkB8qmw=",
12
- "test.json": "E54R+4kVhcW3he06McOJRuVw0a6eReixIbiuLojMX4U=",
13
- "window-cosine-sum.wgsl.jinja": "fzmuM0dllTv+JLWkn5ADVz1ixY31R6N0IVyyJGutuls="
14
  }
15
  },
16
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
17
- "webgpu": { "manifestSpec": "2.0", "variants": { "generic": ["window-cosine-sum.wgsl.jinja"] } }
18
  }
 
1
  {
2
  "name": "ai.onnx.HammingWindow",
3
+ "id": "_ai_onnx_hammingwindow_webgpu_92e6839",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "rq1fsSujNDEu/bTsI7EzNk7jC3ZKHOll8prqUEfSwQs=",
11
+ "manifest.json": "PKKBP/gSN516qO7APeMiYAIKB+Gb4VcLgOl7utG8/yk=",
12
+ "test.json": "jyWr72jJNClbCf/VfsHjwjxlEabO1EXnX5FKx1NJ/9s=",
13
+ "window-cosine-sum.wgsl.jinja": "R1t8x/hjbgdPE61JguNx7bvv5wak7wUuLGkPWFIlbjg="
14
  }
15
  },
16
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
17
+ "webgpu": { "manifestSpec": "2.1", "variants": { "generic": ["window-cosine-sum.wgsl.jinja"] } }
18
  }
build/webgpu/test.json CHANGED
@@ -68,7 +68,7 @@
68
  {
69
  "name": "ort_periodic_int16_output_size8",
70
  "provenance": {
71
- "source": "ONNX Runtime CPUExecutionProvider",
72
  "notes": "ONNX HammingWindow computes in float and applies output_datatype=INT16 as a final cast."
73
  },
74
  "attrs": { "output_datatype": 5 },
@@ -85,7 +85,7 @@
85
  {
86
  "name": "ort_symmetric_uint32_output_size1_zero_cast",
87
  "provenance": {
88
- "source": "ONNX Runtime CPUExecutionProvider",
89
  "notes": "The symmetric size-one formula produces NaN in float; ONNX Cast to UINT32 maps that value to zero."
90
  },
91
  "attrs": { "output_datatype": 12, "periodic": 0 },
@@ -198,6 +198,15 @@
198
  "requires": { "features": ["shader-f16"] },
199
  "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16] } } },
200
  "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0.002 } }
 
 
 
 
 
 
 
 
 
201
  }
202
  ]
203
  }
 
68
  {
69
  "name": "ort_periodic_int16_output_size8",
70
  "provenance": {
71
+ "source": "ONNX Runtime's CPU provider",
72
  "notes": "ONNX HammingWindow computes in float and applies output_datatype=INT16 as a final cast."
73
  },
74
  "attrs": { "output_datatype": 5 },
 
85
  {
86
  "name": "ort_symmetric_uint32_output_size1_zero_cast",
87
  "provenance": {
88
+ "source": "ONNX Runtime's CPU provider",
89
  "notes": "The symmetric size-one formula produces NaN in float; ONNX Cast to UINT32 maps that value to zero."
90
  },
91
  "attrs": { "output_datatype": 12, "periodic": 0 },
 
198
  "requires": { "features": ["shader-f16"] },
199
  "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16] } } },
200
  "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0.002 } }
201
+ },
202
+ {
203
+ "name": "periodic_f32_clamped_dispatch_65537",
204
+ "tunables": { "WORKGROUP_SIZE": 1 },
205
+ "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [65537] } } },
206
+ "outputs": { "y": { "dtype": "float32", "shape": [65537], "tolerance": 0.00001 } },
207
+ "provenance": {
208
+ "notes": "A non-power-of-two period whose one-lane workgroups reach the dispatch cap and reuse the first invocations for the final two output elements."
209
+ }
210
  }
211
  ]
212
  }
build/webgpu/window-cosine-sum.wgsl.jinja CHANGED
@@ -9,12 +9,14 @@
9
  // implementation-defined, and a rounded multiply followed by subtraction would
10
  // double-round. Grid strides cover large windows with a clamped dispatch.
11
  {{ env.wgsl.resourceDeclarations }}
 
12
  // Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
13
  // rational turn and route polynomial evaluation through one trigonometric core
14
  // for FFT, direct DFT, STFT, and cosine-window calculations.
15
  //
16
  // phase is in [0, period). sign is -1 for a forward transform and +1 for an
17
  // inverse transform. The result is (cos, sin).
 
18
  // Backend-stable f32 sine/cosine core.
19
  //
20
  // Shader transcendental accuracy is implementation-defined. Each path retains
@@ -63,7 +65,6 @@ fn precise_sincos_centered(x: f32) -> vec2<f32> {
63
  return vec2<f32>(cosineSign * value.x, value.y);
64
  }
65
 
66
-
67
  fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
68
  let reducedPhase = phase % period;
69
  if (reducedPhase == 0u) {
@@ -88,11 +89,10 @@ fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
88
  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
89
  }
90
 
91
-
92
  const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
93
 
94
  fn cos_turns(phase: u32, period: u32) -> f32 {
95
- return dft_precise_twiddle(phase % period, period, 1.0).x;
96
  }
97
 
98
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 
9
  // implementation-defined, and a rounded multiply followed by subtraction would
10
  // double-round. Grid strides cover large windows with a clamped dispatch.
11
  {{ env.wgsl.resourceDeclarations }}
12
+
13
  // Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
14
  // rational turn and route polynomial evaluation through one trigonometric core
15
  // for FFT, direct DFT, STFT, and cosine-window calculations.
16
  //
17
  // phase is in [0, period). sign is -1 for a forward transform and +1 for an
18
  // inverse transform. The result is (cos, sin).
19
+
20
  // Backend-stable f32 sine/cosine core.
21
  //
22
  // Shader transcendental accuracy is implementation-defined. Each path retains
 
65
  return vec2<f32>(cosineSign * value.x, value.y);
66
  }
67
 
 
68
  fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
69
  let reducedPhase = phase % period;
70
  if (reducedPhase == 0u) {
 
89
  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
90
  }
91
 
 
92
  const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
93
 
94
  fn cos_turns(phase: u32, period: u32) -> f32 {
95
+ return dft_precise_twiddle(phase, period, 1.0).x;
96
  }
97
 
98
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})