Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.BitShift

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

sync 91d990483a17

Browse files
Files changed (7) hide show
  1. README.md +14 -10
  2. build/webgpu/bench.json +0 -1
  3. build/webgpu/bitshift-vec4.wgsl.jinja +3 -3
  4. build/webgpu/bitshift.wgsl.jinja +8 -8
  5. build/webgpu/manifest.json +38 -92
  6. build/webgpu/metadata.json +11 -8
  7. build/webgpu/test.json +16 -14
README.md CHANGED
@@ -18,16 +18,16 @@ See the [ONNX `BitShift` spec](https://onnx.ai/onnx/operators/onnx__BitShift.htm
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
 
@@ -45,7 +45,7 @@ Attributes and default values (overridable per request):
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
@@ -54,10 +54,14 @@ Attributes and default values (overridable per request):
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
 
@@ -65,7 +69,7 @@ Replace each `*Data` placeholder with a typed array containing the corresponding
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
  ```
 
18
 
19
  ## Inputs
20
 
21
+ | Name | Upstream name | 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 | Upstream name | 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
 
 
45
 
46
  ## Files
47
 
48
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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
 
54
 
55
  ## Use with `@huggingface/kernels`
56
 
57
+ ```sh
58
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.2
59
+ ```
60
+
61
+ Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
62
 
63
  The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
64
+ It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
65
 
66
  Replace each `*Data` placeholder with a typed array containing the corresponding input data.
67
 
 
69
  import { getKernel } from "@huggingface/kernels";
70
 
71
  const kernel = await getKernel("webgpu-kernels/ai.onnx.BitShift", { version: 1 });
72
+ const { z } = await kernel({ x: { data: xData, shape: [1] }, y: { data: yData, shape: [3] } }, {
73
+ attrs: { direction: "LEFT" },
74
  });
75
  ```
build/webgpu/bench.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "ai.onnx.BitShift",
3
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
  "cases": [
5
  {
 
1
  {
 
2
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
3
  "cases": [
4
  {
build/webgpu/bitshift-vec4.wgsl.jinja CHANGED
@@ -18,10 +18,10 @@ fn shift_one(in_value: u32, amount: u32) -> u32 {
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
  }
 
18
  }
19
 
20
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
21
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
22
  // 2D-folded flat index: gid.y carries the high bits past the
23
+ // per-axis dispatch fold width.
24
+ let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
25
  if (i >= params.count) {
26
  return;
27
  }
build/webgpu/bitshift.wgsl.jinja CHANGED
@@ -65,25 +65,25 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
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;
 
65
  const BIT_WIDTH: u32 = {{ bitWidth }}u;
66
 
67
 
68
+ {{ broadcast_offset_fn("x_offset", xShape, xRank, zShape, zRank) }}
69
 
70
+ {{ broadcast_offset_fn("y_offset", yShape, yRank, zShape, zRank) }}
71
 
72
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
73
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
74
  // 2D-folded flat index: gid.y carries the high bits past the
75
+ // per-axis dispatch fold width.
76
+ let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
77
  if (i >= params.count) {
78
  return;
79
  }
80
+ let amount = shift[{{ broadcast_offset_call("y_offset", yShape, zShape, "i") }}];
81
  var value = 0u;
82
  if (amount < BIT_WIDTH) {
83
  {% if leftShift %}
84
+ value = x[{{ broadcast_offset_call("x_offset", xShape, zShape, "i") }}] << amount;
85
  {% else %}
86
+ value = x[{{ broadcast_offset_call("x_offset", xShape, zShape, "i") }}] >> amount;
87
  {% endif %}
88
  {% if logicalDtype == "uint8" %}
89
  value = value & 0xffu;
build/webgpu/manifest.json CHANGED
@@ -2,125 +2,71 @@
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
  }
 
2
  "domain": "ai.onnx",
3
  "name": "BitShift",
4
  "sinceVersion": 11,
5
+ "inputs": { "x": { "onnx": "X", "dtype": "T" }, "y": { "onnx": "Y", "dtype": "T" } },
6
+ "outputs": {
7
+ "z": { "onnx": "Z", "dtype": "T", "rank": "max(ranks.x, ranks.y)", "shape": "broadcastShape(shapes.x, shapes.y)" }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
8
  },
9
+ "attributes": { "direction": {} },
10
  "attributeConstraints": { "direction": { "required": true, "values": ["LEFT", "RIGHT"] } },
11
  "typeConstraints": { "T": ["uint32", "uint8"] },
12
+ "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
13
+ "derive": {
 
 
 
 
 
 
14
  "leftShift": "attrs.direction == \"LEFT\"",
15
+ "logicalDtype": "tensorDtypes.z",
16
+ "bitWidth": "8 if tensorDtypes.z == \"uint8\" else 32"
17
  },
18
  "variants": [
19
  {
20
  "id": "same_shape_vec4",
21
  "priority": 20,
22
+ "when": ["sameShape(shapes.x, shapes.z)", "sameShape(shapes.y, shapes.z)", "numel(shapes.z) > 0", "numel(shapes.z) % 4 == 0"],
23
+ "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
24
  "passes": [
25
  {
26
  "id": "main",
27
  "name": "BitShift.vec4",
28
+ "shader": "bitshift-vec4.wgsl.jinja",
29
  "bindings": [
30
+ { "arg": "x", "elementType": "$vectorScalar" },
31
+ { "arg": "y", "name": "shift", "elementType": "$vectorScalar" },
32
+ { "arg": "z", "elementType": "$vectorScalar" },
33
+ { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.z) / 4" }] }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
34
  ],
35
+ "dispatch": {
36
+ "x": "min(ceilDiv((numel(shapes.z) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
37
+ "y": "ceilDiv(ceilDiv((numel(shapes.z) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
38
+ "z": 1
39
+ }
40
  }
41
  ]
42
  },
43
  {
44
  "id": "broadcast",
45
+ "when": ["ranks.x >= 0", "ranks.y >= 0", "ranks.z >= 0", "ranks.x <= ranks.z", "ranks.y <= ranks.z"],
46
  "passes": [
47
  {
48
  "id": "main",
49
  "name": "BitShift",
50
+ "shader": "bitshift.wgsl.jinja",
51
+ "derive": {
52
+ "xShape": "shapes.x",
53
+ "yShape": "shapes.y",
54
+ "zShape": "shapes.z",
55
+ "xRank": "ranks.x",
56
+ "yRank": "ranks.y",
57
+ "zRank": "ranks.z"
 
 
58
  },
59
  "bindings": [
60
+ "x",
61
+ { "arg": "y", "name": "shift" },
62
+ "z",
63
+ { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.z)" }] }
 
 
 
 
 
 
 
 
 
 
 
64
  ],
65
+ "dispatch": {
66
+ "x": "min(ceilDiv((numel(shapes.z)), (tunables.WORKGROUP_SIZE)), 65535)",
67
+ "y": "ceilDiv(ceilDiv((numel(shapes.z)), (tunables.WORKGROUP_SIZE)), 65535)",
68
+ "z": 1
69
+ }
70
  }
71
  ]
72
  }
build/webgpu/metadata.json CHANGED
@@ -1,19 +1,22 @@
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
  }
 
1
  {
2
  "name": "ai.onnx.BitShift",
3
+ "id": "_ai_onnx_bitshift_webgpu_b738951",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "pCMEZApV+9hKgXsgY8QFt30AzJB6jYpkK9vUoboxcv4=",
11
+ "bitshift-vec4.wgsl.jinja": "t6oiDS+UC3rc8/U4pGJABtsCNpNI5Uh+EFEePIcoEUc=",
12
+ "bitshift.wgsl.jinja": "w5jaDNLdgC/51OhsXtWNIdZTBf7mFnx/et01uoYsFQY=",
13
+ "manifest.json": "I9WKns2iHR6NHc7PkEz4QzVO8xQibqM8ZZdedOSA4v0=",
14
+ "test.json": "AaUcb5BUmYyvUXpso+6EIwZo3/LLieW55+9CV2pPXlg="
15
  }
16
  },
17
+ "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
18
+ "webgpu": {
19
+ "manifestSpec": "2.0",
20
+ "variants": { "same_shape_vec4": ["bitshift-vec4.wgsl.jinja"], "broadcast": ["bitshift.wgsl.jinja"] }
21
+ }
22
  }
build/webgpu/test.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "ai.onnx.BitShift",
3
  "cases": [
4
  {
5
  "name": "simple_left_u32",
@@ -45,9 +44,12 @@
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",
@@ -188,7 +190,7 @@
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": {
@@ -202,7 +204,7 @@
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": {
@@ -218,29 +220,29 @@
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
  {
@@ -261,7 +263,7 @@
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": {
@@ -275,7 +277,7 @@
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": {
 
1
  {
 
2
  "cases": [
3
  {
4
  "name": "simple_left_u32",
 
44
  "attrs": { "direction": "RIGHT" },
45
  "inputs": {
46
  "x": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1000] } },
47
+ "y": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [32, 33, 64, 4] } }
48
  },
49
+ "outputs": { "z": { "dtype": "uint32", "shape": [4] } },
50
+ "provenance": {
51
+ "notes": "Shifts at and past the 32-bit width must clear the value; the trailing in-range shift is the control that a kernel zeroing every shift would fail."
52
+ }
53
  },
54
  {
55
  "name": "left_shift_by_logical_width_u8",
 
190
  "provenance": {
191
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
192
  "test": "BitShiftOpTest.LeftShiftByBitWidth_Uint64",
193
+ "notes": "Projects the uint64 source case to supported uint32 storage while preserving shift-by-width semantics."
194
  },
195
  "attrs": { "direction": "LEFT" },
196
  "inputs": {
 
204
  "provenance": {
205
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
206
  "test": "BitShiftOpTest.LeftShiftByBitWidth_Uint64",
207
+ "notes": "Projects the uint64 source case to logical uint8 storage while preserving the shift-by-width invariant."
208
  },
209
  "attrs": { "direction": "LEFT" },
210
  "inputs": {
 
220
  "provenance": {
221
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
222
  "test": "BitShiftOpTest.RightShiftByMoreThanBitWidth_Uint64",
223
+ "notes": "Diverges from the upstream test's inputs (inputs.x shape [2] values [1000, 42] -> shape [3] values [1000, 42, 1000]; inputs.y shape [2] values [33, 64] -> shape [3] values [33, 64, 3]); the expected output is recomputed by the CPU reference for the new inputs. Projects the uint64 source case to supported uint32 storage while preserving shift-past-width semantics, with a trailing in-range shift as the control that a kernel zeroing every shift would fail."
224
  },
225
  "attrs": { "direction": "RIGHT" },
226
  "inputs": {
227
+ "x": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1000, 42, 1000] } },
228
+ "y": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [33, 64, 3] } }
229
  },
230
+ "outputs": { "z": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
231
  },
232
  {
233
  "name": "ort_right_shift_by_more_than_bit_width_u8",
234
  "provenance": {
235
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
236
  "test": "BitShiftOpTest.RightShiftByMoreThanBitWidth_Uint64",
237
+ "notes": "Diverges from the upstream test's inputs (inputs.y values [9, 16, 31, 255] -> values [9, 16, 31, 3]); the expected output is recomputed by the CPU reference for the new inputs. Projects the uint64 source case to logical uint8 storage while preserving the shift-past-width invariant, with a trailing in-range shift as the control that a kernel zeroing every shift would fail."
238
  },
239
  "attrs": { "direction": "RIGHT" },
240
  "inputs": {
241
  "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [255, 128, 1, 42] } },
242
+ "y": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [9, 16, 31, 3] } }
243
  },
244
  "outputs": {
245
+ "z": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 0, 0, 5] }, "tolerance": 0 }
246
  }
247
  },
248
  {
 
263
  "provenance": {
264
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
265
  "test": "BitShiftOpTest.BroadcastXRight",
266
+ "notes": "Projects the uint64 source values to supported uint32 storage while preserving the broadcast pattern."
267
  },
268
  "attrs": { "direction": "RIGHT" },
269
  "inputs": {
 
277
  "provenance": {
278
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
279
  "test": "BitShiftOpTest.BroadcastYLeft",
280
+ "notes": "Projects the uint64 source values to supported uint32 storage while preserving the broadcast pattern."
281
  },
282
  "attrs": { "direction": "LEFT" },
283
  "inputs": {