Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Col2Im

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
980291f
·
verified ·
1 Parent(s): a7f0ece

sync 6fdf6301e2bb

Browse files
Files changed (5) hide show
  1. README.md +2 -2
  2. build/webgpu/bench.json +1 -1
  3. build/webgpu/col2im-nd.wgsl.jinja +5 -32
  4. build/webgpu/manifest.json +7 -7
  5. build/webgpu/metadata.json +6 -6
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
  - [`col2im-nd.wgsl.jinja`](build/webgpu/col2im-nd.wgsl.jinja)
64
  - [`datamove-elementwise-copy.wgsl.jinja`](build/webgpu/datamove-elementwise-copy.wgsl.jinja)
65
  - [`datamove-flat-copy.wgsl.jinja`](build/webgpu/datamove-flat-copy.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
  - [`col2im-nd.wgsl.jinja`](build/webgpu/col2im-nd.wgsl.jinja)
64
  - [`datamove-elementwise-copy.wgsl.jinja`](build/webgpu/datamove-elementwise-copy.wgsl.jinja)
65
  - [`datamove-flat-copy.wgsl.jinja`](build/webgpu/datamove-flat-copy.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, 16, 32, 32] } }
13
  },
14
  {
15
- "name": "nchw_4m_k3_overlap_healthy",
16
  "preset": "smoke",
17
  "attrs": {},
18
  "inputs": {
 
12
  "outputs": { "output": { "dtype": "float32", "shape": [1, 16, 32, 32] } }
13
  },
14
  {
15
+ "name": "nchw_4m_k3_overlap_control",
16
  "preset": "smoke",
17
  "attrs": {},
18
  "inputs": {
build/webgpu/col2im-nd.wgsl.jinja CHANGED
@@ -1,36 +1,9 @@
1
- {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
2
- {% if note == "dispatch-limit" %}
3
- // 2D-folded flat index: gid.y carries the high bits past the dispatch's
4
- // per-axis workgroup fold width (outputs > 16.7M elements).
5
- {% elif note == "limit" %}
6
- // 2D-folded flat index: gid.y carries the high bits past the dispatch's
7
- // per-axis workgroup fold width.
8
- {% elif note == "device-axis" %}
9
- // The flat dispatch is folded across x/y at a fixed per-axis workgroup
10
- // width; gid.y carries the high portion of the output index.
11
- {% elif note == "vec4-limit" %}
12
- // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
13
- // per-axis workgroup fold width (the dispatch caps x and spills into y).
14
- {% elif note == "element-limit" %}
15
- // 2D-folded flat element index: gid.y carries the high bits past the
16
- // dispatch's per-axis workgroup fold width.
17
- {% elif note == "dispatch" %}
18
  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
19
  // per-axis workgroup fold width.
20
- {% endif %}
21
- {% if bound == "" %}
22
- let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
23
- {%- elif guardInline %}
24
- let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
25
- if ({{ name }} >= {{ bound }}) { return; }
26
- {%- else %}
27
- let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
28
- if ({{ name }} >= {{ bound }}) {
29
- return;
30
- }
31
- {%- endif %}
32
- {% endmacro %}
33
-
34
  {{ env.wgsl.resourceDeclarations }}
35
 
36
  // Image geometry, strides, pads, and dilations are shader constants, so the
@@ -63,7 +36,7 @@ fn divisible(value: i32, divisor: u32) -> bool {
63
 
64
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
65
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
66
- {{ flat_index_2d("index", guardInline=true, note="") }}
67
 
68
  {% if spatialRankSpec == 3 %}
69
  let blockD = block_shape[0];
 
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 }}) { return; }{% endmacro %}
 
 
 
 
 
 
 
 
 
 
 
 
7
  {{ env.wgsl.resourceDeclarations }}
8
 
9
  // Image geometry, strides, pads, and dilations are shader constants, so the
 
36
 
37
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
38
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
39
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE, "index", guardInline=true) }}
40
 
41
  {% if spatialRankSpec == 3 %}
42
  let blockD = block_shape[0];
build/webgpu/manifest.json CHANGED
@@ -33,11 +33,11 @@
33
  },
34
  "when": ["colBaseContract"],
35
  "bindings": {
36
- "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] },
37
- "input": { "buffer": "read-only-storage", "elementType": "f32" },
38
- "output": { "buffer": "storage", "elementType": "f32" },
39
- "image_shape": { "buffer": "read-only-storage", "elementType": "u32" },
40
- "block_shape": { "buffer": "read-only-storage", "elementType": "u32" }
41
  },
