Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.BitwiseOr

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

sync 2e7068faf55e

Browse files
Files changed (7) hide show
  1. README.md +58 -0
  2. build/webgpu/bench.json +27 -0
  3. build/webgpu/bitwise-binary-broadcast.wgsl.jinja +110 -0
  4. build/webgpu/bitwise-binary-vec4.wgsl.jinja +26 -0
  5. build/webgpu/manifest.json +148 -0
  6. build/webgpu/metadata.json +19 -0
  7. build/webgpu/test.json +315 -0
README.md CHANGED
@@ -1,3 +1,61 @@
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.BitwiseOr
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 18
12
+
13
+ ## Description
14
+
15
+ 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.
16
+
17
+ See the [ONNX `BitwiseOr` spec](https://onnx.ai/onnx/operators/onnx__BitwiseOr.html) for the reference semantics.
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
+
34
+ | Variable | Allowed dtypes |
35
+ | --- | --- |
36
+ | `T` | `uint32`, `int32`, `int16`, `uint8`, `int8` |
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
44
+ - [`bitwise-binary-broadcast.wgsl.jinja`](build/webgpu/bitwise-binary-broadcast.wgsl.jinja)
45
+ - [`bitwise-binary-vec4.wgsl.jinja`](build/webgpu/bitwise-binary-vec4.wgsl.jinja)
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
+
56
+ ```js
57
+ import { getKernel } from "@huggingface/kernels";
58
+
59
+ const kernel = await getKernel("webgpu-kernels/ai.onnx.BitwiseOr", { version: 1 });
60
+ const { c } = await kernel({ a: { data: aData, shape: [3] }, b: { data: bData, shape: [3] } });
61
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,27 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.BitwiseOr",
3
+ "cases": [
4
+ {
5
+ "name": "u32_1m",
6
+ "preset": "smoke",
7
+ "inputs": {
8
+ "a": { "dtype": "uint32", "shape": [1048576], "data": { "kind": "constant", "value": 305419896 } },
9
+ "b": { "dtype": "uint32", "shape": [1048576], "data": { "kind": "constant", "value": 252645135 } }
10
+ },
11
+ "outputs": { "c": { "dtype": "uint32", "shape": [1048576] } },
12
+ "bench": {
13
+ "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.a) + numel(shapes.b) + numel(shapes.c)) * 4" }]
14
+ }
15
+ },
16
+ {
17
+ "name": "u32_same_shape_odd_numel",
18
+ "preset": "stress",
19
+ "inputs": {
20
+ "a": { "dtype": "uint32", "shape": [16777217], "dist": "uniformBits", "seed": 3 },
21
+ "b": { "dtype": "uint32", "shape": [16777217], "dist": "uniformBits", "seed": 4 }
22
+ },
23
+ "outputs": { "c": { "dtype": "uint32", "shape": [16777217], "dist": "empty" } },
24
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "16777217 * 4 * 3" }] }
25
+ }
26
+ ]
27
+ }
build/webgpu/bitwise-binary-broadcast.wgsl.jinja ADDED
@@ -0,0 +1,110 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
+ {% if source.itemsPerInvocation is defined %}
8
+ // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
9
+ // the logical tensor length (or its storage binding) to be vec4 aligned.
10
+ let begin = invocation * {{ source.itemsPerInvocation }}u;
11
+ let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
12
+ for (var i = begin; i < end; i = i + 1u) {
13
+ {%- else %}
14
+ let i = invocation;
15
+ if (i >= params.count) {
16
+ return;
17
+ }
18
+ {%- endif %}
19
+ {% endmacro %}
20
+ {% macro flat_tail_close() %}
21
+ {% if source.itemsPerInvocation is defined %}
22
+ }
23
+ {% endif %}
24
+ {% endmacro %}
25
+
26
+ {% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
27
+ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif %}) -> u32 {
28
+ {% if out_numel == 0 %}
29
+ return 0u;
30
+ {% elif op_numel == 1 %}
31
+ return 0u;
32
+ {% elif op_same %}
33
+ return out_index;
34
+ {% else %}
35
+ var offset = 0u;
36
+ {% for axis in range(outRank) %}
37
+ {% set op_axis = axis - (outRank - opRank) %}
38
+ {% if op_axis >= 0 and opShape[op_axis] != 1 %}
39
+ {% set c_stride = namespace(value=1) %}
40
+ {% for j in range(axis + 1, outRank) %}
41
+ {% set c_stride.value = c_stride.value * outShape[j] %}
42
+ {% endfor %}
43
+ {% set op_stride = namespace(value=1) %}
44
+ {% for j in range(op_axis + 1, opRank) %}
45
+ {% set op_stride.value = op_stride.value * opShape[j] %}
46
+ {% endfor %}
47
+ {% if c_stride.value == 1 %}
48
+ let coord{{ axis }} = out_index % {{ outShape[axis] }}u;
49
+ {% else %}
50
+ let coord{{ axis }} = (out_index / {{ c_stride.value }}u) % {{ outShape[axis] }}u;
51
+ {% endif %}
52
+ {% if op_stride.value == 1 %}
53
+ offset = offset + coord{{ axis }};
54
+ {% else %}
55
+ offset = offset + coord{{ axis }} * {{ op_stride.value }}u;
56
+ {% endif %}
57
+ {% endif %}
58
+ {% endfor %}
59
+ return offset;
60
+ {% endif %}
61
+ }
62
+ {%- endmacro %}{% macro broadcast_offset_call(fn_name, opShape, outShape, out_index) %}
63
+ {% set op_numel = namespace(value=1) %}
64
+ {% for d in opShape %}{% set op_numel.value = op_numel.value * d %}{% endfor %}
65
+ {% set out_numel = namespace(value=1) %}
66
+ {% for d in outShape %}{% set out_numel.value = out_numel.value * d %}{% endfor %}
67
+ {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
68
+ {%- endmacro %}{% macro broadcast_offset_fn(fn_name, opShape, opRank, outShape, outRank) %}
69
+ {% set op_numel = namespace(value=1) %}
70
+ {% for d in opShape %}
71
+ {% set op_numel.value = op_numel.value * d %}
72
+ {% endfor %}
73
+ {% set out_numel = namespace(value=1) %}
74
+ {% for d in outShape %}
75
+ {% set out_numel.value = out_numel.value * d %}
76
+ {% endfor %}
77
+ {% set op_same = namespace(value=(opRank == outRank)) %}
78
+ {% if op_same.value %}
79
+ {% for axis in range(outRank) %}
80
+ {% if opShape[axis] != outShape[axis] %}
81
+ {% set op_same.value = false %}
82
+ {% endif %}
83
+ {% endfor %}
84
+ {% endif %}
85
+ {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
86
+ {%- endmacro %}{% macro binary_broadcast_offsets() %}
87
+ {{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.cShape, source.cRank) }}
88
+
89
+ {{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.cShape, source.cRank) }}
90
+ {%- endmacro %}
91
+
92
+ {{ env.wgsl.resourceDeclarations }}
93
+
94
+
95
+ {{ binary_broadcast_offsets() }}
96
+
97
+ {{ flat_tail_open() }}
98
+ {% if bitwiseOp == "and" %}
99
+ var value = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] & b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
100
+ {% elif bitwiseOp == "or" %}
101
+ var value = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] | b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
102
+ {% else %}
103
+ var value = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] ^ b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
104
+ {% endif %}
105
+ {% if logicalDtype == "uint8" %}
106
+ value = value & 0xffu;
107
+ {% endif %}
108
+ c[i] = value;
109
+ {{ flat_tail_close() -}}
110
+ }
build/webgpu/bitwise-binary-vec4.wgsl.jinja ADDED
@@ -0,0 +1,26 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,148 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
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
+ }
95
+ },
96
+ "bindings": [
97
+ { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
98
+ { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
99
+ { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$T" },
100
+ {
101
+ "name": "params",
102
+ "semantic": "kernel.params",
103
+ "buffer": { "type": "uniform" },
104
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
105
+ }
106
+ ],
107
+ "dispatch": { "threads": "numel(shapes.C)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
108
+ }
109
+ ]
110
+ },
111
+ {
112
+ "id": "same_shape_scalar_x4",
113
+ "priority": 15,
114
+ "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"],
115
+ "constants": { "bitwiseOp": "\"or\"", "logicalDtype": "tensorDtypes.C" },
116
+ "passes": [
117
+ {
118
+ "id": "main",
119
+ "name": "BitwiseOr",
120
+ "source": {
121
+ "shader": "bitwise-binary-broadcast.wgsl.jinja",
122
+ "inputs": {
123
+ "aShape": "shapes.A",
124
+ "bShape": "shapes.B",
125
+ "cShape": "shapes.C",
126
+ "aRank": "ranks.A",
127
+ "bRank": "ranks.B",
128
+ "cRank": "ranks.C",
129
+ "itemsPerInvocation": 4
130
+ }
131
+ },
132
+ "bindings": [
133
+ { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
134
+ { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
135
+ { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$T" },
136
+ {
137
+ "name": "params",
138
+ "semantic": "kernel.params",
139
+ "buffer": { "type": "uniform" },
140
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
141
+ }
142
+ ],
143
+ "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
144
+ }
145
+ ]
146
+ }
147
+ ]
148
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,19 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "ai.onnx.BitwiseOr",
3
+ "id": "_ai_onnx_bitwiseor_webgpu_ee18180",
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": "N06Hs2oYVObcwu7yEYqmZQBNqGsjZGlVeWMhjxT23Hw=",
12
+ "bitwise-binary-vec4.wgsl.jinja": "IO1Z3esuK46FeyJHI7IeZb9mo2iV7Ugp4we/V1WbKO4=",
13
+ "manifest.json": "ljvAVQG0VC+AnTApqhyNIPLJchu9CbJ5P/M6gwZba3Q=",
14
+ "test.json": "2X3NVUGCqJ/yQZVEB//MJEbf2eMEuOu+x7FQZ+stgiw="
15
+ }
16
+ },
17
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
18
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.BitwiseOr" }
19
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,315 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
+ },
6
+ "cases": [
7
+ {
8
+ "name": "int32_signed_ort",
9
+ "inputs": {
10
+ "a": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-1, -2, 3] } },
11
+ "b": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 0, 4] } }
12
+ },
13
+ "outputs": { "c": { "dtype": "int32", "shape": [3] } },
14
+ "provenance": {
15
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
16
+ "test": "MathOpTest.BitwiseOr"
17
+ }
18
+ },
19
+ {
20
+ "name": "ort_int16_signed",
21
+ "provenance": {
22
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
23
+ "test": "MathOpTest.BitwiseOr_int8",
24
+ "notes": "Signed ORT bitwise pattern widened to ONNX-valid int16 with min/max sentinels."
25
+ },
26
+ "inputs": {
27
+ "a": { "dtype": "int16", "shape": [5], "data": { "kind": "values", "values": [-32768, -1, -2, 3, 32767] } },
28
+ "b": { "dtype": "int16", "shape": [5], "data": { "kind": "values", "values": [32767, 1, 0, 4, -32768] } }
29
+ },
30
+ "outputs": { "c": { "dtype": "int16", "shape": [5], "tolerance": 0 } }
31
+ },
32
+ {
33
+ "name": "uint8_logical_width",
34
+ "inputs": {
35
+ "a": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [1, 4, 5, 3] } },
36
+ "b": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [1, 2, 2, 1] } }
37
+ },
38
+ "outputs": { "c": { "dtype": "uint8", "shape": [4] } },
39
+ "provenance": {
40
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
41
+ "test": "MathOpTest.BitwiseOr_uint8"
42
+ }
43
+ },
44
+ {
45
+ "name": "broadcast_ab",
46
+ "inputs": {
47
+ "a": {
48
+ "dtype": "int32",
49
+ "shape": [4, 2],
50
+ "data": { "kind": "values", "values": [10, 11, 12, 13, 14, 15, 16, 17] }
51
+ },
52
+ "b": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [15, 7] } }
53
+ },
54
+ "outputs": { "c": { "dtype": "int32", "shape": [4, 2] } },
55
+ "provenance": {
56
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
57
+ "test": "MathOpTest.BitwiseOr_broadcastAB"
58
+ }
59
+ },
60
+ {
61
+ "name": "broadcast_ba",
62
+ "inputs": {
63
+ "a": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [15, 7] } },
64
+ "b": {
65
+ "dtype": "int32",
66
+ "shape": [4, 2],
67
+ "data": { "kind": "values", "values": [10, 11, 12, 13, 14, 15, 16, 17] }
68
+ }
69
+ },
70
+ "outputs": { "c": { "dtype": "int32", "shape": [4, 2] } },
71
+ "provenance": {
72
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
73
+ "test": "MathOpTest.BitwiseOr_broadcastBA"
74
+ }
75
+ },
76
+ {
77
+ "name": "multidirectional_broadcast",
78
+ "inputs": {
79
+ "a": { "dtype": "int32", "shape": [4, 1], "data": { "kind": "values", "values": [10, 11, 12, 13] } },
80
+ "b": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [15, 7] } }
81
+ },
82
+ "outputs": { "c": { "dtype": "int32", "shape": [4, 2] } },
83
+ "provenance": {
84
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
85
+ "test": "MathOpTest.BitwiseOr_multidirectional_broadcastAB"
86
+ }
87
+ },
88
+ {
89
+ "name": "multidirectional_broadcast_ba",
90
+ "inputs": {
91
+ "a": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [15, 7] } },
92
+ "b": { "dtype": "int32", "shape": [4, 1], "data": { "kind": "values", "values": [10, 11, 12, 13] } }
93
+ },
94
+ "outputs": { "c": { "dtype": "int32", "shape": [4, 2] } },
95
+ "provenance": {
96
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
97
+ "test": "MathOpTest.BitwiseOr_multidirectional_broadcastBA"
98
+ }
99
+ },
100
+ {
101
+ "name": "int32_high_bits_no_float_cast",
102
+ "inputs": {
103
+ "a": {
104
+ "dtype": "int32",
105
+ "shape": [6],
106
+ "data": {
107
+ "kind": "values",
108
+ "values": [2147483647, -2147483648, 16777217, -16777217, 1431655765, -1431655766]
109
+ }
110
+ },
111
+ "b": {
112
+ "dtype": "int32",
113
+ "shape": [6],
114
+ "data": { "kind": "values", "values": [252645135, 252645135, 16777216, 16777216, 858993459, 858993459] }
115
+ }
116
+ },
117
+ "outputs": { "c": { "dtype": "int32", "shape": [6] } }
118
+ },
119
+ {
120
+ "name": "uint32_high_bits_no_float_cast",
121
+ "inputs": {
122
+ "a": {
123
+ "dtype": "uint32",
124
+ "shape": [5],
125
+ "data": { "kind": "values", "values": [4294967295, 4000000001, 2147483648, 16777217, 305419896] }
126
+ },
127
+ "b": {
128
+ "dtype": "uint32",
129
+ "shape": [5],
130
+ "data": { "kind": "values", "values": [252645135, 4042322160, 2147483648, 16777216, 16711935] }
131
+ }
132
+ },
133
+ "outputs": { "c": { "dtype": "uint32", "shape": [5] } }
134
+ },
135
+ {
136
+ "name": "uint8_backend_bcast_4v3d",
137
+ "inputs": {
138
+ "a": {
139
+ "dtype": "uint8",
140
+ "shape": [2, 2, 2, 3],
141
+ "data": {
142
+ "kind": "values",
143
+ "values": [255, 170, 85, 15, 240, 51, 204, 129, 126, 1, 2, 4, 8, 16, 32, 64, 128, 127, 3, 5, 9, 17, 33, 65]
144
+ }
145
+ },
146
+ "b": {
147
+ "dtype": "uint8",
148
+ "shape": [2, 2, 3],
149
+ "data": { "kind": "values", "values": [15, 240, 170, 85, 51, 204, 3, 5, 9, 17, 33, 65] }
150
+ }
151
+ },
152
+ "outputs": { "c": { "dtype": "uint8", "shape": [2, 2, 2, 3] } }
153
+ },
154
+ {
155
+ "name": "ort_int8_vector",
156
+ "provenance": {
157
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
158
+ "test": "MathOpTest.BitwiseOr_int8"
159
+ },
160
+ "inputs": {
161
+ "a": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-1, -2, 3] } },
162
+ "b": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [1, 0, 4] } }
163
+ },
164
+ "outputs": { "c": { "dtype": "int8", "shape": [3], "tolerance": 0 } }
165
+ },
166
+ {
167
+ "name": "onnx_backend_bitwise_or_i32_2d",
168
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitwise_or_i32_2d" },
169
+ "inputs": {
170
+ "a": {
171
+ "dtype": "int32",
172
+ "shape": [3, 4],
173
+ "data": {
174
+ "kind": "values",
175
+ "values": [-356387803, 2135392491, 946286476, 1857819720, -2146992385, -1597193335, -848975157, 2143362693, -1517171889, -1133489216, -1750892400, -444182399]
176
+ }
177
+ },
178
+ "b": {
179
+ "dtype": "int32",
180
+ "shape": [3, 4],
181
+ "data": {
182
+ "kind": "values",
183
+ "values": [-356387803, 2135392491, 946286476, 1857819720, -2146992385, -1597193335, -848975157, 2143362693, -1517171889, -1133489216, -1750892400, -444182399]
184
+ }
185
+ }
186
+ },
187
+ "outputs": { "c": { "dtype": "int32", "shape": [3, 4], "tolerance": 0 } }
188
+ },
189
+ {
190
+ "name": "onnx_backend_bitwise_or_i16_4d",
191
+ "provenance": {
192
+ "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitwise_or_i16_4d",
193
+ "test": "test_bitwise_or_i16_4d"
194
+ },
195
+ "inputs": {
196
+ "a": {
197
+ "dtype": "int8",
198
+ "shape": [3, 4, 5, 6],
199
+ "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_bitwise_or_i16_4d_values" } }
200
+ },
201
+ "b": {
202
+ "dtype": "int8",
203
+ "shape": [3, 4, 5, 6],
204
+ "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_bitwise_or_i16_4d_values" } }
205
+ }
206
+ },
207
+ "outputs": { "c": { "dtype": "int8", "shape": [3, 4, 5, 6], "tolerance": 0 } }
208
+ },
209
+ {
210
+ "name": "onnx_backend_bitwise_or_ui8_bcast_4v3d",
211
+ "provenance": {
212
+ "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitwise_or_ui8_bcast_4v3d",
213
+ "test": "test_bitwise_or_ui8_bcast_4v3d"
214
+ },
215
+ "inputs": {
216
+ "a": {
217
+ "dtype": "uint8",
218
+ "shape": [3, 4, 5, 6],
219
+ "data": {
220
+ "kind": "values",
221
+ "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]
222
+ }
223
+ },
224
+ "b": {
225
+ "dtype": "uint8",
226
+ "shape": [4, 5, 6],
227
+ "data": {
228
+ "kind": "values",
229
+ "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]
230
+ }
231
+ }
232
+ },
233
+ "outputs": { "c": { "dtype": "uint8", "shape": [3, 4, 5, 6], "tolerance": 0 } }
234
+ },
235
+ {
236
+ "name": "empty_input_zero_dim",
237
+ "inputs": {
238
+ "a": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } },
239
+ "b": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } }
240
+ },
241
+ "outputs": { "c": { "dtype": "int32", "shape": [0], "tolerance": 0 } }
242
+ },
243
+ {
244
+ "name": "broadcast_scalar_a_uint32",
245
+ "inputs": {
246
+ "a": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [4042322160] } },
247
+ "b": {
248
+ "dtype": "uint32",
249
+ "shape": [2, 3],
250
+ "data": { "kind": "values", "values": [65535, 16711935, 4278255360, 4042322160, 252645135, 2863311530] }
251
+ }
252
+ },
253
+ "outputs": { "c": { "dtype": "uint32", "shape": [2, 3] } }
254
+ },
255
+ {
256
+ "name": "broadcast_rank6_uint32",
257
+ "inputs": {
258
+ "a": {
259
+ "dtype": "uint32",
260
+ "shape": [3, 1, 1, 1, 1, 1],
261
+ "data": { "kind": "values", "values": [252645135, 4042322160, 2863311530] }
262
+ },
263
+ "b": {
264
+ "dtype": "uint32",
265
+ "shape": [1, 1, 1, 1, 1, 3],
266
+ "data": { "kind": "values", "values": [16776960, 4278190335, 305419896] }
267
+ }
268
+ },
269
+ "outputs": {
270
+ "c": {
271
+ "dtype": "uint32",
272
+ "shape": [3, 1, 1, 1, 1, 3],
273
+ "data": {
274
+ "kind": "values",
275
+ "values": [268435215, 4279177215, 524246911, 4043309040, 4293980415, 4076140280, 2868903850, 4289374975, 3133079290]
276
+ },
277
+ "tolerance": 0
278
+ }
279
+ }
280
+ },
281
+ {
282
+ "name": "int8_sign_boundary_broadcast",
283
+ "inputs": {
284
+ "a": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [-128, -1, 0, 127, -85, 85] } },
285
+ "b": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [-1, -128, 127, 0, 85, -85] } }
286
+ },
287
+ "outputs": {
288
+ "c": {
289
+ "dtype": "int8",
290
+ "shape": [6],
291
+ "data": { "kind": "values", "values": [-1, -1, 127, 127, -1, -1] },
292
+ "tolerance": 0
293
+ }
294
+ }
295
+ },
296
+ {
297
+ "name": "empty_multidim_broadcast",
298
+ "inputs": {
299
+ "a": { "dtype": "int32", "shape": [2, 0, 4], "data": { "kind": "values", "values": [] } },
300
+ "b": { "dtype": "int32", "shape": [2, 0, 4], "data": { "kind": "values", "values": [] } }
301
+ },
302
+ "outputs": {
303
+ "c": { "dtype": "int32", "shape": [2, 0, 4], "data": { "kind": "values", "values": [] }, "tolerance": 0 }
304
+ }
305
+ },
306
+ {
307
+ "name": "rank7_broadcast_scalar_tail",
308
+ "inputs": {
309
+ "a": { "dtype": "uint32", "shape": [2, 1, 2, 1, 2, 1, 3], "data": { "kind": "cycle", "values": [1, 3, 7, 15] } },
310
+ "b": { "dtype": "uint32", "shape": [1, 2, 1, 2, 1, 2, 1], "data": { "kind": "cycle", "values": [2, 5, 10] } }
311
+ },
312
+ "outputs": { "c": { "dtype": "uint32", "shape": [2, 2, 2, 2, 2, 2, 3], "tolerance": 0 } }
313
+ }
314
+ ]
315
+ }