Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.BitShift

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

sync 2e7068faf55e

Browse files
Files changed (7) hide show
  1. README.md +68 -0
  2. build/webgpu/bench.json +30 -0
  3. build/webgpu/bitshift-vec4.wgsl.jinja +36 -0
  4. build/webgpu/bitshift.wgsl.jinja +93 -0
  5. build/webgpu/manifest.json +128 -0
  6. build/webgpu/metadata.json +19 -0
  7. build/webgpu/test.json +426 -0
README.md CHANGED
@@ -1,3 +1,71 @@
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.BitShift
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 11
12
+
13
+ ## Description
14
+
15
+ Performs an elementwise bitwise shift on unsigned integer tensors. `X` is shifted left or right by the amounts in `Y`, with the direction controlled by the `direction` attribute. Supports multidirectional (NumPy-style) broadcasting between `X` and `Y`.
16
+
17
+ See the [ONNX `BitShift` spec](https://onnx.ai/onnx/operators/onnx__BitShift.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 tensor to be shifted. | required |
24
+ | `Y` | `y` | `T` | — | — | Tensor specifying the number of bit positions to shift each element of X. | required |
25
+
26
+ ## Outputs
27
+
28
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
29
+ | --- | --- | --- | --- | --- | --- | --- |
30
+ | `Z` | `z` | `T` | derived | broadcast result of `X` and `Y` | Output tensor with the same shape as the broadcast result of X and Y. | required |
31
+
32
+ ## Attributes
33
+
34
+ Attributes and default values (overridable per request):
35
+
36
+ | Attribute | Default | Description |
37
+ | --- | --- | --- |
38
+ | `direction` | — | Direction of the bit shift: `LEFT` shifts bits toward higher significance (increasing value), while `RIGHT` shifts toward lower significance (decreasing value). |
39
+
40
+ ## Type constraints
41
+
42
+ | Variable | Allowed dtypes |
43
+ | --- | --- |
44
+ | `T` | `uint32`, `uint8` |
45
+
46
+ ## Files
47
+
48
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
49
+ - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
50
+ - [`test.json`](build/webgpu/test.json) — correctness cases
51
+ - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
52
+ - [`bitshift-vec4.wgsl.jinja`](build/webgpu/bitshift-vec4.wgsl.jinja)
53
+ - [`bitshift.wgsl.jinja`](build/webgpu/bitshift.wgsl.jinja)
54
+
55
+ ## Use with `@huggingface/kernels`
56
+
57
+ The loader derives every required output's shape and logical dtype from the manifest contract and this call.
58
+ It then allocates the result tensors automatically.
59
+
60
+ The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
61
+
62
+ Replace each `*Data` placeholder with a typed array containing the corresponding input data.
63
+
64
+ ```js
65
+ import { getKernel } from "@huggingface/kernels";
66
+
67
+ const kernel = await getKernel("webgpu-kernels/ai.onnx.BitShift", { version: 1 });
68
+ const { z } = await kernel({ x: { data: xData, shape: [2] }, y: { data: yData, shape: [2] } }, {
69
+ attrs: { direction: "RIGHT" },
70
+ });
71
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,30 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.BitShift",
3
+ "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
+ "cases": [
5
+ {
6
+ "name": "u32_left_1m",
7
+ "preset": "smoke",
8
+ "attrs": { "direction": "LEFT" },
9
+ "inputs": {
10
+ "x": { "dtype": "uint32", "shape": [1048576], "data": { "kind": "constant", "value": 1 } },
11
+ "y": { "dtype": "uint32", "shape": [1048576], "data": { "kind": "constant", "value": 3 } }
12
+ },
13
+ "outputs": { "z": { "dtype": "uint32", "shape": [1048576] } },
14
+ "bench": {
15
+ "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y) + numel(shapes.z)) * 4" }]
16
+ }
17
+ },
18
+ {
19
+ "name": "broadcast_scalar_x_4m_left_u32",
20
+ "preset": "edge",
21
+ "attrs": { "direction": "LEFT" },
22
+ "inputs": {
23
+ "x": { "dtype": "uint32", "shape": [], "dist": "uniform", "seed": 42, "min": 1, "max": 1000 },
24
+ "y": { "dtype": "uint32", "shape": [4194304], "dist": "uniform", "seed": 43, "min": 0, "max": 31 }
25
+ },
26
+ "outputs": { "z": { "dtype": "uint32", "shape": [4194304] } },
27
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.z) * 8" }] }
28
+ }
29
+ ]
30
+ }
build/webgpu/bitshift-vec4.wgsl.jinja ADDED
@@ -0,0 +1,36 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {{ env.wgsl.resourceDeclarations }}
2
+
3
+ const BIT_WIDTH: u32 = {{ bitWidth }}u;
4
+
5
+ fn shift_one(in_value: u32, amount: u32) -> u32 {
6
+ var value = 0u;
7
+ if (amount < BIT_WIDTH) {
8
+ {% if leftShift %}
9
+ value = in_value << amount;
10
+ {% else %}
11
+ value = in_value >> amount;
12
+ {% endif %}
13
+ {% if logicalDtype == "uint8" %}
14
+ value = value & 0xffu;
15
+ {% endif %}
16
+ }
17
+ return value;
18
+ }
19
+
20
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
21
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
22
+ // 2D-folded flat index: gid.y carries the high bits past the
23
+ // maxComputeWorkgroupsPerDimension limit.
24
+ let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
25
+ if (i >= params.count) {
26
+ return;
27
+ }
28
+ let xv = x[i];
29
+ let amount = shift[i];
30
+ z[i] = vec4<u32>(
31
+ shift_one(xv.x, amount.x),
32
+ shift_one(xv.y, amount.y),
33
+ shift_one(xv.z, amount.z),
34
+ shift_one(xv.w, amount.w)
35
+ );
36
+ }
build/webgpu/bitshift.wgsl.jinja ADDED
@@ -0,0 +1,93 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
2
+ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif %}) -> u32 {
3
+ {% if out_numel == 0 %}
4
+ return 0u;
5
+ {% elif op_numel == 1 %}
6
+ return 0u;
7
+ {% elif op_same %}
8
+ return out_index;
9
+ {% else %}
10
+ var offset = 0u;
11
+ {% for axis in range(outRank) %}
12
+ {% set op_axis = axis - (outRank - opRank) %}
13
+ {% if op_axis >= 0 and opShape[op_axis] != 1 %}
14
+ {% set c_stride = namespace(value=1) %}
15
+ {% for j in range(axis + 1, outRank) %}
16
+ {% set c_stride.value = c_stride.value * outShape[j] %}
17
+ {% endfor %}
18
+ {% set op_stride = namespace(value=1) %}
19
+ {% for j in range(op_axis + 1, opRank) %}
20
+ {% set op_stride.value = op_stride.value * opShape[j] %}
21
+ {% endfor %}
22
+ {% if c_stride.value == 1 %}
23
+ let coord{{ axis }} = out_index % {{ outShape[axis] }}u;
24
+ {% else %}
25
+ let coord{{ axis }} = (out_index / {{ c_stride.value }}u) % {{ outShape[axis] }}u;
26
+ {% endif %}
27
+ {% if op_stride.value == 1 %}
28
+ offset = offset + coord{{ axis }};
29
+ {% else %}
30
+ offset = offset + coord{{ axis }} * {{ op_stride.value }}u;
31
+ {% endif %}
32
+ {% endif %}
33
+ {% endfor %}
34
+ return offset;
35
+ {% endif %}
36
+ }
37
+ {%- endmacro %}{% macro broadcast_offset_call(fn_name, opShape, outShape, out_index) %}
38
+ {% set op_numel = namespace(value=1) %}
39
+ {% for d in opShape %}{% set op_numel.value = op_numel.value * d %}{% endfor %}
40
+ {% set out_numel = namespace(value=1) %}
41
+ {% for d in outShape %}{% set out_numel.value = out_numel.value * d %}{% endfor %}
42
+ {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
43
+ {%- endmacro %}{% macro broadcast_offset_fn(fn_name, opShape, opRank, outShape, outRank) %}
44
+ {% set op_numel = namespace(value=1) %}
45
+ {% for d in opShape %}
46
+ {% set op_numel.value = op_numel.value * d %}
47
+ {% endfor %}
48
+ {% set out_numel = namespace(value=1) %}
49
+ {% for d in outShape %}
50
+ {% set out_numel.value = out_numel.value * d %}
51
+ {% endfor %}
52
+ {% set op_same = namespace(value=(opRank == outRank)) %}
53
+ {% if op_same.value %}
54
+ {% for axis in range(outRank) %}
55
+ {% if opShape[axis] != outShape[axis] %}
56
+ {% set op_same.value = false %}
57
+ {% endif %}
58
+ {% endfor %}
59
+ {% endif %}
60
+ {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
61
+ {%- endmacro %}
62
+
63
+ {{ env.wgsl.resourceDeclarations }}
64
+
65
+ const BIT_WIDTH: u32 = {{ bitWidth }}u;
66
+
67
+
68
+ {{ broadcast_offset_fn("x_offset", source.xShape, source.xRank, source.zShape, source.zRank) }}
69
+
70
+ {{ broadcast_offset_fn("y_offset", source.yShape, source.yRank, source.zShape, source.zRank) }}
71
+
72
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
73
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
74
+ // 2D-folded flat index: gid.y carries the high bits past the
75
+ // maxComputeWorkgroupsPerDimension limit.
76
+ let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
77
+ if (i >= params.count) {
78
+ return;
79
+ }
80
+ let amount = shift[{{ broadcast_offset_call("y_offset", source.yShape, source.zShape, "i") }}];
81
+ var value = 0u;
82
+ if (amount < BIT_WIDTH) {
83
+ {% if leftShift %}
84
+ value = x[{{ broadcast_offset_call("x_offset", source.xShape, source.zShape, "i") }}] << amount;
85
+ {% else %}
86
+ value = x[{{ broadcast_offset_call("x_offset", source.xShape, source.zShape, "i") }}] >> amount;
87
+ {% endif %}
88
+ {% if logicalDtype == "uint8" %}
89
+ value = value & 0xffu;
90
+ {% endif %}
91
+ }
92
+ z[i] = value;
93
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,128 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "ai.onnx",
3
+ "name": "BitShift",
4
+ "sinceVersion": 11,
5
+ "description": "Performs an elementwise bitwise shift on unsigned integer tensors. `X` is shifted left or right by the amounts in `Y`, with the direction controlled by the `direction` attribute. Supports multidirectional (NumPy-style) broadcasting between `X` and `Y`.",
6
+ "inputs": [
7
+ { "role": "X", "dtype": "T", "description": "Input tensor to be shifted." },
8
+ {
9
+ "role": "Y",
10
+ "dtype": "T",
11
+ "description": "Tensor specifying the number of bit positions to shift each element of X."
12
+ }
13
+ ],
14
+ "outputs": [
15
+ {
16
+ "role": "Z",
17
+ "dtype": "T",
18
+ "rank": "max(ranks.X, ranks.Y)",
19
+ "description": "Output tensor with the same shape as the broadcast result of X and Y.",
20
+ "shape": "broadcastShape(shapes.X, shapes.Y)"
21
+ }
22
+ ],
23
+ "attributes": {},
24
+ "attributeDescriptions": {
25
+ "direction": "Direction of the bit shift: `LEFT` shifts bits toward higher significance (increasing value), while `RIGHT` shifts toward lower significance (decreasing value)."
26
+ },
27
+ "attributeConstraints": { "direction": { "required": true, "values": ["LEFT", "RIGHT"] } },
28
+ "typeConstraints": { "T": ["uint32", "uint8"] },
29
+ "args": {
30
+ "x": { "kind": "tensor", "semantic": "X", "role": "input" },
31
+ "y": { "kind": "tensor", "semantic": "Y", "role": "input" },
32
+ "z": { "kind": "tensor", "semantic": "Z", "role": "output" }
33
+ },
34
+ "tunables": { "WORKGROUP_SIZE": 256 },
35
+ "derive": {},
36
+ "constants": {
37
+ "leftShift": "attrs.direction == \"LEFT\"",
38
+ "logicalDtype": "tensorDtypes.Z",
39
+ "bitWidth": "8 if tensorDtypes.Z == \"uint8\" else 32"
40
+ },
41
+ "variants": [
42
+ {
43
+ "id": "same_shape_vec4",
44
+ "priority": 20,
45
+ "when": ["sameShape(shapes.X, shapes.Z)", "sameShape(shapes.Y, shapes.Z)", "numel(shapes.Z) > 0", "numel(shapes.Z) % 4 == 0"],
46
+ "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
47
+ "passes": [
48
+ {
49
+ "id": "main",
50
+ "name": "BitShift.vec4",
51
+ "source": { "shader": "bitshift-vec4.wgsl.jinja", "inputs": {} },
52
+ "bindings": [
53
+ {
54
+ "name": "x",
55
+ "arg": "x",
56
+ "semantic": "X",
57
+ "buffer": { "type": "read-only-storage" },
58
+ "elementType": "$vectorScalar"
59
+ },
60
+ {
61
+ "name": "shift",
62
+ "arg": "y",
63
+ "semantic": "Y",
64
+ "buffer": { "type": "read-only-storage" },
65
+ "elementType": "$vectorScalar"
66
+ },
67
+ {
68
+ "name": "z",
69
+ "arg": "z",
70
+ "semantic": "Z",
71
+ "buffer": { "type": "storage" },
72
+ "elementType": "$vectorScalar"
73
+ },
74
+ {
75
+ "name": "params",
76
+ "semantic": "kernel.params",
77
+ "buffer": { "type": "uniform" },
78
+ "struct": {
79
+ "name": "Params",
80
+ "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.Z) / 4" }]
81
+ }
82
+ }
83
+ ],
84
+ "dispatch": { "threads": "numel(shapes.Z) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
85
+ }
86
+ ]
87
+ },
88
+ {
89
+ "id": "broadcast",
90
+ "when": ["ranks.X >= 0", "ranks.Y >= 0", "ranks.Z >= 0", "ranks.X <= ranks.Z", "ranks.Y <= ranks.Z"],
91
+ "passes": [
92
+ {
93
+ "id": "main",
94
+ "name": "BitShift",
95
+ "source": {
96
+ "shader": "bitshift.wgsl.jinja",
97
+ "inputs": {
98
+ "xShape": "shapes.X",
99
+ "yShape": "shapes.Y",
100
+ "zShape": "shapes.Z",
101
+ "xRank": "ranks.X",
102
+ "yRank": "ranks.Y",
103
+ "zRank": "ranks.Z"
104
+ }
105
+ },
106
+ "bindings": [
107
+ { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
108
+ {
109
+ "name": "shift",
110
+ "arg": "y",
111
+ "semantic": "Y",
112
+ "buffer": { "type": "read-only-storage" },
113
+ "elementType": "$T"
114
+ },
115
+ { "name": "z", "arg": "z", "semantic": "Z", "buffer": { "type": "storage" }, "elementType": "$T" },
116
+ {
117
+ "name": "params",
118
+ "semantic": "kernel.params",
119
+ "buffer": { "type": "uniform" },
120
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.Z)" }] }
121
+ }
122
+ ],
123
+ "dispatch": { "threads": "numel(shapes.Z)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
124
+ }
125
+ ]
126
+ }
127
+ ]
128
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,19 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "ai.onnx.BitShift",
3
+ "id": "_ai_onnx_bitshift_webgpu_af917d0",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "Wi3Qq9s02ufgpBsKItUckvB1QuNpLdhjBCShaYh7bxM=",
11
+ "bitshift-vec4.wgsl.jinja": "JHoqmsAmesroDU6P9yHkySgeNojQ3mxsqNpgK4Pubp0=",
12
+ "bitshift.wgsl.jinja": "RvLgjOM0Ww3GczpJfd9J09sXZGeOITnqWiRMPS7AF9w=",
13
+ "manifest.json": "+tJFCJljjZ7mHnNpoU4Sfaw6tb+tlKqgodNgLDGtpd8=",
14
+ "test.json": "42uQqcJTxNDUiR8AYNnUz/IGVjwNJMv6TuTah0LUHDE="
15
+ }
16
+ },
17
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
18
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.BitShift" }
19
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,426 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.BitShift",
3
+ "cases": [
4
+ {
5
+ "name": "simple_left_u32",
6
+ "attrs": { "direction": "LEFT" },
7
+ "inputs": {
8
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
9
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
10
+ },
11
+ "outputs": { "z": { "dtype": "uint32", "shape": [3] } },
12
+ "provenance": {
13
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
14
+ "test": "BitShiftOpTest.SimpleLeft"
15
+ }
16
+ },
17
+ {
18
+ "name": "simple_right_u32",
19
+ "attrs": { "direction": "RIGHT" },
20
+ "inputs": {
21
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
22
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
23
+ },
24
+ "outputs": { "z": { "dtype": "uint32", "shape": [3] } },
25
+ "provenance": {
26
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
27
+ "test": "BitShiftOpTest.SimpleRight"
28
+ }
29
+ },
30
+ {
31
+ "name": "broadcast_y_left_u8",
32
+ "attrs": { "direction": "LEFT" },
33
+ "inputs": {
34
+ "x": { "dtype": "uint8", "shape": [3, 2], "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6] } },
35
+ "y": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [1, 2] } }
36
+ },
37
+ "outputs": { "z": { "dtype": "uint8", "shape": [3, 2] } },
38
+ "provenance": {
39
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
40
+ "test": "BitShiftOpTest.BroadcastYLeft_Uint8"
41
+ }
42
+ },
43
+ {
44
+ "name": "scalar_right_shift_by_width_u32",
45
+ "attrs": { "direction": "RIGHT" },
46
+ "inputs": {
47
+ "x": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1000] } },
48
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [32, 33, 64] } }
49
+ },
50
+ "outputs": { "z": { "dtype": "uint32", "shape": [3] } }
51
+ },
52
+ {
53
+ "name": "left_shift_by_logical_width_u8",
54
+ "attrs": { "direction": "LEFT" },
55
+ "inputs": {
56
+ "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [1, 2, 128, 255] } },
57
+ "y": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [7, 8, 1, 9] } }
58
+ },
59
+ "outputs": { "z": { "dtype": "uint8", "shape": [4] } }
60
+ },
61
+ {
62
+ "name": "backend_right_uint32_high_bits",
63
+ "attrs": { "direction": "RIGHT" },
64
+ "inputs": {
65
+ "x": {
66
+ "dtype": "uint32",
67
+ "shape": [5],
68
+ "data": { "kind": "values", "values": [4294967295, 2147483648, 4000000001, 16777217, 305419896] }
69
+ },
70
+ "y": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [1, 31, 4, 0, 8] } }
71
+ },
72
+ "outputs": { "z": { "dtype": "uint32", "shape": [5] } }
73
+ },
74
+ {
75
+ "name": "backend_left_uint32_high_bits",
76
+ "attrs": { "direction": "LEFT" },
77
+ "inputs": {
78
+ "x": {
79
+ "dtype": "uint32",
80
+ "shape": [5],
81
+ "data": { "kind": "values", "values": [1, 3, 255, 16777217, 305419896] }
82
+ },
83
+ "y": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [31, 30, 24, 8, 4] } }
84
+ },
85
+ "outputs": { "z": { "dtype": "uint32", "shape": [5] } }
86
+ },
87
+ {
88
+ "name": "scalar_left_shift_by_width_u32",
89
+ "attrs": { "direction": "LEFT" },
90
+ "inputs": {
91
+ "x": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
92
+ "y": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [30, 31, 32, 33] } }
93
+ },
94
+ "outputs": { "z": { "dtype": "uint32", "shape": [4] } }
95
+ },
96
+ {
97
+ "name": "right_shift_by_logical_width_u8",
98
+ "attrs": { "direction": "RIGHT" },
99
+ "inputs": {
100
+ "x": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [255, 128, 64, 3, 1] } },
101
+ "y": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [7, 8, 6, 1, 9] } }
102
+ },
103
+ "outputs": { "z": { "dtype": "uint8", "shape": [5] } }
104
+ },
105
+ {
106
+ "name": "ort_scalar_left_x_u32",
107
+ "provenance": {
108
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
109
+ "test": "BitShiftOpTest.ScalarLeftX"
110
+ },
111
+ "attrs": { "direction": "LEFT" },
112
+ "inputs": {
113
+ "x": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [16] } },
114
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
115
+ },
116
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
117
+ },
118
+ {
119
+ "name": "ort_scalar_left_y_u32",
120
+ "provenance": {
121
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
122
+ "test": "BitShiftOpTest.ScalarLeftY"
123
+ },
124
+ "attrs": { "direction": "LEFT" },
125
+ "inputs": {
126
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
127
+ "y": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } }
128
+ },
129
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
130
+ },
131
+ {
132
+ "name": "ort_scalar_right_x_u32",
133
+ "provenance": {
134
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
135
+ "test": "BitShiftOpTest.ScalarRightX"
136
+ },
137
+ "attrs": { "direction": "RIGHT" },
138
+ "inputs": {
139
+ "x": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [16] } },
140
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
141
+ },
142
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
143
+ },
144
+ {
145
+ "name": "ort_scalar_right_y_u32",
146
+ "provenance": {
147
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
148
+ "test": "BitShiftOpTest.ScalarRightY"
149
+ },
150
+ "attrs": { "direction": "RIGHT" },
151
+ "inputs": {
152
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
153
+ "y": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } }
154
+ },
155
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
156
+ },
157
+ {
158
+ "name": "ort_right_shift_by_bit_width_u32",
159
+ "provenance": {
160
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
161
+ "test": "BitShiftOpTest.RightShiftByBitWidth_Uint32"
162
+ },
163
+ "attrs": { "direction": "RIGHT" },
164
+ "inputs": {
165
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
166
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [32, 32, 32] } }
167
+ },
168
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
169
+ },
170
+ {
171
+ "name": "ort_right_shift_by_bit_width_u8",
172
+ "provenance": {
173
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
174
+ "test": "BitShiftOpTest.RightShiftByBitWidth_Uint32",
175
+ "notes": "Same ORT shift-by-width invariant applied to logical uint8, whose bit width is 8."
176
+ },
177
+ "attrs": { "direction": "RIGHT" },
178
+ "inputs": {
179
+ "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [255, 128, 1, 42] } },
180
+ "y": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [8, 8, 8, 8] } }
181
+ },
182
+ "outputs": {
183
+ "z": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 0, 0, 0] }, "tolerance": 0 }
184
+ }
185
+ },
186
+ {
187
+ "name": "ort_left_shift_by_bit_width_u32",
188
+ "provenance": {
189
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
190
+ "test": "BitShiftOpTest.LeftShiftByBitWidth_Uint64",
191
+ "notes": "ORT covers uint64; this framework exercises the same shift-by-width semantics with supported uint32 storage."
192
+ },
193
+ "attrs": { "direction": "LEFT" },
194
+ "inputs": {
195
+ "x": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [1000, 255, 1, 42] } },
196
+ "y": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [32, 32, 32, 32] } }
197
+ },
198
+ "outputs": { "z": { "dtype": "uint32", "shape": [4], "tolerance": 0 } }
199
+ },
200
+ {
201
+ "name": "ort_left_shift_by_bit_width_u8",
202
+ "provenance": {
203
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
204
+ "test": "BitShiftOpTest.LeftShiftByBitWidth_Uint64",
205
+ "notes": "ORT covers uint64; this framework exercises the same shift-by-width invariant with logical uint8."
206
+ },
207
+ "attrs": { "direction": "LEFT" },
208
+ "inputs": {
209
+ "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [255, 128, 1, 42] } },
210
+ "y": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [8, 8, 8, 8] } }
211
+ },
212
+ "outputs": {
213
+ "z": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 0, 0, 0] }, "tolerance": 0 }
214
+ }
215
+ },
216
+ {
217
+ "name": "ort_right_shift_by_more_than_bit_width_u32",
218
+ "provenance": {
219
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
220
+ "test": "BitShiftOpTest.RightShiftByMoreThanBitWidth_Uint64",
221
+ "notes": "ORT covers uint64; this framework exercises the same shift-past-width semantics with supported uint32 storage."
222
+ },
223
+ "attrs": { "direction": "RIGHT" },
224
+ "inputs": {
225
+ "x": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1000, 42] } },
226
+ "y": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [33, 64] } }
227
+ },
228
+ "outputs": { "z": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
229
+ },
230
+ {
231
+ "name": "ort_right_shift_by_more_than_bit_width_u8",
232
+ "provenance": {
233
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
234
+ "test": "BitShiftOpTest.RightShiftByMoreThanBitWidth_Uint64",
235
+ "notes": "ORT covers uint64; this framework exercises the same shift-past-width invariant with logical uint8."
236
+ },
237
+ "attrs": { "direction": "RIGHT" },
238
+ "inputs": {
239
+ "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [255, 128, 1, 42] } },
240
+ "y": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [9, 16, 31, 255] } }
241
+ },
242
+ "outputs": {
243
+ "z": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 0, 0, 0] }, "tolerance": 0 }
244
+ }
245
+ },
246
+ {
247
+ "name": "ort_broadcast_x_right_u8",
248
+ "provenance": {
249
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
250
+ "test": "BitShiftOpTest.BroadcastXRight_Uint8"
251
+ },
252
+ "attrs": { "direction": "RIGHT" },
253
+ "inputs": {
254
+ "x": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [64, 32] } },
255
+ "y": { "dtype": "uint8", "shape": [3, 2], "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6] } }
256
+ },
257
+ "outputs": { "z": { "dtype": "uint8", "shape": [3, 2], "tolerance": 0 } }
258
+ },
259
+ {
260
+ "name": "ort_broadcast_x_right_u32",
261
+ "provenance": {
262
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
263
+ "test": "BitShiftOpTest.BroadcastXRight",
264
+ "notes": "ORT uses uint64 values; this framework exercises the same broadcast pattern with supported uint32 storage."
265
+ },
266
+ "attrs": { "direction": "RIGHT" },
267
+ "inputs": {
268
+ "x": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [64, 32] } },
269
+ "y": { "dtype": "uint32", "shape": [3, 2], "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6] } }
270
+ },
271
+ "outputs": { "z": { "dtype": "uint32", "shape": [3, 2], "tolerance": 0 } }
272
+ },
273
+ {
274
+ "name": "ort_broadcast_y_left_u32",
275
+ "provenance": {
276
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
277
+ "test": "BitShiftOpTest.BroadcastYLeft",
278
+ "notes": "ORT uses uint64 values; this framework exercises the same broadcast pattern with supported uint32 storage."
279
+ },
280
+ "attrs": { "direction": "LEFT" },
281
+ "inputs": {
282
+ "x": { "dtype": "uint32", "shape": [3, 2], "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6] } },
283
+ "y": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } }
284
+ },
285
+ "outputs": { "z": { "dtype": "uint32", "shape": [3, 2], "tolerance": 0 } }
286
+ },
287
+ {
288
+ "name": "onnx_backend_bitshift_left_uint32",
289
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitshift_left_uint32" },
290
+ "attrs": { "direction": "LEFT" },
291
+ "inputs": {
292
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
293
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
294
+ },
295
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
296
+ },
297
+ {
298
+ "name": "onnx_backend_bitshift_right_uint32",
299
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitshift_right_uint32" },
300
+ "attrs": { "direction": "RIGHT" },
301
+ "inputs": {
302
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
303
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
304
+ },
305
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
306
+ },
307
+ {
308
+ "name": "onnx_backend_bitshift_left_uint8",
309
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitshift_left_uint8" },
310
+ "attrs": { "direction": "LEFT" },
311
+ "inputs": {
312
+ "x": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
313
+ "y": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
314
+ },
315
+ "outputs": { "z": { "dtype": "uint8", "shape": [3], "tolerance": 0 } }
316
+ },
317
+ {
318
+ "name": "onnx_backend_bitshift_right_uint8",
319
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitshift_right_uint8" },
320
+ "attrs": { "direction": "RIGHT" },
321
+ "inputs": {
322
+ "x": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [16, 4, 1] } },
323
+ "y": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } }
324
+ },
325
+ "outputs": { "z": { "dtype": "uint8", "shape": [3], "tolerance": 0 } }
326
+ },
327
+ {
328
+ "name": "vec4_right_u32_with_overwide_shifts",
329
+ "attrs": { "direction": "RIGHT" },
330
+ "inputs": {
331
+ "x": {
332
+ "dtype": "uint32",
333
+ "shape": [8],
334
+ "data": { "kind": "values", "values": [4294967295, 1024, 7, 4294967295, 256, 65535, 8, 12345678] }
335
+ },
336
+ "y": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [4, 10, 1, 32, 33, 0, 3, 8] } }
337
+ },
338
+ "outputs": {
339
+ "z": {
340
+ "dtype": "uint32",
341
+ "shape": [8],
342
+ "data": { "kind": "values", "values": [268435455, 1, 3, 0, 0, 65535, 1, 48225] },
343
+ "tolerance": 0
344
+ }
345
+ }
346
+ },
347
+ {
348
+ "name": "vec4_right_u8_with_overwide_shifts",
349
+ "attrs": { "direction": "RIGHT" },
350
+ "inputs": {
351
+ "x": {
352
+ "dtype": "uint8",
353
+ "shape": [8],
354
+ "data": { "kind": "values", "values": [255, 128, 200, 7, 255, 16, 3, 9] }
355
+ },
356
+ "y": { "dtype": "uint8", "shape": [8], "data": { "kind": "values", "values": [1, 7, 8, 0, 9, 2, 1, 3] } }
357
+ },
358
+ "outputs": {
359
+ "z": {
360
+ "dtype": "uint8",
361
+ "shape": [8],
362
+ "data": { "kind": "values", "values": [127, 1, 0, 7, 0, 4, 1, 1] },
363
+ "tolerance": 0
364
+ }
365
+ }
366
+ },
367
+ {
368
+ "name": "empty_input_zero_dim",
369
+ "attrs": { "direction": "LEFT" },
370
+ "inputs": {
371
+ "x": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } },
372
+ "y": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
373
+ },
374
+ "outputs": { "z": { "dtype": "uint32", "shape": [0], "tolerance": 0 } }
375
+ },
376
+ {
377
+ "name": "broadcast_2d_u8_non_mult4_left",
378
+ "attrs": { "direction": "LEFT" },
379
+ "inputs": {
380
+ "x": { "dtype": "uint8", "shape": [1, 3], "data": { "kind": "values", "values": [1, 2, 3] } },
381
+ "y": { "dtype": "uint8", "shape": [2, 1], "data": { "kind": "values", "values": [1, 2] } }
382
+ },
383
+ "outputs": {
384
+ "z": {
385
+ "dtype": "uint8",
386
+ "shape": [2, 3],
387
+ "tolerance": 0,
388
+ "data": { "kind": "values", "values": [2, 4, 6, 4, 8, 12] }
389
+ }
390
+ }
391
+ },
392
+ {
393
+ "name": "empty_z_broadcast_scalar_x",
394
+ "attrs": { "direction": "RIGHT" },
395
+ "inputs": {
396
+ "x": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [255] } },
397
+ "y": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
398
+ },
399
+ "outputs": {
400
+ "z": { "dtype": "uint32", "shape": [0], "tolerance": 0, "data": { "kind": "values", "values": [] } }
401
+ }
402
+ },
403
+ {
404
+ "name": "rank7_broadcast_scalar_tail",
405
+ "attrs": { "direction": "LEFT" },
406
+ "inputs": {
407
+ "x": { "dtype": "uint32", "shape": [2, 1, 2, 1, 2, 1, 3], "data": { "kind": "cycle", "values": [1, 3, 7, 15] } },
408
+ "y": { "dtype": "uint32", "shape": [1, 2, 1, 2, 1, 2, 1], "data": { "kind": "cycle", "values": [0, 1, 2] } }
409
+ },
410
+ "outputs": { "z": { "dtype": "uint32", "shape": [2, 2, 2, 2, 2, 2, 3], "tolerance": 0 } }
411
+ },
412
+ {
413
+ "name": "rank8_broadcast_alternating",
414
+ "attrs": { "direction": "LEFT" },
415
+ "inputs": {
416
+ "x": {
417
+ "dtype": "uint32",
418
+ "shape": [2, 1, 2, 1, 2, 1, 2, 3],
419
+ "data": { "kind": "cycle", "values": [1, 3, 5, 7] }
420
+ },
421
+ "y": { "dtype": "uint32", "shape": [1, 2, 1, 2, 1, 2, 1, 1], "data": { "kind": "cycle", "values": [0, 1, 2] } }
422
+ },
423
+ "outputs": { "z": { "dtype": "uint32", "shape": [2, 2, 2, 2, 2, 2, 2, 3], "tolerance": 0 } }
424
+ }
425
+ ]
426
+ }