Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.BitwiseOr

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

sync 91d990483a17

Browse files
Files changed (7) hide show
  1. README.md +12 -8
  2. build/webgpu/bench.json +0 -1
  3. build/webgpu/bitwise-binary-broadcast.wgsl.jinja +15 -14
  4. build/webgpu/bitwise-binary-vec4.wgsl.jinja +28 -12
  5. build/webgpu/manifest.json +68 -111
  6. build/webgpu/metadata.json +15 -8
  7. build/webgpu/test.json +0 -1
README.md CHANGED
@@ -18,16 +18,16 @@ See the [ONNX `BitwiseOr` spec](https://onnx.ai/onnx/operators/onnx__BitwiseOr.h
18
 
19
  ## Inputs
20
 
21
- | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
  | --- | --- | --- | --- | --- | --- | --- |
23
- | `A` | `a` | `T` | — | — | First input operand for the bitwise OR. | required |
24
- | `B` | `b` | `T` | — | — | Second input operand for the bitwise OR. | required |
25
 
26
  ## Outputs
27
 
28
- | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
29
  | --- | --- | --- | --- | --- | --- | --- |
30
- | `C` | `c` | `T` | derived | broadcast result of `A` and `B` | Result tensor containing the elementwise bitwise OR of A and B. | required |
31
 
32
  ## Type constraints
33
 
@@ -37,7 +37,7 @@ See the [ONNX `BitwiseOr` spec](https://onnx.ai/onnx/operators/onnx__BitwiseOr.h
37
 
38
  ## Files
39
 
40
- - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
41
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
42
  - [`test.json`](build/webgpu/test.json) — correctness cases
43
  - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
@@ -46,10 +46,14 @@ See the [ONNX `BitwiseOr` spec](https://onnx.ai/onnx/operators/onnx__BitwiseOr.h
46
 
47
  ## Use with `@huggingface/kernels`
48
 
49
- The loader derives every required output's shape and logical dtype from the manifest contract and this call.
50
- It then allocates the result tensors automatically.
 
 
 
51
 
52
  The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 
53
 
54
  Replace each `*Data` placeholder with a typed array containing the corresponding input data.
55
 
 
18
 
19
  ## Inputs
20
 
21
+ | Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
22
  | --- | --- | --- | --- | --- | --- | --- |
23
+ | `a` | `A` | `T` | — | — | First input operand for the bitwise OR. | required |
24
+ | `b` | `B` | `T` | — | — | Second input operand for the bitwise OR. | required |
25
 
26
  ## Outputs
27
 
28
+ | Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
29
  | --- | --- | --- | --- | --- | --- | --- |
30
+ | `c` | `C` | `T` | derived | broadcast result of `a` and `b` | Result tensor containing the elementwise bitwise OR of A and B. | required |
31
 
32
  ## Type constraints
33
 
 
37
 
38
  ## Files
39
 
40
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
41
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
42
  - [`test.json`](build/webgpu/test.json) — correctness cases
43
  - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 
46
 
47
  ## Use with `@huggingface/kernels`
48
 
49
+ ```sh
50
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.2
51
+ ```
52
+
53
+ Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
54
 
55
  The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
56
+ It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
57
 
58
  Replace each `*Data` placeholder with a typed array containing the corresponding input data.
59
 
build/webgpu/bench.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "ai.onnx.BitwiseOr",
3
  "cases": [
4
  {
5
  "name": "u32_1m",
 
1
  {
 
2
  "cases": [
3
  {
4
  "name": "u32_1m",
build/webgpu/bitwise-binary-broadcast.wgsl.jinja CHANGED
@@ -1,13 +1,14 @@
1
  {% macro flat_tail_open() %}
2
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
3
- fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
4
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
5
- // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
6
- let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
7
  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
8
  // the logical tensor length (or its storage binding) to be vec4 aligned.
9
- let begin = invocation * {{ source.itemsPerInvocation }}u;
10
- let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
 
11
  for (var i = begin; i < end; i = i + 1u) {
12
  {%- endmacro %}
13
  {% macro flat_tail_close() %}
@@ -75,9 +76,15 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
75
  {% endif %}
76
  {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
77
  {%- endmacro %}{% macro binary_broadcast_offsets() %}
78
- {{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.cShape, source.cRank) }}
 
 
 
 
 
 
79
 
80
- {{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.cShape, source.cRank) }}
81
  {%- endmacro %}
82
 
83
  {{ env.wgsl.resourceDeclarations }}
@@ -86,13 +93,7 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
86
  {{ binary_broadcast_offsets() }}
87
 
88
  {{ flat_tail_open() }}
89
- {% if bitwiseOp == "and" %}
90
- var value = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] & b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
91
- {% elif bitwiseOp == "or" %}
92
- var value = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] | b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
93
- {% else %}
94
- var value = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] ^ b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
95
- {% endif %}
96
  {% if logicalDtype == "uint8" %}
97
  value = value & 0xffu;
98
  {% endif %}
 
1
  {% macro flat_tail_open() %}
2
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
3
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
4
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
5
+ // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
6
+ let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
7
  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
8
  // the logical tensor length (or its storage binding) to be vec4 aligned.
9
+ {% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
10
+ let begin = invocation * {{ itemsPerInvocation }}u;
11
+ let end = min(begin + {{ itemsPerInvocation }}u, params.count);
12
  for (var i = begin; i < end; i = i + 1u) {
13
  {%- endmacro %}
14
  {% macro flat_tail_close() %}
 
76
  {% endif %}
77
  {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
78
  {%- endmacro %}{% macro binary_broadcast_offsets() %}
79
+ {% set aShape = aShape | default([]) %}
80
+ {% set aRank = aRank | default(0) %}
81
+ {% set bShape = bShape | default([]) %}
82
+ {% set bRank = bRank | default(0) %}
83
+ {% set cShape = cShape | default([]) %}
84
+ {% set cRank = cRank | default(0) %}
85
+ {{ broadcast_offset_fn("a_offset", aShape, aRank, cShape, cRank) }}
86
 
87
+ {{ broadcast_offset_fn("b_offset", bShape, bRank, cShape, cRank) }}
88
  {%- endmacro %}
89
 
90
  {{ env.wgsl.resourceDeclarations }}
 
93
  {{ binary_broadcast_offsets() }}
94
 
95
  {{ flat_tail_open() }}
96
+ var value = a[{{ broadcast_offset_call("a_offset", aShape, cShape, "i") }}] | b[{{ broadcast_offset_call("b_offset", bShape, cShape, "i") }}];
 
 
 
 
 
 
97
  {% if logicalDtype == "uint8" %}
98
  value = value & 0xffu;
99
  {% endif %}
build/webgpu/bitwise-binary-vec4.wgsl.jinja CHANGED
@@ -1,26 +1,42 @@
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
- // Same-shape vec4 bitwise binary (and/or/xor): 4 contiguous elements per thread
4
  // (128-bit loads/stores). uint8 storage uses one u32 slot per element, so
5
  // the result is masked to the low byte per lane. Same semantics as the scalar
6
  // broadcast kernel when A, B, C share a shape.
 
 
 
 
 
7
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
8
- fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
9
- // 2D-folded flat vec4 index: gid.y carries the high bits past the
10
- // maxComputeWorkgroupsPerDimension limit.
11
- let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
 
 
 
 
 
 
 
 
 
 
 
 
12
  if (i >= params.count) {
13
  return;
14
  }
15
- {% if bitwiseOp == "and" %}
16
- var value = a[i] & b[i];
17
- {% elif bitwiseOp == "or" %}
18
- var value = a[i] | b[i];
19
- {% else %}
20
- var value = a[i] ^ b[i];
21
  {% endif %}
22
- {% if logicalDtype == "uint8" %}
 
 
23
  value = value & vec4<u32>(0xffu);
24
  {% endif %}
25
  c[i] = value;
 
 
 
26
  }
 
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
+ // Same-shape vec4 bitwise binary (and/or/xor): 4 contiguous elements per lane
4
  // (128-bit loads/stores). uint8 storage uses one u32 slot per element, so
5
  // the result is masked to the low byte per lane. Same semantics as the scalar
6
  // broadcast kernel when A, B, C share a shape.
7
+ {% set vec4PerThread = vec4PerThread %}
8
+ {% if vec4PerThread > 1 %}
9
+ const ITEMS: u32 = {{ vec4PerThread }}u;
10
+ {% endif %}
11
+
12
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
13
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
14
+ // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
15
+ // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
16
+ {% if vec4PerThread > 1 %}
17
+ // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
18
+ // access consecutive words on every step, while each lane can keep several
19
+ // independent loads in flight.
20
+ let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
21
+ let span = (params.count + ITEMS - 1u) / ITEMS;
22
+ for (var j = 0u; j < ITEMS; j = j + 1u) {
23
+ let i = tid + j * span;
24
+ if (i >= params.count) {
25
+ break;
26
+ }
27
+ {% else %}
28
+ let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
29
  if (i >= params.count) {
30
  return;
31
  }
 
 
 
 
 
 
32
  {% endif %}
33
+
34
+ var value = a[i] | b[i];
35
+ {% if cDtype == "uint8" %}
36
  value = value & vec4<u32>(0xffu);
37
  {% endif %}
38
  c[i] = value;
39
+ {% if vec4PerThread > 1 %}
40
+ }
41
+ {% endif %}
42
  }
build/webgpu/manifest.json CHANGED
@@ -2,148 +2,105 @@
2
  "domain": "ai.onnx",
3
  "name": "BitwiseOr",
4
  "sinceVersion": 18,
5
- "description": "Computes the elementwise bitwise `or` of two integer tensors `A` and `B`, with multidirectional (NumPy-style) broadcasting. The output `C` has the broadcasted shape and the same dtype as the inputs.",
6
- "inputs": [
7
- { "role": "A", "dtype": "T", "description": "First input operand for the bitwise OR." },
8
- { "role": "B", "dtype": "T", "description": "Second input operand for the bitwise OR." }
9
- ],
10
- "outputs": [
11
- {
12
- "role": "C",
13
- "dtype": "T",
14
- "rank": "max(ranks.A, ranks.B)",
15
- "description": "Result tensor containing the elementwise bitwise OR of A and B.",
16
- "shape": "broadcastShape(shapes.A, shapes.B)"
17
- }
18
- ],
19
- "typeConstraints": { "T": ["uint32", "int32", "int16", "uint8", "int8"] },
20
- "args": {
21
- "a": { "kind": "tensor", "semantic": "A", "role": "input" },
22
- "b": { "kind": "tensor", "semantic": "B", "role": "input" },
23
- "c": { "kind": "tensor", "semantic": "C", "role": "output" }
24
  },
25
- "tunables": { "WORKGROUP_SIZE": 256 },
 
26
  "variants": [
27
  {
28
  "id": "same_shape_vec4",
29
  "priority": 20,
30
- "when": ["tensorDtypes.A == tensorDtypes.B", "tensorDtypes.A == tensorDtypes.C", "sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 == 0"],
31
- "constants": {
32
- "bitwiseOp": "\"or\"",
33
- "logicalDtype": "tensorDtypes.C",
34
- "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
35
- },
36
  "passes": [
37
  {
38
  "id": "main",
39
  "name": "BitwiseOr.vec4",
40
- "source": { "shader": "bitwise-binary-vec4.wgsl.jinja" },
41
- "bindings": [
42
- {
43
- "name": "a",
44
- "arg": "a",
45
- "semantic": "A",
46
- "buffer": { "type": "read-only-storage" },
47
- "elementType": "$vectorScalar"
48
- },
49
- {
50
- "name": "b",
51
- "arg": "b",
52
- "semantic": "B",
53
- "buffer": { "type": "read-only-storage" },
54
- "elementType": "$vectorScalar"
55
- },
56
- {
57
- "name": "c",
58
- "arg": "c",
59
- "semantic": "C",
60
- "buffer": { "type": "storage" },
61
- "elementType": "$vectorScalar"
62
- },
63
- {
64
- "name": "params",
65
- "semantic": "kernel.params",
66
- "buffer": { "type": "uniform" },
67
- "struct": {
68
- "name": "Params",
69
- "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
70
- }
71
- }
72
- ],
73
- "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
74
  }
75
  ]
76
  },
77
  {
78
  "id": "broadcast",
79
- "when": ["tensorDtypes.A == tensorDtypes.B", "tensorDtypes.A == tensorDtypes.C", "ranks.A <= ranks.C", "ranks.B <= ranks.C"],
80
- "constants": { "bitwiseOp": "\"or\"", "logicalDtype": "tensorDtypes.C" },
81
  "passes": [
82
  {
83
  "id": "main",
84
  "name": "BitwiseOr",
85
- "source": {
86
- "shader": "bitwise-binary-broadcast.wgsl.jinja",
87
- "inputs": {
88
- "aShape": "shapes.A",
89
- "bShape": "shapes.B",
90
- "cShape": "shapes.C",
91
- "aRank": "ranks.A",
92
- "bRank": "ranks.B",
93
- "cRank": "ranks.C",
94
- "itemsPerInvocation": 4
95
- }
96
  },
97
- "bindings": [
98
- { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
99
- { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
100
- { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$T" },
101
- {
102
- "name": "params",
103
- "semantic": "kernel.params",
104
- "buffer": { "type": "uniform" },
105
- "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
106
- }
107
- ],
108
- "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
109
  }
110
  ]
111
  },
112
  {
113
  "id": "same_shape_scalar_x4",
114
  "priority": 15,
115
- "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 != 0", "tensorDtypes.A == tensorDtypes.B", "tensorDtypes.A == tensorDtypes.C", "ranks.A <= ranks.C", "ranks.B <= ranks.C"],
116
- "constants": { "bitwiseOp": "\"or\"", "logicalDtype": "tensorDtypes.C" },
117
  "passes": [
118
  {
119
  "id": "main",
120
  "name": "BitwiseOr",
121
- "source": {
122
- "shader": "bitwise-binary-broadcast.wgsl.jinja",
123
- "inputs": {
124
- "aShape": "shapes.A",
125
- "bShape": "shapes.B",
126
- "cShape": "shapes.C",
127
- "aRank": "ranks.A",
128
- "bRank": "ranks.B",
129
- "cRank": "ranks.C",
130
- "itemsPerInvocation": 4
131
- }
132
  },
133
- "bindings": [
134
- { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
135
- { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
136
- { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$T" },
137
- {
138
- "name": "params",
139
- "semantic": "kernel.params",
140
- "buffer": { "type": "uniform" },
141
- "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
142
- }
143
- ],
144
- "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
145
  }
146
  ]
147
  }
148
- ]
 
 
 
 
 
 
 
 
 
 
 
 
 
 
149
  }
 
2
  "domain": "ai.onnx",
3
  "name": "BitwiseOr",
4
  "sinceVersion": 18,
5
+ "inputs": { "a": { "onnx": "A", "dtype": "T" }, "b": { "onnx": "B", "dtype": "T" } },
6
+ "outputs": {
7
+ "c": { "onnx": "C", "dtype": "T", "rank": "max(ranks.a, ranks.b)", "shape": "broadcastShape(shapes.a, shapes.b)" }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
8
  },
9
+ "typeConstraints": { "T": ["uint32", "int32", "int16", "uint8", "int8"] },
10
+ "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
11
  "variants": [
12
  {
13
  "id": "same_shape_vec4",
14
  "priority": 20,
15
+ "when": ["tensorDtypes.a == tensorDtypes.b", "tensorDtypes.a == tensorDtypes.c", "sameShape(shapes.a, shapes.c)", "sameShape(shapes.b, shapes.c)", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0"],
16
+ "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
 
 
 
 
17
  "passes": [
18
  {
19
  "id": "main",
20
  "name": "BitwiseOr.vec4",
21
+ "shader": "bitwise-binary-vec4.wgsl.jinja",
22
+ "derive": {
23
+ "op": "\"or\"",
24
+ "cDtype": "tensorDtypes.c",
25
+ "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
26
+ },
27
+ "bindings": ["a", "b", "c_bitwise", "params"],
28
+ "dispatch": {
29
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
30
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
31
+ "z": 1
32
+ }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
33
  }
34
  ]
35
  },
36
  {
37
  "id": "broadcast",
38
+ "when": ["tensorDtypes.a == tensorDtypes.b", "tensorDtypes.a == tensorDtypes.c", "ranks.a <= ranks.c", "ranks.b <= ranks.c"],
39
+ "derive": { "bitwiseOp": "\"or\"", "logicalDtype": "tensorDtypes.c" },
40
  "passes": [
41
  {
42
  "id": "main",
43
  "name": "BitwiseOr",
44
+ "shader": "bitwise-binary-broadcast.wgsl.jinja",
45
+ "derive": {
46
+ "aShape": "shapes.a",
47
+ "bShape": "shapes.b",
48
+ "cShape": "shapes.c",
49
+ "aRank": "ranks.a",
50
+ "bRank": "ranks.b",
51
+ "cRank": "ranks.c",
52
+ "itemsPerInvocation": 4
 
 
53
  },
54
+ "bindings": ["a_2_bitwise", "b_2_bitwise", "c_2_bitwise", "params_2"],
55
+ "dispatch": {
56
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
57
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
58
+ "z": 1
59
+ }
 
 
 
 
 
 
60
  }
61
  ]
62
  },
63
  {
64
  "id": "same_shape_scalar_x4",
65
  "priority": 15,
66
+ "when": ["sameShape(shapes.a, shapes.c)", "sameShape(shapes.b, shapes.c)", "numel(shapes.c) > 0", "numel(shapes.c) % 4 != 0", "tensorDtypes.a == tensorDtypes.b", "tensorDtypes.a == tensorDtypes.c", "ranks.a <= ranks.c", "ranks.b <= ranks.c"],
67
+ "derive": { "bitwiseOp": "\"or\"", "logicalDtype": "tensorDtypes.c" },
68
  "passes": [
69
  {
70
  "id": "main",
71
  "name": "BitwiseOr",
72
+ "shader": "bitwise-binary-broadcast.wgsl.jinja",
73
+ "derive": {
74
+ "aShape": "shapes.a",
75
+ "bShape": "shapes.b",
76
+ "cShape": "shapes.c",
77
+ "aRank": "ranks.a",
78
+ "bRank": "ranks.b",
79
+ "cRank": "ranks.c",
80
+ "itemsPerInvocation": 4
 
 
81
  },
82
+ "bindings": ["a_2_bitwise", "b_2_bitwise", "c_2_bitwise", "params_2"],
83
+ "dispatch": {
84
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
85
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
86
+ "z": 1
87
+ }
 
 
 
 
 
 
88
  }
89
  ]
90
  }
91
+ ],
92
+ "bindings": {
93
+ "a": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
94
+ "b": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
95
+ "c_bitwise": { "buffer": "storage", "elementType": "$vectorScalar", "name": "c" },
96
+ "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c) / 4" }] },
97
+ "a_2_bitwise": { "buffer": "read-only-storage", "name": "a", "elementType": "$T" },
98
+ "b_2_bitwise": { "buffer": "read-only-storage", "name": "b", "elementType": "$T" },
99
+ "c_2_bitwise": { "buffer": "storage", "name": "c", "elementType": "$T" },
100
+ "params_2": {
101
+ "buffer": "uniform",
102
+ "name": "params",
103
+ "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c)" }]
104
+ }
105
+ }
106
  }
build/webgpu/metadata.json CHANGED
@@ -1,19 +1,26 @@
1
  {
2
  "name": "ai.onnx.BitwiseOr",
3
- "id": "_ai_onnx_bitwiseor_webgpu_6c70151",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "nrZEiMJy807p7i9B7Q+nRd39S2JEyYXUdRkD8Gy/MFI=",
11
- "bitwise-binary-broadcast.wgsl.jinja": "RCjn7BVX/bwpqqBWr5VMlxLy7mqeOJptl9+uJazGZ/w=",
12
- "bitwise-binary-vec4.wgsl.jinja": "IO1Z3esuK46FeyJHI7IeZb9mo2iV7Ugp4we/V1WbKO4=",
13
- "manifest.json": "0niDDQt0f7lDPm4hR444mwB4BF5biuxdfO0qJaccXio=",
14
- "test.json": "2X3NVUGCqJ/yQZVEB//MJEbf2eMEuOu+x7FQZ+stgiw="
15
  }
16
  },
17
- "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
18
- "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.BitwiseOr" }
 
 
 
 
 
 
 
19
  }
 
1
  {
2
  "name": "ai.onnx.BitwiseOr",
3
+ "id": "_ai_onnx_bitwiseor_webgpu_8e87398",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "mERKODQEQmulEut7roCKqmrmyS+N1AQRVcIWVUTpnwc=",
11
+ "bitwise-binary-broadcast.wgsl.jinja": "92rC9SLChfbxowLyL3FBnv8aF7q3Qz1MzTRuMJA0M9g=",
12
+ "bitwise-binary-vec4.wgsl.jinja": "Ya480vMzE1FogZwUV0yNxmmoN8txJGwMw0hYks0P8V8=",
13
+ "manifest.json": "g0N7NLtcwWg/wq+LSVjhGiCwg8rQcXh+St10paafCHI=",
14
+ "test.json": "kEOiVM09gGxUF6F5gGIdczfCsGnnzBc0ZlPlcGFceHM="
15
  }
16
  },
17
+ "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
18
+ "webgpu": {
19
+ "manifestSpec": "2.0",
20
+ "variants": {
21
+ "same_shape_vec4": ["bitwise-binary-vec4.wgsl.jinja"],
22
+ "broadcast": ["bitwise-binary-broadcast.wgsl.jinja"],
23
+ "same_shape_scalar_x4": ["bitwise-binary-broadcast.wgsl.jinja"]
24
+ }
25
+ }
26
  }
