sync 91d990483a17
Browse files- README.md +13 -9
- build/webgpu/bench.json +0 -1
- build/webgpu/bitcast.wgsl.jinja +6 -6
- build/webgpu/manifest.json +33 -85
- build/webgpu/metadata.json +10 -7
- build/webgpu/test.json +2 -3
README.md
CHANGED
|
@@ -18,15 +18,15 @@ See the [ONNX `BitCast` spec](https://onnx.ai/onnx/operators/onnx__BitCast.html)
|
|
| 18 |
|
| 19 |
## Inputs
|
| 20 |
|
| 21 |
-
| Name |
|
| 22 |
-
| --- | --- | --- | --- | --- | --- |
|
| 23 |
-
| `input` | `
|
| 24 |
|
| 25 |
## Outputs
|
| 26 |
|
| 27 |
-
| Name |
|
| 28 |
-
| --- | --- | --- | --- | --- | --- |
|
| 29 |
-
| `output` | `
|
| 30 |
|
| 31 |
## Attributes
|
| 32 |
|
|
@@ -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
|
|
@@ -53,10 +53,14 @@ Attributes and default values (overridable per request):
|
|
| 53 |
|
| 54 |
## Use with `@huggingface/kernels`
|
| 55 |
|
| 56 |
-
|
| 57 |
-
|
|
|
|
|
|
|
|
|
|
| 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 |
|
|
|
|
| 18 |
|
| 19 |
## Inputs
|
| 20 |
|
| 21 |
+
| Name | Logical dtype | Rank | Shape | Description | Presence |
|
| 22 |
+
| --- | --- | --- | --- | --- | --- |
|
| 23 |
+
| `input` | `T` | — | — | Input tensor to be bitwise reinterpreted. | required |
|
| 24 |
|
| 25 |
## Outputs
|
| 26 |
|
| 27 |
+
| Name | Logical dtype | Rank | Shape | Description | Presence |
|
| 28 |
+
| --- | --- | --- | --- | --- | --- |
|
| 29 |
+
| `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 |
|
|
|
|
| 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
|
|
|
|
| 53 |
|
| 54 |
## Use with `@huggingface/kernels`
|
| 55 |
|
| 56 |
+
```sh
|
| 57 |
+
npm install --save-exact @huggingface/kernels@0.0.1-preview.2
|
| 58 |
+
```
|
| 59 |
+
|
| 60 |
+
Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
|
| 61 |
|
| 62 |
The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
|
| 63 |
+
It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
|
| 64 |
|
| 65 |
Replace each `*Data` placeholder with a typed array containing the corresponding input data.
|
| 66 |
|
build/webgpu/bench.json
CHANGED
|
@@ -1,5 +1,4 @@
|
|
| 1 |
{
|
| 2 |
-
"op": "ai.onnx.BitCast",
|
| 3 |
"cases": [
|
| 4 |
{
|
| 5 |
"name": "bitcast-f32-to-u32-33m",
|
|
|
|
| 1 |
{
|
|
|
|
| 2 |
"cases": [
|
| 3 |
{
|
| 4 |
"name": "bitcast-f32-to-u32-33m",
|
build/webgpu/bitcast.wgsl.jinja
CHANGED
|
@@ -1,14 +1,14 @@
|
|
| 1 |
{{ env.wgsl.resourceDeclarations }}
|
| 2 |
|
| 3 |
-
{% set vectorized =
|
| 4 |
// Each invocation reinterprets one 32-bit slot, or four contiguous slots on
|
| 5 |
-
// the vectorized route.
|
| 6 |
-
// explicit low-byte masking and sign extension.
|
| 7 |
-
@compute @workgroup_size({{
|
| 8 |
-
fn main(@builtin(global_invocation_id) gid: 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 *
|
| 12 |
if (i >= params.count) {
|
| 13 |
return;
|
| 14 |
}
|
|
|
|
| 1 |
{{ env.wgsl.resourceDeclarations }}
|
| 2 |
|
| 3 |
+
{% set vectorized = vectorizedSpec if vectorizedSpec is defined else false %}
|
| 4 |
// Each invocation reinterprets one 32-bit slot, or four contiguous slots on
|
| 5 |
+
// the vectorized route. Logical int8/uint8 values stored in 32-bit slots retain
|
| 6 |
+
// their explicit low-byte masking and sign extension.
|
| 7 |
+
@compute @workgroup_size({{ workgroupSizeSpec }})
|
| 8 |
+
fn main(@builtin(global_invocation_id) gid: 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 * {{ DISPATCH_FOLD_WIDTH }}u * {{ workgroupSizeSpec }}u;
|
| 12 |
if (i >= params.count) {
|
| 13 |
return;
|
| 14 |
}
|
build/webgpu/manifest.json
CHANGED
|
@@ -2,40 +2,22 @@
|
|
| 2 |
"domain": "ai.onnx",
|
| 3 |
"name": "BitCast",
|
| 4 |
"sinceVersion": 26,
|
| 5 |
-
"
|
| 6 |
-
"
|
| 7 |
-
"
|
| 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 |
-
"
|
| 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\"",
|
|
@@ -43,84 +25,50 @@
|
|
| 43 |
"outputIsInt8": "tensorDtypes.output == \"int8\"",
|
| 44 |
"outputIsUint8": "tensorDtypes.output == \"uint8\""
|
| 45 |
},
|
| 46 |
-
"
|
| 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"
|
| 100 |
-
"
|
| 101 |
"passes": [
|
| 102 |
{
|
| 103 |
"id": "main",
|
| 104 |
"name": "BitCast.vec4",
|
| 105 |
-
"
|
| 106 |
-
|
| 107 |
-
|
| 108 |
-
|
| 109 |
-
|
| 110 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 111 |
}
|
| 112 |
]
|
| 113 |
},
|
| 114 |
{
|
| 115 |
"id": "slot32",
|
| 116 |
-
"when": ["bitcastBaseOk"],
|
| 117 |
"passes": [
|
| 118 |
{
|
| 119 |
"id": "main",
|
| 120 |
"name": "BitCast",
|
| 121 |
-
"
|
| 122 |
-
"
|
| 123 |
-
"
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 124 |
}
|
| 125 |
]
|
| 126 |
}
|
|
|
|
| 2 |
"domain": "ai.onnx",
|
| 3 |
"name": "BitCast",
|
| 4 |
"sinceVersion": 26,
|
| 5 |
+
"inputs": { "input": { "dtype": "T" } },
|
| 6 |
+
"outputs": { "output": { "dtype": "U", "rank": "ranks.input", "shape": "shapes.input" } },
|
| 7 |
+
"attributes": { "to": {} },
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 8 |
"attributeConstraints": { "to": { "required": true } },
|
| 9 |
"typeConstraints": {
|
| 10 |
"T": ["float32", "int8", "int32", "uint8", "uint32"],
|
| 11 |
"U": ["float32", "int8", "int32", "uint8", "uint32"]
|
| 12 |
},
|
| 13 |
+
"tunables": { "WORKGROUP_SIZE": { "default": 256 } },
|
|
|
|
|
|
|
|
|
|
|
|
|
| 14 |
"derive": {
|
| 15 |
"deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
|
| 16 |
"bitcastShapeOk": "ranks.input >= 0 and ranks.output == ranks.input and numel(shapes.input) == numel(shapes.output)",
|
| 17 |
"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\"))",
|
| 18 |
"bitcastWorkgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
|
| 19 |
+
"bitcastDispatchFits": "ceilDiv(ceilDiv(numel(shapes.output), bitcastWorkgroupSize), min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
|
| 20 |
+
"bitcastBaseOk": "bitcastShapeOk and bitcastTypeOk and bitcastDispatchFits",
|
|
|
|
|
|
|
| 21 |
"inScalar": "dtypes.T",
|
| 22 |
"outScalar": "dtypes.U",
|
| 23 |
"inputIsInt8": "tensorDtypes.input == \"int8\"",
|
|
|
|
| 25 |
"outputIsInt8": "tensorDtypes.output == \"int8\"",
|
| 26 |
"outputIsUint8": "tensorDtypes.output == \"uint8\""
|
| 27 |
},
|
| 28 |
+
"when": ["bitcastBaseOk"],
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 29 |
"variants": [
|
| 30 |
{
|
| 31 |
"id": "slot32_vec4",
|
| 32 |
"priority": 20,
|
| 33 |
+
"when": ["numel(shapes.input) % 4 == 0"],
|
| 34 |
+
"derive": { "inVec4": "\"vec4<\" ~ dtypes.T ~ \">\"", "outVec4": "\"vec4<\" ~ dtypes.U ~ \">\"" },
|
| 35 |
"passes": [
|
| 36 |
{
|
| 37 |
"id": "main",
|
| 38 |
"name": "BitCast.vec4",
|
| 39 |
+
"shader": "bitcast.wgsl.jinja",
|
| 40 |
+
"derive": { "vectorizedSpec": true, "workgroupSizeSpec": "bitcastWorkgroupSize" },
|
| 41 |
+
"bindings": [
|
| 42 |
+
{ "arg": "input", "elementType": "$inVec4" },
|
| 43 |
+
{ "arg": "output", "elementType": "$outVec4" },
|
| 44 |
+
{ "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }] }
|
| 45 |
+
],
|
| 46 |
+
"dispatch": {
|
| 47 |
+
"x": "min(ceilDiv((numel(shapes.output) / 4), (bitcastWorkgroupSize)), 65535)",
|
| 48 |
+
"y": "ceilDiv(ceilDiv((numel(shapes.output) / 4), (bitcastWorkgroupSize)), 65535)",
|
| 49 |
+
"z": 1
|
| 50 |
+
}
|
| 51 |
}
|
| 52 |
]
|
| 53 |
},
|
| 54 |
{
|
| 55 |
"id": "slot32",
|
|
|
|
| 56 |
"passes": [
|
| 57 |
{
|
| 58 |
"id": "main",
|
| 59 |
"name": "BitCast",
|
| 60 |
+
"shader": "bitcast.wgsl.jinja",
|
| 61 |
+
"derive": { "workgroupSizeSpec": "bitcastWorkgroupSize" },
|
| 62 |
+
"bindings": [
|
| 63 |
+
"input",
|
| 64 |
+
"output",
|
| 65 |
+
{ "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }] }
|
| 66 |
+
],
|
| 67 |
+
"dispatch": {
|
| 68 |
+
"x": "min(ceilDiv((numel(shapes.output)), (bitcastWorkgroupSize)), 65535)",
|
| 69 |
+
"y": "ceilDiv(ceilDiv((numel(shapes.output)), (bitcastWorkgroupSize)), 65535)",
|
| 70 |
+
"z": 1
|
| 71 |
+
}
|
| 72 |
}
|
| 73 |
]
|
| 74 |
}
|
build/webgpu/metadata.json
CHANGED
|
@@ -1,18 +1,21 @@
|
|
| 1 |
{
|
| 2 |
"name": "ai.onnx.BitCast",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"backend": { "type": "webgpu" },
|
| 7 |
"digest": {
|
| 8 |
"algorithm": "sha256",
|
| 9 |
"files": {
|
| 10 |
-
"bench.json": "
|
| 11 |
-
"bitcast.wgsl.jinja": "
|
| 12 |
-
"manifest.json": "
|
| 13 |
-
"test.json": "
|
| 14 |
}
|
| 15 |
},
|
| 16 |
-
"provenance": { "kernel": { "sha": "
|
| 17 |
-
"webgpu": {
|
|
|
|
|
|
|
|
|
|
| 18 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "ai.onnx.BitCast",
|
| 3 |
+
"id": "_ai_onnx_bitcast_webgpu_ab905c6",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"backend": { "type": "webgpu" },
|
| 7 |
"digest": {
|
| 8 |
"algorithm": "sha256",
|
| 9 |
"files": {
|
| 10 |
+
"bench.json": "7aNcJ8LB0eyjvFh7jXkwAw1p2vgZejuWw6H9iwqbgqY=",
|
| 11 |
+
"bitcast.wgsl.jinja": "1vfOUaYUxSx36bWOnKDPyXUPGprDxhPVEv10804xfVI=",
|
| 12 |
+
"manifest.json": "UplDa7horbJakWMkW0Gkocun6e1kcewHPZAvi6weNHs=",
|
| 13 |
+
"test.json": "FXueeEbCTZUCVYOfr0NjJrQAoVhhvtwqADXKZI6829c="
|
| 14 |
}
|
| 15 |
},
|
| 16 |
+
"provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
|
| 17 |
+
"webgpu": {
|
| 18 |
+
"manifestSpec": "2.0",
|
| 19 |
+
"variants": { "slot32_vec4": ["bitcast.wgsl.jinja"], "slot32": ["bitcast.wgsl.jinja"] }
|
| 20 |
+
}
|
| 21 |
}
|
build/webgpu/test.json
CHANGED
|
@@ -1,5 +1,4 @@
|
|
| 1 |
{
|
| 2 |
-
"op": "ai.onnx.BitCast",
|
| 3 |
"cases": [
|
| 4 |
{
|
| 5 |
"name": "float32_to_int32",
|
|
@@ -227,7 +226,7 @@
|
|
| 227 |
{
|
| 228 |
"name": "int8_identity_pairwise_route",
|
| 229 |
"provenance": {
|
| 230 |
-
"notes": "Identity reinterpretation
|
| 231 |
},
|
| 232 |
"attrs": { "to": 3 },
|
| 233 |
"inputs": {
|
|
@@ -247,7 +246,7 @@
|
|
| 247 |
"provenance": {
|
| 248 |
"source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
|
| 249 |
"test": "BitCastTest.Int8ToUInt8",
|
| 250 |
-
"notes": "
|
| 251 |
},
|
| 252 |
"attrs": { "to": 2 },
|
| 253 |
"inputs": {
|
|
|
|
| 1 |
{
|
|
|
|
| 2 |
"cases": [
|
| 3 |
{
|
| 4 |
"name": "float32_to_int32",
|
|
|
|
| 226 |
{
|
| 227 |
"name": "int8_identity_pairwise_route",
|
| 228 |
"provenance": {
|
| 229 |
+
"notes": "Identity reinterpretation exercises signed 8-bit input and output selection independently while preserving every payload bit."
|
| 230 |
},
|
| 231 |
"attrs": { "to": 3 },
|
| 232 |
"inputs": {
|
|
|
|
| 246 |
"provenance": {
|
| 247 |
"source": "onnxruntime/test/providers/cpu/tensor/bitcast_op_test.cc",
|
| 248 |
"test": "BitCastTest.Int8ToUInt8",
|
| 249 |
+
"notes": "Four elements exercise vectorized signed-to-unsigned byte reinterpretation while preserving every payload bit."
|
| 250 |
},
|
| 251 |
"attrs": { "to": 2 },
|
| 252 |
"inputs": {
|