42
  "variants": [
43
  {
@@ -48,7 +48,7 @@
48
  "passes": [
49
  {
50
  "id": "main",
51
- "name": "Col2Im.identityCopyVec4",
52
  "shader": "datamove-flat-copy.wgsl.jinja",
53
  "derive": { "count": "numel(shapes.output) / 4" },
54
  "bindings": [
@@ -72,7 +72,7 @@
72
  "passes": [
73
  {
74
  "id": "main",
75
- "name": "Col2Im.identityCopy",
76
  "shader": "datamove-elementwise-copy.wgsl.jinja",
77
  "bindings": [
78
  { "arg": "input", "name": "x", "elementType": "$scalar" },
 
33
  },
34
  "when": ["colBaseContract"],
35
  "bindings": {
36
+ "params": { "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] },
37
+ "input": { "elementType": "f32" },
38
+ "output": { "elementType": "f32" },
39
+ "image_shape": { "elementType": "u32" },
40
+ "block_shape": { "elementType": "u32" }
41
  },
42
  "variants": [
43
  {
 
48
  "passes": [
49
  {
50
  "id": "main",
51
+ "name": "Col2Im.IdentityCopyVec4",
52
  "shader": "datamove-flat-copy.wgsl.jinja",
53
  "derive": { "count": "numel(shapes.output) / 4" },
54
  "bindings": [
 
72
  "passes": [
73
  {
74
  "id": "main",
75
+ "name": "Col2Im.IdentityCopy",
76
  "shader": "datamove-elementwise-copy.wgsl.jinja",
77
  "bindings": [
78
  { "arg": "input", "name": "x", "elementType": "$scalar" },
build/webgpu/metadata.json CHANGED
@@ -1,23 +1,23 @@
1
  {
2
  "name": "ai.onnx.Col2Im",
3
- "id": "_ai_onnx_col2im_webgpu_3a8a360",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "HtzA9ij1QBc6FxpgD06P7vnwZ0c1nPbhkIqeLBGOBSg=",
11
- "col2im-nd.wgsl.jinja": "dLnJ/gW2AYMMgEZUHSce6p6tJ9UwLxOiaViX+TvdUXM=",
12
  "datamove-elementwise-copy.wgsl.jinja": "Q1WuZCqcDf9b6cbT4rLilpeWj6327t/bSFmKr6IVBoQ=",
13
  "datamove-flat-copy.wgsl.jinja": "g9d62mer2bmHfbScqX5CIT5Zq/Pdinac0W2dH0CD+5s=",
14
- "manifest.json": "kWA2otoEQ/KiG6nDA8on3UQ9faEaSYg/RAFK2j9UDPM=",
15
  "test.json": "9QFgftL+lxg0S+3h+8siyB5eqkKVj6dLkVgG9dlpJbk="
16
  }
17
  },
18
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
19
  "webgpu": {
20
- "manifestSpec": "2.0",
21
  "variants": {
22
  "identity_copy_vec4": ["datamove-flat-copy.wgsl.jinja"],
23
  "identity_copy": ["datamove-elementwise-copy.wgsl.jinja"],
 
1
  {
2
  "name": "ai.onnx.Col2Im",
3
+ "id": "_ai_onnx_col2im_webgpu_27afd68",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "hE8h1BZxfWU3ROBUVeV838U2jMfYtiRkHPO9/PjoRR0=",
11
+ "col2im-nd.wgsl.jinja": "WbeYdOoXPry5JgxMiHeda646DZ9UY7geW5bs81DssJo=",
12
  "datamove-elementwise-copy.wgsl.jinja": "Q1WuZCqcDf9b6cbT4rLilpeWj6327t/bSFmKr6IVBoQ=",
13
  "datamove-flat-copy.wgsl.jinja": "g9d62mer2bmHfbScqX5CIT5Zq/Pdinac0W2dH0CD+5s=",
14
+ "manifest.json": "u9M5ClP4qQgkLlrJKWhZoB0m7tecRYc5X+Axgmrltl0=",
15
  "test.json": "9QFgftL+lxg0S+3h+8siyB5eqkKVj6dLkVgG9dlpJbk="
16
  }
17
  },
18
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
19
  "webgpu": {
20
+ "manifestSpec": "2.1",
21
  "variants": {
22
  "identity_copy_vec4": ["datamove-flat-copy.wgsl.jinja"],
23
  "identity_copy": ["datamove-elementwise-copy.wgsl.jinja"],