build/webgpu/test.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "ai.onnx.BitwiseOr",
3
  "fixtureArrays": {
4
  "onnx_backend_bitwise_or_i16_4d_values": [-91, 107, 12, -56, 9, 75, 5, -49, 64, 16, 1, 76, -57, 109, 124, 6, -103, 50, -108, 126, -27, 18, 84, 11, 124, 106, 28, 29, 14, -78, -60, 87, 87, 105, 113, 119, 94, -32, -42, 13, 105, 9, -121, -65, -67, -106, -71, -127, 0, -68, 81, -120, 88, 13, -13, 47, 106, -7, 72, -98, -57, 3, 70, 21, -79, -71, -125, 68, -104, 113, -85, -52, -102, -76, -48, -19, -13, -87, 82, -113, -64, 68, -103, -17, 98, 87, 7, -102, 25, -24, -106, -119, 67, 103, -2, -105, -3, -28, 27, 37, -71, -45, 38, 8, -96, 34, -118, -105, 15, 111, -41, -103, 71, 115, -36, -54, 62, -82, 32, 88, 23, 55, -63, -15, -51, -125, 0, 120, 125, 77, -122, -76, 85, 70, -126, -52, 91, 21, 75, -121, -51, 72, -53, -52, -85, -108, -98, -92, -25, -121, -83, 68, -71, 112, 124, 82, -32, -115, -118, -105, -4, -47, 7, -7, 24, 74, 92, 20, 32, 12, 65, -34, -68, 105, 24, -46, -13, -31, 2, 108, 92, -25, -30, -118, 54, -32, 105, -46, 86, 70, 66, -57, 103, 48, -74, -113, 5, 17, 42, -108, -10, 48, -106, 101, 13, 113, -14, -31, 53, -44, -118, -32, 55, -67, -72, 89, 21, 103, -32, 121, 83, -103, 113, 14, 13, 84, -12, -85, 6, 77, 56, 59, 15, -104, 123, 9, 66, 71, -75, -59, 36, -27, 120, -107, -88, -51, 91, 49, -15, -81, 77, 40, 78, -83, -41, 16, 28, 106, -83, 67, -12, -62, 78, -82, 0, 29, -65, -53, -93, 53, 93, -95, 2, 84, -45, -80, -74, -96, 125, 28, -73, 82, -97, -100, 94, -54, 8, -19, -29, -96, -120, -44, 77, -78, -49, 41, -64, -20, 83, -104, -15, 106, -108, -84, -113, 30, 91, 14, 115, -109, 123, 26, 107, -42, 7, 99, -75, 47, -18, 60, 115, -94, -28, 100, -96, 19, 67, -104, 83, 101, -34, 38, 47, 103, 5, -49, -65, -15, -41, 32, -86, 74, 66, 88, 98, 30, 17, -60, -64, 60, 116, 78, 17, 39, 35, 81, 28, 22, -90, 41]
5
  },
 
