Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.MaxUnpool

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
7f1e571
·
verified ·
1 Parent(s): e9991d9

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 +3 -4
  4. build/webgpu/maxunpool-elect.wgsl.jinja +9 -4
  5. build/webgpu/metadata.json +7 -7
  6. build/webgpu/test.json +1 -1
README.md CHANGED
@@ -59,7 +59,7 @@ One implementation is selected per call from the device capabilities, the reques
59
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
60
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
61
  - [`test.json`](build/webgpu/test.json) — correctness cases
62
- - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
63
  - [`maxunpool-elect.wgsl.jinja`](build/webgpu/maxunpool-elect.wgsl.jinja)
64
  - [`maxunpool-gather.wgsl.jinja`](build/webgpu/maxunpool-gather.wgsl.jinja)
65
  - [`maxunpool-zerofill.wgsl.jinja`](build/webgpu/maxunpool-zerofill.wgsl.jinja)
@@ -67,7 +67,7 @@ One implementation is selected per call from the device capabilities, the reques
67
  ## Use with `@huggingface/kernels`
68
 
69
  ```sh
70
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
71
  ```
72
 
73
  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.
 
59
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
60
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
61
  - [`test.json`](build/webgpu/test.json) — correctness cases
62
+ - [`bench.json`](build/webgpu/bench.json) — benchmark cases
63
  - [`maxunpool-elect.wgsl.jinja`](build/webgpu/maxunpool-elect.wgsl.jinja)
64
  - [`maxunpool-gather.wgsl.jinja`](build/webgpu/maxunpool-gather.wgsl.jinja)
65
  - [`maxunpool-zerofill.wgsl.jinja`](build/webgpu/maxunpool-zerofill.wgsl.jinja)
 
67
  ## Use with `@huggingface/kernels`
68
 
69
  ```sh
70
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
71
  ```
72
 
73
  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
@@ -12,7 +12,7 @@
12
  "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 128, 128] } }
13
  },
