Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.BitwiseNot

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

sync 6fdf6301e2bb

Browse files
Files changed (4) hide show
  1. README.md +2 -2
  2. build/webgpu/bitwise-not.wgsl.jinja +5 -30
  3. build/webgpu/manifest.json +1 -1
  4. build/webgpu/metadata.json +5 -5
README.md CHANGED
@@ -39,13 +39,13 @@ See the [ONNX `BitwiseNot` spec](https://onnx.ai/onnx/operators/onnx__BitwiseNot
39
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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
  - [`bitwise-not.wgsl.jinja`](build/webgpu/bitwise-not.wgsl.jinja)
44
 
45
  ## Use with `@huggingface/kernels`
46
 
47
  ```sh
48
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
49
  ```
50
 
51
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 
39
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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 cases
43
  - [`bitwise-not.wgsl.jinja`](build/webgpu/bitwise-not.wgsl.jinja)
44
 
45
  ## Use with `@huggingface/kernels`
46
 
47
  ```sh
48
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
49
  ```
50
 
51
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
build/webgpu/bitwise-not.wgsl.jinja CHANGED
@@ -1,41 +1,16 @@
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
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
37
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
38
- {{ flat_index_2d(note="element-limit") }}
39
  var value = ~x[i];
40
  {% if logicalDtype == "uint8" %}
41
  {% if vectorizedSpec %}
 
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
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
12
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
13
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE) }}
14
  var value = ~x[i];
15
  {% if logicalDtype == "uint8" %}
16
  {% if vectorizedSpec %}
build/webgpu/manifest.json CHANGED
@@ -17,7 +17,7 @@
17
  "passes": [
18
  {
19
  "id": "main",
20
- "name": "BitwiseNot.vec4",
21
  "shader": "bitwise-not.wgsl.jinja",
22
  "derive": { "vectorizedSpec": true },
23
  "bindings": [
 
17
  "passes": [
18
  {
19
  "id": "main",
20
+ "name": "BitwiseNot.Vec4",
21
  "shader": "bitwise-not.wgsl.jinja",
22
  "derive": { "vectorizedSpec": true },
23
  "bindings": [
build/webgpu/metadata.json CHANGED
@@ -1,6 +1,6 @@
1
  {
2
  "name": "ai.onnx.BitwiseNot",
3
- "id": "_ai_onnx_bitwisenot_webgpu_853bcee",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
@@ -8,14 +8,14 @@
8
  "algorithm": "sha256",
9
  "files": {
10
  "bench.json": "MD1dhXerFPrKL1pwLulPwQoIcIJYxQRi6DZEqLmoETg=",
11
- "bitwise-not.wgsl.jinja": "ZtA3TBUEhXt39AUn36XrrbII2ndOJO6W+D3gVBiSMp8=",
12
- "manifest.json": "ifbhqCLw5vH+vD1dd5GtcdijLT+JDK1N98R/WnGEiq8=",
13
  "test.json": "o049sQp+WvVyGnQxzbNjh0At7j0Yn/3g6v73ROxa1iY="
14
  }
15
  },
16
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
17
  "webgpu": {
18
- "manifestSpec": "2.0",
19
  "variants": { "elementwise_vec4": ["bitwise-not.wgsl.jinja"], "elementwise": ["bitwise-not.wgsl.jinja"] }
20
  }
21
  }
 
1
  {
2
  "name": "ai.onnx.BitwiseNot",
3
+ "id": "_ai_onnx_bitwisenot_webgpu_92dddde",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
 
8
  "algorithm": "sha256",
9
  "files": {
10
  "bench.json": "MD1dhXerFPrKL1pwLulPwQoIcIJYxQRi6DZEqLmoETg=",
11
+ "bitwise-not.wgsl.jinja": "ZcwXk4pAdCTO0djDhsPjoCuRxNcKK48wEIZsj3bBtbo=",
12
+ "manifest.json": "kZR42iResDgXQxssUHM4EwXRC+UoNGnRWgYInybvm5A=",
13
  "test.json": "o049sQp+WvVyGnQxzbNjh0At7j0Yn/3g6v73ROxa1iY="
14
  }
15
  },
16
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
17
  "webgpu": {
18
+ "manifestSpec": "2.1",
19
  "variants": { "elementwise_vec4": ["bitwise-not.wgsl.jinja"], "elementwise": ["bitwise-not.wgsl.jinja"] }
20
  }
21
  }