Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.BitCast

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
9d0cf92
·
verified ·
1 Parent(s): 9e0b43e

sync 2e7068faf55e

Browse files
Files changed (6) hide show
  1. README.md +67 -0
  2. build/webgpu/bench.json +30 -0
  3. build/webgpu/bitcast.wgsl.jinja +38 -0
  4. build/webgpu/manifest.json +128 -0
  5. build/webgpu/metadata.json +18 -0
  6. build/webgpu/test.json +385 -0
README.md CHANGED
@@ -1,3 +1,70 @@
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.BitCast
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 26
12
+
13
+ ## Description
14
+
15
+ Reinterprets the raw bit pattern of a tensor as a different data type without any value conversion. The target type must have the same bit-width as the input type, and the output tensor has the same shape as the input.
16
+
17
+ See the [ONNX `BitCast` spec](https://onnx.ai/onnx/operators/onnx__BitCast.html) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `input` | `input` | `T` | — | — | Input tensor to be bitwise reinterpreted. | required |
24
+
25
+ ## Outputs
26
+
27
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
28
+ | --- | --- | --- | --- | --- | --- | --- |
29
+ | `output` | `output` | `U` | same as `input` | same as `input` | Output tensor with the same shape as the input, reinterpreted as the target type. | required |
30
+
31
+ ## Attributes
32
+
33
+ Attributes and default values (overridable per request):
34
+
35
+ | Attribute | Default | Description |
36
+ | --- | --- | --- |
37
+ | `to` | — | Required TensorProto DataType enum integer naming the output dtype; the target type must have the same bit-width as the input type. |
38
+
39
+ ## Type constraints
40
+
41
+ | Variable | Allowed dtypes |
42
+ | --- | --- |
43
+ | `T` | `float32`, `int8`, `int32`, `uint8`, `uint32` |
44
+ | `U` | `float32`, `int8`, `int32`, `uint8`, `uint32` |
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
+ - [`bitcast.wgsl.jinja`](build/webgpu/bitcast.wgsl.jinja)
53
+
54
+ ## Use with `@huggingface/kernels`
55
+
56
+ The loader derives every required output's shape and logical dtype from the manifest contract and this call.
57
+ It then allocates the result tensors automatically.
58
+
59
+ The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
60
+
61
+ Replace each `*Data` placeholder with a typed array containing the corresponding input data.
62
+
63
+ ```js
64
+ import { getKernel } from "@huggingface/kernels";
65
+
66
+ const kernel = await getKernel("webgpu-kernels/ai.onnx.BitCast", { version: 1 });
67
+ const { output } = await kernel({ input: { data: inputData, shape: [] } }, {
68
+ attrs: { to: 6 },
69
+ });
70
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,30 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.BitCast",
3
+ "cases": [
4
+ {
5
+ "name": "bitcast-f32-to-u32-33m",
6
+ "preset": "stress",
7
+ "provenance": { "notes": "Stress-only capacity case: input plus output occupy exactly 256 MiB of GPU storage." },
8
+ "attrs": { "to": 12 },
9
+ "inputs": { "input": { "dtype": "float32", "shape": [4096, 8192], "dist": "normal", "seed": 915, "scale": 1 } },
10
+ "outputs": { "output": { "dtype": "uint32", "shape": [4096, 8192] } },
11
+ "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "4096 * 8192 * 4 * 2" }] }
12
+ },
13
+ {
14
+ "name": "bitcast-f32-to-u32-1m",
15
+ "preset": "smoke",
16
+ "attrs": { "to": 12 },
17
+ "inputs": { "input": { "dtype": "float32", "shape": [1048576], "dist": "normal", "seed": 913, "scale": 1 } },
18
+ "outputs": { "output": { "dtype": "uint32", "shape": [1048576] } },
19
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.input) + numel(shapes.output)) * 4" }] }
20
+ },
21
+ {
22
+ "name": "bitcast-f32-to-u32-scalar-fallback-4m",
23
+ "preset": "edge",
24
+ "attrs": { "to": 12 },
25
+ "inputs": { "input": { "dtype": "float32", "shape": [4194305], "dist": "normal", "seed": 917, "scale": 1 } },
26
+ "outputs": { "output": { "dtype": "uint32", "shape": [4194305] } },
27
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "4194305 * 4 * 2" }] }
28
+ }
29
+ ]
30
+ }
build/webgpu/bitcast.wgsl.jinja ADDED
@@ -0,0 +1,38 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {{ env.wgsl.resourceDeclarations }}
2
+
3
+ {% set vectorized = source.vectorized if source.vectorized is defined else false %}
4
+ // Each invocation reinterprets one 32-bit slot, or four contiguous slots on
5
+ // the vectorized route. Framework-widened int8/uint8 values retain their
6
+ // explicit low-byte masking and sign extension.
7
+ @compute @workgroup_size({{ source.workgroupSize }})
8
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
9
+ // The flat dispatch is folded across x/y at the device limit, so gid.y
10
+ // carries the high portion of the slot index.
11
+ let i = gid.x + gid.y * nwg.x * {{ source.workgroupSize }}u;
12
+ if (i >= params.count) {
13
+ return;
14
+ }
15
+ {% if inScalar == outScalar %}
16
+ output[i] = input[i];
17
+ {% elif inputIsInt8 and outputIsUint8 %}
18
+ {% if vectorized %}
19
+ output[i] = bitcast<vec4<u32>>(input[i]) & vec4<u32>(0xffu);
20
+ {% else %}
21
+ output[i] = bitcast<u32>(input[i]) & 0xffu;
22
+ {% endif %}
23
+ {% elif inputIsUint8 and outputIsInt8 %}
24
+ {% if vectorized %}
25
+ let byte_value = input[i] & vec4<u32>(0xffu);
26
+ output[i] = vec4<i32>(byte_value) - select(vec4<i32>(0), vec4<i32>(256), byte_value >= vec4<u32>(128u));
27
+ {% else %}
28
+ let byte_value = input[i] & 0xffu;
29
+ output[i] = i32(byte_value) - select(0i, 256i, byte_value >= 128u);
30
+ {% endif %}
31
+ {% else %}
32
+ {% if vectorized %}
33
+ output[i] = bitcast<vec4<{{ outScalar }}>>(input[i]);
34
+ {% else %}
35
+ output[i] = bitcast<{{ outScalar }}>(input[i]);
36
+ {% endif %}
37
+ {% endif %}
38
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,128 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "ai.onnx",
3
+ "name": "BitCast",
4
+ "sinceVersion": 26,
5
+ "description": "Reinterprets the raw bit pattern of a tensor as a different data type without any value conversion. The target type must have the same bit-width as the input type, and the output tensor has the same shape as the input.",
6
+ "inputs": [{ "role": "input", "dtype": "T", "description": "Input tensor to be bitwise reinterpreted." }],
7
+ "outputs": [
8
+ {
9
+ "role": "output",
10
+ "dtype": "U",
11
+ "rank": "ranks.input",
12
+ "description": "Output tensor with the same shape as the input, reinterpreted as the target type.",
13
+ "shape": "shapes.input"
14
+ }
15
+ ],
16
+ "attributes": {},
17
+ "attributeDescriptions": {
18
+ "to": "Required TensorProto DataType enum integer naming the output dtype; the target type must have the same bit-width as the input type."
19
+ },
20
+ "attributeConstraints": { "to": { "required": true } },
21
+ "typeConstraints": {
22
+ "T": ["float32", "int8", "int32", "uint8", "uint32"],
23
+ "U": ["float32", "int8", "int32", "uint8", "uint32"]
24
+ },
25
+ "args": {
26
+ "input": { "kind": "tensor", "semantic": "input", "role": "input" },
27
+ "output": { "kind": "tensor", "semantic": "output", "role": "output" }
28
+ },
29
+ "tunables": { "WORKGROUP_SIZE": 256 },
30
+ "derive": {
31
+ "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
32
+ "bitcastShapeOk": "ranks.input >= 0 and ranks.output == ranks.input and numel(shapes.input) == numel(shapes.output)",
33
+ "bitcastTypeOk": "attrs.to == onnxDtypeCode(logicalDtypes.U) and (((tensorDtypes.input == \"int8\" or tensorDtypes.input == \"uint8\") and (tensorDtypes.output == \"int8\" or tensorDtypes.output == \"uint8\")) or (tensorDtypes.input != \"int8\" and tensorDtypes.input != \"uint8\" and tensorDtypes.output != \"int8\" and tensorDtypes.output != \"uint8\"))",
34
+ "bitcastWorkgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
35
+ "bitcastDispatchFits": "ceilDiv(ceilDiv(numel(shapes.output), bitcastWorkgroupSize), device.limits.maxComputeWorkgroupsPerDimension) <= device.limits.maxComputeWorkgroupsPerDimension",
36
+ "bitcastBaseOk": "bitcastShapeOk and bitcastTypeOk and bitcastDispatchFits"
37
+ },
38
+ "constants": {
39
+ "inScalar": "dtypes.T",
40
+ "outScalar": "dtypes.U",
41
+ "inputIsInt8": "tensorDtypes.input == \"int8\"",
42
+ "inputIsUint8": "tensorDtypes.input == \"uint8\"",
43
+ "outputIsInt8": "tensorDtypes.output == \"int8\"",
44
+ "outputIsUint8": "tensorDtypes.output == \"uint8\""
45
+ },
46
+ "bindingSets": {
47
+ "vec4Slots": [
48
+ {
49
+ "name": "input",
50
+ "arg": "input",
51
+ "semantic": "input",
52
+ "buffer": { "type": "read-only-storage" },
53
+ "elementType": "$inVec4"
54
+ },
55
+ {
56
+ "name": "output",
57
+ "arg": "output",
58
+ "semantic": "output",
59
+ "buffer": { "type": "storage" },
60
+ "elementType": "$outVec4"
61
+ },
62
+ {
63
+ "name": "params",
64
+ "semantic": "kernel.params",
65
+ "buffer": { "type": "uniform" },
66
+ "struct": {
67
+ "name": "Params",
68
+ "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
69
+ }
70
+ }
71
+ ],
72
+ "scalarSlots": [
73
+ {
74
+ "name": "input",
75
+ "arg": "input",
76
+ "semantic": "input",
77
+ "buffer": { "type": "read-only-storage" },
78
+ "elementType": "$inScalar"
79
+ },
80
+ {
81
+ "name": "output",
82
+ "arg": "output",
83
+ "semantic": "output",
84
+ "buffer": { "type": "storage" },
85
+ "elementType": "$outScalar"
86
+ },
87
+ {
88
+ "name": "params",
89
+ "semantic": "kernel.params",
90
+ "buffer": { "type": "uniform" },
91
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
92
+ }
93
+ ]
94
+ },
95
+ "variants": [
96
+ {
97
+ "id": "slot32_vec4",
98
+ "priority": 20,
99
+ "when": ["numel(shapes.input) % 4 == 0", "bitcastBaseOk"],
100
+ "constants": { "inVec4": "\"vec4<\" ~ dtypes.T ~ \">\"", "outVec4": "\"vec4<\" ~ dtypes.U ~ \">\"" },
101
+ "passes": [
102
+ {
103
+ "id": "main",
104
+ "name": "BitCast.vec4",
105
+ "source": {
106
+ "shader": "bitcast.wgsl.jinja",
107
+ "inputs": { "vectorized": true, "workgroupSize": "bitcastWorkgroupSize" }
108
+ },
109
+ "bindings": "vec4Slots",
110
+ "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "bitcastWorkgroupSize" }
111
+ }
112
+ ]
113
+ },
114
+ {
115
+ "id": "slot32",
116
+ "when": ["bitcastBaseOk"],
117
+ "passes": [
118
+ {
119
+ "id": "main",
120
+ "name": "BitCast",
121
+ "source": { "shader": "bitcast.wgsl.jinja", "inputs": { "workgroupSize": "bitcastWorkgroupSize" } },
122
+ "bindings": "scalarSlots",
123
+ "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "bitcastWorkgroupSize" }
124
+ }
125
+ ]
126
+ }
127
+ ]
128
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,18 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "ai.onnx.BitCast",
3
+ "id": "_ai_onnx_bitcast_webgpu_953c2be",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "lVPh5n2BwzVpRwCSVTTJRFxMXRTPmVeRM1Iy1ZVmjag=",
11
+ "bitcast.wgsl.jinja": "lXNQpNlvG/xrJll5JpnrFRGYtGajmNiP3uv9e846ObQ=",
12
+ "manifest.json": "gj3f/7xbrjxRy/dTZy9BN+nmJNOqanpBjiwjeIoFuFk=",
13
+ "test.json": "bz92JkCXkvPo9upgVfsjbWlUBWrz1pIqKHUyZS/K2uA="
14
+ }
15
+ },
16
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
17
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.BitCast" }
18
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,385 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.BitCast",
3
+ "cases": [
4
+ {
5
+ "name": "float32_to_int32",
6
+ "provenance": {
7
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
8
+ "test": "BitCastTest.Float32ToInt32"
9
+ },
10
+ "attrs": { "to": 6 },
11
+ "inputs": {
12
+ "input": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.5] } }
13
+ },
14
+ "outputs": { "output": { "dtype": "int32", "shape": [4] } }
15
+ },
16
+ {
17
+ "name": "int32_to_float32",
18
+ "provenance": {
19
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
20
+ "test": "BitCastTest.Int32ToFloat32"
21
+ },
22
+ "attrs": { "to": 1 },
23
+ "inputs": {
24
+ "input": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 1065353216, 1073741824] } }
25
+ },
26
+ "outputs": { "output": { "dtype": "float32", "shape": [3] } }
27
+ },
28
+ {
29
+ "name": "uint32_to_float32",
30
+ "provenance": {
31
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
32
+ "test": "BitCastTest.UInt32ToFloat32"
33
+ },
34
+ "attrs": { "to": 1 },
35
+ "inputs": {
36
+ "input": {
37
+ "dtype": "uint32",
38
+ "shape": [3],
39
+ "data": { "kind": "values", "values": [0, 1065353216, 1073741824] }
40
+ }
41
+ },
42
+ "outputs": { "output": { "dtype": "float32", "shape": [3] } }
43
+ },
44
+ {
45
+ "name": "float32_to_uint32_rank2",
46
+ "attrs": { "to": 12 },
47
+ "inputs": {
48
+ "input": {
49
+ "dtype": "float32",
50
+ "shape": [2, 3],
51
+ "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
52
+ }
53
+ },
54
+ "outputs": { "output": { "dtype": "uint32", "shape": [2, 3] } }
55
+ },
56
+ {
57
+ "name": "scalar_float32_to_int32",
58
+ "provenance": {
59
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
60
+ "test": "BitCastTest.ScalarTensor"
61
+ },
62
+ "attrs": { "to": 6 },
63
+ "inputs": { "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [42.0] } } },
64
+ "outputs": { "output": { "dtype": "int32", "shape": [] } }
65
+ },
66
+ {
67
+ "name": "ort_float32_to_int32_rank2",
68
+ "provenance": {
69
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
70
+ "test": "BitCastTest.Float32ToInt32_2D"
71
+ },
72
+ "attrs": { "to": 6 },
73
+ "inputs": {
74
+ "input": {
75
+ "dtype": "float32",
76
+ "shape": [2, 3],
77
+ "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
78
+ }
79
+ },
80
+ "outputs": { "output": { "dtype": "int32", "shape": [2, 3] } }
81
+ },
82
+ {
83
+ "name": "ort_empty_float32_to_int32",
84
+ "provenance": {
85
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
86
+ "test": "BitCastTest.EmptyTensor"
87
+ },
88
+ "attrs": { "to": 6 },
89
+ "inputs": { "input": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
90
+ "outputs": { "output": { "dtype": "int32", "shape": [0] } }
91
+ },
92
+ {
93
+ "name": "ort_float32_to_int32_rank3",
94
+ "provenance": {
95
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
96
+ "test": "BitCastTest.Float32ToInt32_3D"
97
+ },
98
+ "attrs": { "to": 6 },
99
+ "inputs": {
100
+ "input": {
101
+ "dtype": "float32",
102
+ "shape": [2, 2, 3],
103
+ "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0] }
104
+ }
105
+ },
106
+ "outputs": { "output": { "dtype": "int32", "shape": [2, 2, 3] } }
107
+ },
108
+ {
109
+ "name": "ort_float32_to_float32_identity",
110
+ "provenance": {
111
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
112
+ "test": "BitCastTest.Float32ToFloat32"
113
+ },
114
+ "attrs": { "to": 1 },
115
+ "inputs": {
116
+ "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
117
+ },
118
+ "outputs": { "output": { "dtype": "float32", "shape": [3] } }
119
+ },
120
+ {
121
+ "name": "onnx_backend_uint32_to_int32_high_bits",
122
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitcast_uint32_to_int32" },
123
+ "attrs": { "to": 6 },
124
+ "inputs": {
125
+ "input": {
126
+ "dtype": "uint32",
127
+ "shape": [3],
128
+ "data": { "kind": "values", "values": [4294967295, 2147483648, 2147483647] }
129
+ }
130
+ },
131
+ "outputs": { "output": { "dtype": "int32", "shape": [3], "tolerance": 0 } }
132
+ },
133
+ {
134
+ "name": "onnx_backend_float32_to_int32",
135
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitcast_float32_to_int32" },
136
+ "attrs": { "to": 6 },
137
+ "inputs": {
138
+ "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, -2.5, 3.75] } }
139
+ },
140
+ "outputs": { "output": { "dtype": "int32", "shape": [3], "tolerance": 0 } }
141
+ },
142
+ {
143
+ "name": "onnx_backend_float32_to_int32_rank2",
144
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitcast_2d_float32_to_int32" },
145
+ "attrs": { "to": 6 },
146
+ "inputs": {
147
+ "input": {
148
+ "dtype": "float32",
149
+ "shape": [2, 3],
150
+ "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
151
+ }
152
+ },
153
+ "outputs": { "output": { "dtype": "int32", "shape": [2, 3], "tolerance": 0 } }
154
+ },
155
+ {
156
+ "name": "onnx_backend_scalar_float32_to_int32",
157
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitcast_scalar_float32_to_int32" },
158
+ "attrs": { "to": 6 },
159
+ "inputs": { "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } } },
160
+ "outputs": { "output": { "dtype": "int32", "shape": [], "tolerance": 0 } }
161
+ },
162
+ {
163
+ "name": "onnx_backend_int32_to_float32",
164
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_bitcast_int32_to_float32" },
165
+ "attrs": { "to": 1 },
166
+ "inputs": {
167
+ "input": {
168
+ "dtype": "int32",
169
+ "shape": [3],
170
+ "data": { "kind": "values", "values": [1065353216, -1071644672, 1081081856] }
171
+ }
172
+ },
173
+ "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
174
+ },
175
+ {
176
+ "name": "float32_to_uint32_special_bits",
177
+ "provenance": {
178
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
179
+ "test": "BitCastTest.Float32ToInt32",
180
+ "notes": "Exercises the same byte-reinterpretation semantics on signed zero and infinities, with uint32 output so high bits are explicit."
181
+ },
182
+ "attrs": { "to": 12 },
183
+ "inputs": {
184
+ "input": {
185
+ "dtype": "float32",
186
+ "shape": [6],
187
+ "data": { "kind": "values", "values": [0.0, -0.0, 1.0, -2.5, "Infinity", "-Infinity"] }
188
+ }
189
+ },
190
+ "outputs": {
191
+ "output": {
192
+ "dtype": "uint32",
193
+ "shape": [6],
194
+ "tolerance": 0,
195
+ "data": { "kind": "values", "values": [0, 2147483648, 1065353216, 3223322624, 2139095040, 4286578688] }
196
+ }
197
+ }
198
+ },
199
+ {
200
+ "name": "uint32_to_float32_special_bits",
201
+ "provenance": {
202
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
203
+ "test": "BitCastTest.UInt32ToFloat32",
204
+ "notes": "Pins down non-finite and signed-zero payloads so BitCast cannot silently become numeric conversion."
205
+ },
206
+ "attrs": { "to": 1 },
207
+ "inputs": {
208
+ "input": {
209
+ "dtype": "uint32",
210
+ "shape": [7],
211
+ "data": {
212
+ "kind": "values",
213
+ "values": [0, 2147483648, 2139095040, 4286578688, 1065353216, 3223322624, 2143289344]
214
+ }
215
+ }
216
+ },
217
+ "outputs": {
218
+ "output": {
219
+ "dtype": "float32",
220
+ "shape": [7],
221
+ "tolerance": 0,
222
+ "allowNaN": true,
223
+ "data": { "kind": "values", "values": [0.0, 0.0, "Infinity", "-Infinity", 1.0, -2.5, "NaN"] }
224
+ }
225
+ }
226
+ },
227
+ {
228
+ "name": "int8_identity_pairwise_route",
229
+ "provenance": {
230
+ "notes": "Identity reinterpretation bridges the independently selectable signed and unsigned 8-bit input/output routes; the ORT-derived int8-to-uint8 sibling retains cross-signedness coverage."
231
+ },
232
+ "attrs": { "to": 3 },
233
+ "inputs": {
234
+ "input": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-1, -128, 127, 0] } }
235
+ },
236
+ "outputs": {
237
+ "output": {
238
+ "dtype": "int8",
239
+ "shape": [4],
240
+ "tolerance": 0,
241
+ "data": { "kind": "values", "values": [-1, -128, 127, 0] }
242
+ }
243
+ }
244
+ },
245
+ {
246
+ "name": "int8_to_uint8_vectorized",
247
+ "provenance": {
248
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
249
+ "test": "BitCastTest.Int8ToUInt8",
250
+ "notes": "A four-element sibling pins the vectorized signed-to-unsigned byte reinterpretation as well as the scalar ORT fixture below."
251
+ },
252
+ "attrs": { "to": 2 },
253
+ "inputs": {
254
+ "input": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [0, -1, 127, -128] } }
255
+ },
256
+ "outputs": {
257
+ "output": {
258
+ "dtype": "uint8",
259
+ "shape": [4],
260
+ "tolerance": 0,
261
+ "data": { "kind": "values", "values": [0, 255, 127, 128] }
262
+ }
263
+ }
264
+ },
265
+ {
266
+ "name": "ort_int8_to_uint8",
267
+ "provenance": {
268
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
269
+ "test": "BitCastTest.Int8ToUInt8"
270
+ },
271
+ "attrs": { "to": 2 },
272
+ "inputs": {
273
+ "input": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [0, 1, -1, 127, -128] } }
274
+ },
275
+ "outputs": {
276
+ "output": {
277
+ "dtype": "uint8",
278
+ "shape": [5],
279
+ "tolerance": 0,
280
+ "data": { "kind": "values", "values": [0, 1, 255, 127, 128] }
281
+ }
282
+ }
283
+ },
284
+ {
285
+ "name": "ort_uint8_to_int8",
286
+ "provenance": {
287
+ "source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
288
+ "test": "BitCastTest.UInt8ToInt8"
289
+ },
290
+ "attrs": { "to": 3 },
291
+ "inputs": {
292
+ "input": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [0, 1, 127, 128, 255] } }
293
+ },
294
+ "outputs": {
295
+ "output": {
296
+ "dtype": "int8",
297
+ "shape": [5],
298
+ "tolerance": 0,
299
+ "data": { "kind": "values", "values": [0, 1, 127, -128, -1] }
300
+ }
301
+ }
302
+ },
303
+ {
304
+ "name": "float32_to_int32_rank6",
305
+ "attrs": { "to": 6 },
306
+ "inputs": {
307
+ "input": {
308
+ "dtype": "float32",
309
+ "shape": [2, 2, 2, 2, 2, 2],
310
+ "data": { "kind": "linspace", "start": -1.0, "end": 1.0 }
311
+ }
312
+ },
313
+ "outputs": { "output": { "dtype": "int32", "shape": [2, 2, 2, 2, 2, 2], "tolerance": 0 } }
314
+ },
315
+ {
316
+ "name": "float32_to_int32_rank7",
317
+ "provenance": {
318
+ "source": "ONNX BitCast-26 contract and ONNX Runtime CPUExecutionProvider",
319
+ "notes": "BitCast preserves arbitrary tensor rank and only reinterprets same-width element bits."
320
+ },
321
+ "attrs": { "to": 6 },
322
+ "inputs": {
323
+ "input": {
324
+ "dtype": "float32",
325
+ "shape": [1, 1, 1, 1, 1, 1, 4],
326
+ "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.5] }
327
+ }
328
+ },
329
+ "outputs": {
330
+ "output": {
331
+ "dtype": "int32",
332
+ "shape": [1, 1, 1, 1, 1, 1, 4],
333
+ "data": { "kind": "values", "values": [0, 1065353216, -1082130432, 1056964608] },
334
+ "tolerance": 0
335
+ }
336
+ }
337
+ },
338
+ {
339
+ "name": "int32_to_float32_numel_not_mult4",
340
+ "attrs": { "to": 1 },
341
+ "inputs": {
342
+ "input": {
343
+ "dtype": "int32",
344
+ "shape": [5],
345
+ "data": { "kind": "values", "values": [0, 1065353216, 1073741824, -1082130432, 1086324736] }
346
+ }
347
+ },
348
+ "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0 } }
349
+ },
350
+ {
351
+ "name": "float32_to_int32_rank8",
352
+ "attrs": { "to": 6 },
353
+ "inputs": {
354
+ "input": {
355
+ "dtype": "float32",
356
+ "shape": [2, 1, 1, 1, 1, 1, 2, 4],
357
+ "data": { "kind": "linspace", "start": -4.0, "end": 11.0 }
358
+ }
359
+ },
360
+ "outputs": { "output": { "dtype": "int32", "shape": [2, 1, 1, 1, 1, 1, 2, 4], "tolerance": 0 } }
361
+ },
362
+ {
363
+ "name": "uint8_to_int8_vec4_packed",
364
+ "provenance": {
365
+ "notes": "Eight uint8 elements exercise the packed uint8-to-int8 reinterpretation. Values straddle 128 in every vec4 lane, making sign extension visible lane by lane."
366
+ },
367
+ "attrs": { "to": 3 },
368
+ "inputs": {
369
+ "input": {
370
+ "dtype": "uint8",
371
+ "shape": [8],
372
+ "data": { "kind": "values", "values": [0, 1, 127, 128, 129, 200, 254, 255] }
373
+ }
374
+ },
375
+ "outputs": {
376
+ "output": {
377
+ "dtype": "int8",
378
+ "shape": [8],
379
+ "tolerance": 0,
380
+ "data": { "kind": "values", "values": [0, 1, 127, -128, -127, -56, -2, -1] }
381
+ }
382
+ }
383
+ }
384
+ ]
385
+ }