1
  {
 
2
  "fixtureArrays": {
3
  "onnx_backend_bitwise_or_i16_4d_values": [-91, 107, 12, -56, 9, 75, 5, -49, 64, 16, 1, 76, -57, 109, 124, 6, -103, 50, -108, 126, -27, 18, 84, 11, 124, 106, 28, 29, 14, -78, -60, 87, 87, 105, 113, 119, 94, -32, -42, 13, 105, 9, -121, -65, -67, -106, -71, -127, 0, -68, 81, -120, 88, 13, -13, 47, 106, -7, 72, -98, -57, 3, 70, 21, -79, -71, -125, 68, -104, 113, -85, -52, -102, -76, -48, -19, -13, -87, 82, -113, -64, 68, -103, -17, 98, 87, 7, -102, 25, -24, -106, -119, 67, 103, -2, -105, -3, -28, 27, 37, -71, -45, 38, 8, -96, 34, -118, -105, 15, 111, -41, -103, 71, 115, -36, -54, 62, -82, 32, 88, 23, 55, -63, -15, -51, -125, 0, 120, 125, 77, -122, -76, 85, 70, -126, -52, 91, 21, 75, -121, -51, 72, -53, -52, -85, -108, -98, -92, -25, -121, -83, 68, -71, 112, 124, 82, -32, -115, -118, -105, -4, -47, 7, -7, 24, 74, 92, 20, 32, 12, 65, -34, -68, 105, 24, -46, -13, -31, 2, 108, 92, -25, -30, -118, 54, -32, 105, -46, 86, 70, 66, -57, 103, 48, -74, -113, 5, 17, 42, -108, -10, 48, -106, 101, 13, 113, -14, -31, 53, -44, -118, -32, 55, -67, -72, 89, 21, 103, -32, 121, 83, -103, 113, 14, 13, 84, -12, -85, 6, 77, 56, 59, 15, -104, 123, 9, 66, 71, -75, -59, 36, -27, 120, -107, -88, -51, 91, 49, -15, -81, 77, 40, 78, -83, -41, 16, 28, 106, -83, 67, -12, -62, 78, -82, 0, 29, -65, -53, -93, 53, 93, -95, 2, 84, -45, -80, -74, -96, 125, 28, -73, 82, -97, -100, 94, -54, 8, -19, -29, -96, -120, -44, 77, -78, -49, 41, -64, -20, 83, -104, -15, 106, -108, -84, -113, 30, 91, 14, 115, -109, 123, 26, 107, -42, 7, 99, -75, 47, -18, 60, 115, -94, -28, 100, -96, 19, 67, -104, 83, 101, -34, 38, 47, 103, 5, -49, -65, -15, -41, 32, -86, 74, 66, 88, 98, 30, 17, -60, -64, 60, 116, 78, 17, 39, 35, 81, 28, 22, -90, 41]
4
  },