Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.BitwiseNot

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

sync 2e7068faf55e

Browse files
Files changed (6) hide show
  1. README.md +56 -0
  2. build/webgpu/bench.json +29 -0
  3. build/webgpu/bitwise-not.wgsl.jinja +48 -0
  4. build/webgpu/manifest.json +76 -0
  5. build/webgpu/metadata.json +18 -0
  6. build/webgpu/test.json +135 -0
README.md CHANGED
@@ -1,3 +1,59 @@
1
  ---
 
2
  license: apache-2.0
 
 
 
 
3
  ---
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
  ---
2
+ library_name: kernels
3
  license: apache-2.0
4
+ tags:
5
+ - kernel
6
+ - webgpu
7
+ - wgsl
8
  ---
9
+ # ai.onnx.BitwiseNot
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 18
12
+
13
+ ## Description
14
+
15
+ Applies a bitwise NOT to each element of the input tensor, flipping every bit. The output has the same shape and integer dtype as the input.
16
+
17
+ See the [ONNX `BitwiseNot` spec](https://onnx.ai/onnx/operators/onnx__BitwiseNot.html) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `X` | `x` | `T` | — | — | Input integer tensor. | required |
24
+
25
+ ## Outputs
26
+
27
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
28
+ | --- | --- | --- | --- | --- | --- | --- |
29
+ | `Y` | `y` | `T` | same as `X` | same as `X` | Output tensor with each element bitwise-negated; same shape and dtype as `X`. | required |
30
+
31
+ ## Type constraints
32
+
33
+ | Variable | Allowed dtypes |
34
+ | --- | --- |
35
+ | `T` | `uint32`, `int32`, `int16`, `uint8`, `int8` |
36
+
37
+ ## Files
38
+
39
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, 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
+ The loader derives every required output's shape and logical dtype from the manifest contract and this call.
48
+ It then allocates the result tensors automatically.
49
+
50
+ The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
51
+
52
+ Replace each `*Data` placeholder with a typed array containing the corresponding input data.
53
+
54
+ ```js
55
+ import { getKernel } from "@huggingface/kernels";
56
+
57
+ const kernel = await getKernel("webgpu-kernels/ai.onnx.BitwiseNot", { version: 1 });
58
+ const { y } = await kernel({ x: { data: xData, shape: [3] } });
59
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,29 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.BitwiseNot",
3
+ "cases": [
4
+ {
5
+ "name": "u32_33m",
6
+ "preset": "stress",
7
+ "provenance": { "notes": "Stress-only capacity case: input plus output occupy exactly 256 MiB of GPU storage." },
8
+ "inputs": { "x": { "dtype": "uint32", "shape": [33554432], "data": { "kind": "constant", "value": 305419896 } } },
9
+ "outputs": { "y": { "dtype": "uint32", "shape": [33554432] } },
10
+ "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "33554432 * 4 * 2" }] }
11
+ },
12
+ {
13
+ "name": "u32_1m",
14
+ "preset": "smoke",
15
+ "inputs": { "x": { "dtype": "uint32", "shape": [1048576], "data": { "kind": "constant", "value": 305419896 } } },
16
+ "outputs": { "y": { "dtype": "uint32", "shape": [1048576] } },
17
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }] }
18
+ },
19
+ {
20
+ "name": "u32_33m_scalar_fallback",
21
+ "preset": "stress",
22
+ "inputs": {
23
+ "x": { "dtype": "uint32", "shape": [33554433], "dist": "uniform", "seed": 701, "min": 0, "max": 4294967295 }
24
+ },
25
+ "outputs": { "y": { "dtype": "uint32", "shape": [33554433] } },
26
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "33554433 * 4 * 2", "name": "scalar-fallback BW" }] }
27
+ }
28
+ ]
29
+ }
build/webgpu/bitwise-not.wgsl.jinja ADDED
@@ -0,0 +1,48 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
4
+ // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
5
+ {% elif note == "limit" %}
6
+ // 2D-folded flat index: gid.y carries the high bits past the
7
+ // maxComputeWorkgroupsPerDimension limit.
8
+ {% elif note == "device-axis" %}
9
+ // The flat dispatch is folded across x/y at the device's per-axis workgroup
10
+ // limit; 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
13
+ // maxComputeWorkgroupsPerDimension limit (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
+ // maxComputeWorkgroupsPerDimension limit.
17
+ {% elif note == "dispatch" %}
18
+ // 2D-folded flat index: gid.y carries the high bits past the
19
+ // maxComputeWorkgroupsPerDimension dispatch limit.
20
+ {% endif %}
21
+ {% if bound == "" %}
22
+ let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
23
+ {%- elif guardInline %}
24
+ let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
25
+ if ({{ name }} >= {{ bound }}) { return; }
26
+ {%- else %}
27
+ let {{ name }} = gid.x + gid.y * nwg.x * {{ 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>, @builtin(num_workgroups) nwg: vec3<u32>) {
38
+ {{ flat_index_2d(note="element-limit") }}
39
+ var value = ~x[i];
40
+ {% if logicalDtype == "uint8" %}
41
+ {% if source.vectorized %}
42
+ value = value & vec4<u32>(0xffu);
43
+ {% else %}
44
+ value = value & 0xffu;
45
+ {% endif %}
46
+ {% endif %}
47
+ y[i] = value;
48
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,76 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "ai.onnx",
3
+ "name": "BitwiseNot",
4
+ "sinceVersion": 18,
5
+ "description": "Applies a bitwise NOT to each element of the input tensor, flipping every bit. The output has the same shape and integer dtype as the input.",
6
+ "inputs": [{ "role": "X", "dtype": "T", "description": "Input integer tensor." }],
7
+ "outputs": [
8
+ {
9
+ "role": "Y",
10
+ "dtype": "T",
11
+ "rank": "ranks.X",
12
+ "description": "Output tensor with each element bitwise-negated; same shape and dtype as `X`.",
13
+ "shape": "shapes.X"
14
+ }
15
+ ],
16
+ "typeConstraints": { "T": ["uint32", "int32", "int16", "uint8", "int8"] },
17
+ "args": {
18
+ "x": { "kind": "tensor", "semantic": "X", "role": "input" },
19
+ "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
20
+ },
21
+ "tunables": { "WORKGROUP_SIZE": 256 },
22
+ "derive": { "shapeContract": "numel(shapes.X) == numel(shapes.Y)" },
23
+ "constants": { "logicalDtype": "tensorDtypes.Y" },
24
+ "bindingSets": {
25
+ "vector": [
26
+ { "name": "x", "arg": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$vectorScalar" },
27
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
28
+ {
29
+ "name": "params",
30
+ "semantic": "kernel.params",
31
+ "buffer": { "type": "uniform" },
32
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
33
+ }
34
+ ],
35
+ "scalar": [
36
+ { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
37
+ { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
38
+ {
39
+ "name": "params",
40
+ "semantic": "kernel.params",
41
+ "buffer": { "type": "uniform" },
42
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" }] }
43
+ }
44
+ ]
45
+ },
46
+ "variants": [
47
+ {
48
+ "id": "elementwise_vec4",
49
+ "priority": 20,
50
+ "when": ["numel(shapes.X) % 4 == 0", "shapeContract"],
51
+ "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
52
+ "passes": [
53
+ {
54
+ "id": "main",
55
+ "name": "BitwiseNot.vec4",
56
+ "source": { "shader": "bitwise-not.wgsl.jinja", "inputs": { "vectorized": true } },
57
+ "bindings": "vector",
58
+ "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
59
+ }
60
+ ]
61
+ },
62
+ {
63
+ "id": "elementwise",
64
+ "when": "shapeContract",
65
+ "passes": [
66
+ {
67
+ "id": "main",
68
+ "name": "BitwiseNot",
69
+ "source": { "shader": "bitwise-not.wgsl.jinja", "inputs": { "vectorized": false } },
70
+ "bindings": "scalar",
71
+ "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
72
+ }
73
+ ]
74
+ }
75
+ ]
76
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,18 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "ai.onnx.BitwiseNot",
3
+ "id": "_ai_onnx_bitwisenot_webgpu_32a96bb",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "CczI58x8pA6aakhXdpmw14wRRq8CEArt333TMzvUE5I=",
11
+ "bitwise-not.wgsl.jinja": "f8h+mARfRrpuDY+hBRyHRhegkFnDEct4lQBZx36dcos=",
12
+ "manifest.json": "0dDdOkc503LzGY8f0BAs48R9B/tEFd/aiqfgfYE/5/Y=",
13
+ "test.json": "XY7peYVfnFEA8/3VjLC1K8VgcdVQMfTvctow0rwlmTM="
14
+ }
15
+ },
16
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
17
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.BitwiseNot" }
18
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,135 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.BitwiseNot",
3
+ "cases": [
4
+ {
5
+ "name": "int32_signed_ort",
6
+ "inputs": { "x": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-1, -2, 3] } } },
7
+ "outputs": { "y": { "dtype": "int32", "shape": [3] } },
8
+ "provenance": {
9
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
10
+ "test": "MathOpTest.BitwiseNot"
11
+ }
12
+ },
13
+ {
14
+ "name": "ort_int16_signed",
15
+ "provenance": {
16
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
17
+ "test": "MathOpTest.BitwiseNot_int8",
18
+ "notes": "Signed ORT bitwise pattern widened to ONNX-valid int16 with min/max sentinels."
19
+ },
20
+ "inputs": {
21
+ "x": { "dtype": "int16", "shape": [5], "data": { "kind": "values", "values": [-32768, -2, -1, 0, 32767] } }
22
+ },
23
+ "outputs": { "y": { "dtype": "int16", "shape": [5], "tolerance": 0 } }
24
+ },
25
+ {
26
+ "name": "int8_signed_ort",
27
+ "provenance": {
28
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
29
+ "test": "MathOpTest.BitwiseNot_int8",
30
+ "notes": "Extends ORT's signed int8 values with min/max sentinels."
31
+ },
32
+ "inputs": {
33
+ "x": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [-128, -2, -1, 0, 127] } }
34
+ },
35
+ "outputs": { "y": { "dtype": "int8", "shape": [5] } }
36
+ },
37
+ {
38
+ "name": "uint8_logical_width",
39
+ "inputs": { "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [1, 4, 5, 255] } } },
40
+ "outputs": { "y": { "dtype": "uint8", "shape": [4] } }
41
+ },
42
+ {
43
+ "name": "int32_high_bits_no_float_cast",
44
+ "inputs": {
45
+ "x": {
46
+ "dtype": "int32",
47
+ "shape": [6],
48
+ "data": {
49
+ "kind": "values",
50
+ "values": [2147483647, -2147483648, 16777217, -16777217, 1431655765, -1431655766]
51
+ }
52
+ }
53
+ },
54
+ "outputs": { "y": { "dtype": "int32", "shape": [6] } }
55
+ },
56
+ {
57
+ "name": "uint32_high_bits_no_float_cast",
58
+ "inputs": {
59
+ "x": {
60
+ "dtype": "uint32",
61
+ "shape": [5],
62
+ "data": { "kind": "values", "values": [4294967295, 4000000001, 2147483648, 16777217, 305419896] }
63
+ }
64
+ },
65
+ "outputs": { "y": { "dtype": "uint32", "shape": [5] } }
66
+ },
67
+ {
68
+ "name": "ort_int8_vector",
69
+ "provenance": {
70
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
71
+ "test": "MathOpTest.BitwiseNot_int8"
72
+ },
73
+ "inputs": { "x": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-1, -2, 3] } } },
74
+ "outputs": { "y": { "dtype": "int8", "shape": [3], "tolerance": 0 } }
75
+ },
76
+ {
77
+ "name": "ort_uint8_vector",
78
+ "provenance": {
79
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
80
+ "test": "MathOpTest.BitwiseNot_uint8"
81
+ },
82
+ "inputs": { "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [1, 4, 5, 3] } } },
83
+ "outputs": { "y": { "dtype": "uint8", "shape": [4], "tolerance": 0 } }
84
+ },
85
+ {
86
+ "name": "onnx_backend_bitwise_not_2d",
87
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitwise_not_2d" },
88
+ "inputs": {
89
+ "x": {
90
+ "dtype": "int32",
91
+ "shape": [3, 4],
92
+ "data": {
93
+ "kind": "values",
94
+ "values": [-356387803, 2135392491, 946286476, 1857819720, -2146992385, -1597193335, -848975157, 2143362693, -1517171889, -1133489216, -1750892400, -444182399]
95
+ }
96
+ }
97
+ },
98
+ "outputs": { "y": { "dtype": "int32", "shape": [3, 4], "tolerance": 0 } }
99
+ },
100
+ {
101
+ "name": "onnx_backend_bitwise_not_4d",
102
+ "provenance": {
103
+ "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitwise_not_4d",
104
+ "test": "test_bitwise_not_4d"
105
+ },
106
+ "inputs": {
107
+ "x": {
108
+ "dtype": "uint8",
109
+ "shape": [3, 4, 5, 6],
110
+ "data": {
111
+ "kind": "values",
112
+ "values": [37, 235, 140, 72, 137, 203, 133, 79, 192, 144, 129, 204, 71, 237, 252, 134, 25, 178, 20, 254, 101, 146, 212, 139, 252, 234, 156, 157, 142, 50, 68, 215, 215, 233, 241, 247, 222, 96, 86, 141, 233, 137, 7, 63, 61, 22, 57, 1, 128, 60, 209, 8, 216, 141, 115, 175, 234, 121, 200, 30, 71, 131, 198, 149, 49, 57, 3, 196, 24, 241, 43, 76, 26, 52, 80, 109, 115, 41, 210, 15, 64, 196, 25, 111, 226, 215, 135, 26, 153, 104, 22, 9, 195, 231, 126, 23, 125, 100, 155, 165, 57, 83, 166, 136, 32, 162, 10, 23, 143, 239, 87, 25, 199, 243, 92, 74, 190, 46, 160, 216, 151, 183, 65, 113, 77, 3, 128, 248, 253, 205, 6, 52, 213, 198, 2, 76, 219, 149, 203, 7, 77, 200, 75, 76, 43, 20, 30, 36, 103, 7, 45, 196, 57, 240, 252, 210, 96, 13, 10, 23, 124, 81, 135, 121, 152, 202, 220, 148, 160, 140, 193, 94, 60, 233, 152, 82, 115, 97, 130, 236, 220, 103, 98, 10, 182, 96, 233, 82, 214, 198, 194, 71, 231, 176, 54, 15, 133, 145, 170, 20, 118, 176, 22, 229, 141, 241, 114, 97, 181, 84, 10, 96, 183, 61, 56, 217, 149, 231, 96, 249, 211, 25, 241, 142, 141, 212, 116, 43, 134, 205, 184, 187, 143, 24, 251, 137, 194, 199, 53, 69, 164, 101, 248, 21, 40, 77, 219, 177, 113, 47, 205, 168, 206, 45, 87, 144, 156, 234, 45, 195, 116, 66, 206, 46, 128, 157, 63, 75, 35, 181, 221, 33, 130, 212, 83, 48, 54, 32, 253, 156, 55, 210, 31, 28, 222, 74, 136, 109, 99, 32, 8, 84, 205, 50, 79, 169, 64, 108, 211, 24, 113, 234, 20, 44, 15, 158, 219, 142, 243, 19, 251, 154, 235, 86, 135, 227, 53, 175, 110, 188, 243, 34, 100, 228, 32, 147, 195, 24, 211, 229, 94, 166, 175, 231, 133, 79, 63, 113, 87, 160, 42, 202, 194, 216, 226, 158, 145, 68, 64, 188, 244, 206, 145, 167, 163, 209, 156, 150, 38, 169]
113
+ }
114
+ }
115
+ },
116
+ "outputs": { "y": { "dtype": "uint8", "shape": [3, 4, 5, 6], "tolerance": 0 } }
117
+ },
118
+ {
119
+ "name": "empty_input_zero_dim",
120
+ "inputs": { "x": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } } },
121
+ "outputs": { "y": { "dtype": "int32", "shape": [0], "tolerance": 0 } }
122
+ },
123
+ {
124
+ "name": "int8_vec4_boundary_values",
125
+ "inputs": {
126
+ "x": {
127
+ "dtype": "int8",
128
+ "shape": [8],
129
+ "data": { "kind": "values", "values": [-128, -127, -2, -1, 0, 1, 126, 127] }
130
+ }
131
+ },
132
+ "outputs": { "y": { "dtype": "int8", "shape": [8], "tolerance": 0 } }
133
+ }
134
+ ]
135
+ }