sync 6fdf6301e2bb
Browse files- README.md +2 -2
- build/webgpu/metadata.json +4 -4
- build/webgpu/split-to-sequence.wgsl.jinja +6 -3
README.md
CHANGED
|
@@ -55,13 +55,13 @@ Default values (overridable per request):
|
|
| 55 |
- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
|
| 56 |
- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
|
| 57 |
- [`test.json`](build/webgpu/test.json) — correctness cases
|
| 58 |
-
- [`bench.json`](build/webgpu/bench.json) — benchmark
|
| 59 |
- [`split-to-sequence.wgsl.jinja`](build/webgpu/split-to-sequence.wgsl.jinja)
|
| 60 |
|
| 61 |
## Use with `@huggingface/kernels`
|
| 62 |
|
| 63 |
```sh
|
| 64 |
-
npm install --save-exact @huggingface/kernels@0.0.1-preview.
|
| 65 |
```
|
| 66 |
|
| 67 |
Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
|
|
|
|
| 55 |
- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
|
| 56 |
- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
|
| 57 |
- [`test.json`](build/webgpu/test.json) — correctness cases
|
| 58 |
+
- [`bench.json`](build/webgpu/bench.json) — benchmark cases
|
| 59 |
- [`split-to-sequence.wgsl.jinja`](build/webgpu/split-to-sequence.wgsl.jinja)
|
| 60 |
|
| 61 |
## Use with `@huggingface/kernels`
|
| 62 |
|
| 63 |
```sh
|
| 64 |
+
npm install --save-exact @huggingface/kernels@0.0.1-preview.3
|
| 65 |
```
|
| 66 |
|
| 67 |
Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
|
build/webgpu/metadata.json
CHANGED
|
@@ -1,6 +1,6 @@
|
|
| 1 |
{
|
| 2 |
"name": "ai.onnx.SplitToSequence",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"backend": { "type": "webgpu" },
|
|
@@ -9,13 +9,13 @@
|
|
| 9 |
"files": {
|
| 10 |
"bench.json": "Zs84ZKi7bfjFdmtbBi7218N/ixgvCEbOJMXn6YS5yn4=",
|
| 11 |
"manifest.json": "KLNBWHNVH08iKjw2HfJ33n1lHvbg3fmdzBcp6ioyvgk=",
|
| 12 |
-
"split-to-sequence.wgsl.jinja": "
|
| 13 |
"test.json": "FAy7gFEfxoFsB44fhBjHbN4cWZYYHKlVqz09kYcJMr0="
|
| 14 |
}
|
| 15 |
},
|
| 16 |
-
"provenance": { "kernel": { "sha": "
|
| 17 |
"webgpu": {
|
| 18 |
-
"manifestSpec": "2.
|
| 19 |
"variants": {
|
| 20 |
"two_outputs": ["split-to-sequence.wgsl.jinja"],
|
| 21 |
"three_outputs": ["split-to-sequence.wgsl.jinja"],
|
|
|
|
| 1 |
{
|
| 2 |
"name": "ai.onnx.SplitToSequence",
|
| 3 |
+
"id": "_ai_onnx_splittosequence_webgpu_e9c235d",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"backend": { "type": "webgpu" },
|
|
|
|
| 9 |
"files": {
|
| 10 |
"bench.json": "Zs84ZKi7bfjFdmtbBi7218N/ixgvCEbOJMXn6YS5yn4=",
|
| 11 |
"manifest.json": "KLNBWHNVH08iKjw2HfJ33n1lHvbg3fmdzBcp6ioyvgk=",
|
| 12 |
+
"split-to-sequence.wgsl.jinja": "GMGhxy1517EY/lb9BHhTdlhsu6/jinINkwFH9LLx0rE=",
|
| 13 |
"test.json": "FAy7gFEfxoFsB44fhBjHbN4cWZYYHKlVqz09kYcJMr0="
|
| 14 |
}
|
| 15 |
},
|
| 16 |
+
"provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
|
| 17 |
"webgpu": {
|
| 18 |
+
"manifestSpec": "2.1",
|
| 19 |
"variants": {
|
| 20 |
"two_outputs": ["split-to-sequence.wgsl.jinja"],
|
| 21 |
"three_outputs": ["split-to-sequence.wgsl.jinja"],
|
build/webgpu/split-to-sequence.wgsl.jinja
CHANGED
|
@@ -1,3 +1,8 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
{{ env.wgsl.resourceDeclarations }}
|
| 2 |
|
| 3 |
// This internal lowering moves contiguous slices along the split axis into
|
|
@@ -48,9 +53,7 @@ fn input_offset_y{{ k }}(out_index: u32) -> u32 {
|
|
| 48 |
|
| 49 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 50 |
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 51 |
-
|
| 52 |
-
// per-axis dispatch fold width (outputs > 16.7M elements).
|
| 53 |
-
let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
|
| 54 |
{% for k in range(numOutputs) %}
|
| 55 |
if (i < params.y{{ k }}Count) {
|
| 56 |
y{{ k }}[i] = input[input_offset_y{{ k }}(i)];
|
|
|
|
| 1 |
+
{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
|
| 2 |
+
{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
|
| 3 |
+
// 2D-folded flat index: gid.y carries the high bits past the dispatch's
|
| 4 |
+
// per-axis workgroup fold width.
|
| 5 |
+
let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% endmacro %}
|
| 6 |
{{ env.wgsl.resourceDeclarations }}
|
| 7 |
|
| 8 |
// This internal lowering moves contiguous slices along the split axis into
|
|
|
|
| 53 |
|
| 54 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 55 |
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 56 |
+
{{ flat_index_2d(tunables.WORKGROUP_SIZE, "i", "") }}
|
|
|
|
|
|
|
| 57 |
{% for k in range(numOutputs) %}
|
| 58 |
if (i < params.y{{ k }}Count) {
|
| 59 |
y{{ k }}[i] = input[input_offset_y{{ k }}(i)];
|