14
  {
15
- "name": "f16_dense_2x_unpool_1x16x128x128_unmeasured",
16
  "preset": "smoke",
17
  "attrs": { "kernel_shape": [2, 2], "strides": [2, 2] },
18
  "inputs": {
 
12
  "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 128, 128] } }
13
  },
14
  {
15
+ "name": "f16_dense_2x_unpool_1x16x128x128",
16
  "preset": "smoke",
17
  "attrs": { "kernel_shape": [2, 2], "strides": [2, 2] },
18
  "inputs": {
build/webgpu/manifest.json CHANGED
@@ -23,10 +23,9 @@
23
  "winnerElement": "\"vec4<u32>\" if outputVec4 else \"u32\""
24
  },
25
  "bindings": {
26
- "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] },
27
- "params_2": {
28
  "name": "params",
29
- "buffer": "uniform",
30
  "struct": [
31
  { "name": "count", "type": "u32", "value": "numel(shapes.x)" },
32
  { "name": "outCount", "type": "u32", "value": "numel(shapes.output)" }
@@ -73,7 +72,7 @@
73
  "id": "elect",
74
  "name": "MaxUnpool.Elect",
75
  "shader": "maxunpool-elect.wgsl.jinja",
76
- "bindings": ["indices", { "name": "winner", "buffer": "storage", "elementType": "atomic<u32>" }, "params_2"],
77
  "dispatch": {
78
  "x": "min(ceilDiv((numel(shapes.x)), (tunables.WORKGROUP_SIZE)), 65535)",
79
  "y": "ceilDiv(ceilDiv((numel(shapes.x)), (tunables.WORKGROUP_SIZE)), 65535)",
 
23
  "winnerElement": "\"vec4<u32>\" if outputVec4 else \"u32\""
24
  },
25
  "bindings": {
26
+ "params": { "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] },
27
+ "params_elect": {
28
  "name": "params",
 
29
  "struct": [
30
  { "name": "count", "type": "u32", "value": "numel(shapes.x)" },
31
  { "name": "outCount", "type": "u32", "value": "numel(shapes.output)" }
 
72
  "id": "elect",
73
  "name": "MaxUnpool.Elect",
74
  "shader": "maxunpool-elect.wgsl.jinja",
75
+ "bindings": ["indices", { "name": "winner", "elementType": "atomic<u32>" }, "params_elect"],
76
  "dispatch": {
77
  "x": "min(ceilDiv((numel(shapes.x)), (tunables.WORKGROUP_SIZE)), 65535)",
78
  "y": "ceilDiv(ceilDiv((numel(shapes.x)), (tunables.WORKGROUP_SIZE)), 65535)",
build/webgpu/maxunpool-elect.wgsl.jinja CHANGED
@@ -1,3 +1,11 @@
 
 
 
 
 
 
 
 
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
  // Elect one writer per destination. Overlapping pooling windows can name the
@@ -5,10 +13,7 @@
5
  // elected index is stored as index + 1, leaving zero for "no writer".
6
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
7
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
8
- let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
9
- if (i >= params.count) {
10
- return;
11
- }
12
  let dst = indices[i];
13
  if (dst < params.outCount) {
14
  atomicMax(&winner[dst], i + 1u);
 
1
+ {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
2
+ {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
3
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
4
+ // per-axis workgroup fold width.
5
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
6
+ if ({{ name }} >= {{ bound }}) {
7
+ return;
8
+ }{% endmacro %}
9
  {{ env.wgsl.resourceDeclarations }}
10
 
11
  // Elect one writer per destination. Overlapping pooling windows can name the
 
13
  // elected index is stored as index + 1, leaving zero for "no writer".
14
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
15
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
16
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE) }}
 
 
 
17
  let dst = indices[i];
18
  if (dst < params.outCount) {
19
  atomicMax(&winner[dst], i + 1u);
build/webgpu/metadata.json CHANGED
@@ -1,23 +1,23 @@
1
  {
2
  "name": "ai.onnx.MaxUnpool",
3
- "id": "_ai_onnx_maxunpool_webgpu_36f6795",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "9yCIOHgrjtcYoRPoEyS6pMXvgi5T4Swd9uj13r8aAb0=",
11
- "manifest.json": "/8AQvbb67ZIIJVY8MDN7oU9dhHYoIVxVLNuHyONjOYY=",
12
- "maxunpool-elect.wgsl.jinja": "RyImAcLUoLZ321Q7rbqsT6SuWuogJixMTlrNxVXZfkU=",
13
  "maxunpool-gather.wgsl.jinja": "7RCk6AvfAGxYgwaL6MdR0KKy7vceCqP3OYd0CYI362Q=",
14
  "maxunpool-zerofill.wgsl.jinja": "WG9FDOFO0Y1MDxrU95AyYYP/uvx+IEQ570iLHBzsbZU=",
15
- "test.json": "K6KpsyfyLT2UekgN7enhBp82RJnn3Ou09Rx9Dq3mj1w="
16
  }
17
  },
18
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
19
  "webgpu": {
20
- "manifestSpec": "2.0",
21
  "variants": {
22
  "empty": ["maxunpool-zerofill.wgsl.jinja"],
23
  "generic": ["maxunpool-elect.wgsl.jinja", "maxunpool-gather.wgsl.jinja", "maxunpool-zerofill.wgsl.jinja"]
 
1
  {
2
  "name": "ai.onnx.MaxUnpool",
3
+ "id": "_ai_onnx_maxunpool_webgpu_a88d7ae",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "MIhNaEcYHUvQ4bjoppnCFyf4PhXPEZfECfgJaD5dpFs=",
11
+ "manifest.json": "FiWAOc8WJuNIOANKVJKnls5Djc17c8UUodp++hDFBiA=",
12
+ "maxunpool-elect.wgsl.jinja": "co6TIrJKCiL03xOfcuRd7VqDMKQszx8+YbfM8eUaQOA=",
13
  "maxunpool-gather.wgsl.jinja": "7RCk6AvfAGxYgwaL6MdR0KKy7vceCqP3OYd0CYI362Q=",
14
  "maxunpool-zerofill.wgsl.jinja": "WG9FDOFO0Y1MDxrU95AyYYP/uvx+IEQ570iLHBzsbZU=",
15
+ "test.json": "/DrtR6K4Ivz6sxabQikXxCM6ZxT2v69lq2vQgGlWyrU="
16
  }
17
  },
18
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
19
  "webgpu": {
20
+ "manifestSpec": "2.1",
21
  "variants": {
22
  "empty": ["maxunpool-zerofill.wgsl.jinja"],
23
  "generic": ["maxunpool-elect.wgsl.jinja", "maxunpool-gather.wgsl.jinja", "maxunpool-zerofill.wgsl.jinja"]
build/webgpu/test.json CHANGED
@@ -14,7 +14,7 @@
14
  "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 4097, 4097], "tolerance": 0 } },
15
  "tunables": { "WORKGROUP_SIZE": 64 },
16
  "provenance": {
17
- "notes": "The 64-thread workgroup makes both four-output clearing/publication and per-input election cross the dispatch fold, exercising each pass on the existing large shape."
18
  }
19
  },
20
  {
 
14
  "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 4097, 4097], "tolerance": 0 } },
15
  "tunables": { "WORKGROUP_SIZE": 64 },
16
  "provenance": {
17
+ "notes": "The 64-thread workgroup makes both four-output clearing/publication and per-input election cross the dispatch fold, exercising each pass on a large shape."
18
  }
19
  },
20
  {