sync 6fdf6301e2bb

Browse files

Files changed (5) hide show

README.md +2 -2
build/webgpu/bench.json +5 -5
build/webgpu/loop-add-step.wgsl.jinja +1 -1
build/webgpu/manifest.json +1 -1
build/webgpu/metadata.json +6 -6

README.md CHANGED Viewed

@@ -45,13 +45,13 @@ See the [standard ONNX `Loop` spec](https://onnx.ai/onnx/operators/onnx__Loop.ht
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
-- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`loop-add-step.wgsl.jinja`](build/webgpu/loop-add-step.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
-npm install --save-exact @huggingface/kernels@0.0.1-preview.2
 ```
 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.

 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
 - [`loop-add-step.wgsl.jinja`](build/webgpu/loop-add-step.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 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/bench.json CHANGED Viewed

@@ -1,7 +1,7 @@
 {
   "cases": [
     {
-      "name": "lowered_HEALTHY_baseline_dim128_steps128",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [128] } },
@@ -16,7 +16,7 @@
       "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
-      "name": "lowered_PATHO_singlelane_serial_dim2048_steps512",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [512] } },
@@ -31,7 +31,7 @@
       "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
-      "name": "lowered_PATHO_singlelane_serial_dim4096_steps1024",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1024] } },
@@ -46,7 +46,7 @@
       "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
-      "name": "lowered_PATHO_smalldim2_singlelane_steps8192",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8192] } },
@@ -61,7 +61,7 @@
       "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
-      "name": "lowered_PATHO_nonpow2_unaligned_dim1000_steps1000",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1000] } },

 {
   "cases": [
     {
+      "name": "lowered_control_baseline_dim128_steps128",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [128] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
+      "name": "lowered_stress_singlelane_serial_dim2048_steps512",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [512] } },
       "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
+      "name": "lowered_stress_singlelane_serial_dim4096_steps1024",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1024] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
+      "name": "lowered_stress_smalldim2_singlelane_steps8192",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8192] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
     },
     {
+      "name": "lowered_stress_nonpow2_unaligned_dim1000_steps1000",
       "preset": "smoke",
       "inputs": {
         "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1000] } },

build/webgpu/loop-add-step.wgsl.jinja CHANGED Viewed

@@ -3,7 +3,7 @@
 // One invocation owns each loop-carried element and walks its iterations
 // serially. Columns run independently, but the per-column additions and stores
 // stay sequential, preserving the f32 accumulation order.
-@compute @workgroup_size(256)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let i = gid.x;
   if (i >= params.dim) { return; }

 // One invocation owns each loop-carried element and walks its iterations
 // serially. Columns run independently, but the per-column additions and stores
 // stay sequential, preserving the f32 accumulation order.
+@compute @workgroup_size({{ workgroupSize }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let i = gid.x;
   if (i >= params.dim) { return; }

build/webgpu/manifest.json CHANGED Viewed

@@ -18,7 +18,7 @@
     {
       "id": "lowered_add_step",
       "when": ["ranks.m <= 1", "numel(shapes.m) == 1", "ranks.cond <= 1", "numel(shapes.cond) == 1", "ranks.v_initial == 1", "ranks.step == 1", "ranks.v_final == 1", "ranks.scan_output == 2", "dim(shapes.step, 0) == dim(shapes.v_initial, 0)", "dim(shapes.v_final, 0) == dim(shapes.v_initial, 0)", "dim(shapes.scan_output, 1) == dim(shapes.v_initial, 0)"],
-      "derive": { "dimBlocks": "ceil(dim(shapes.v_initial, 0) / 256)" },
       "passes": [
         {
           "id": "main",

     {
       "id": "lowered_add_step",
       "when": ["ranks.m <= 1", "numel(shapes.m) == 1", "ranks.cond <= 1", "numel(shapes.cond) == 1", "ranks.v_initial == 1", "ranks.step == 1", "ranks.v_final == 1", "ranks.scan_output == 2", "dim(shapes.step, 0) == dim(shapes.v_initial, 0)", "dim(shapes.v_final, 0) == dim(shapes.v_initial, 0)", "dim(shapes.scan_output, 1) == dim(shapes.v_initial, 0)"],
+      "derive": { "workgroupSize": 256, "dimBlocks": "ceil(dim(shapes.v_initial, 0) / workgroupSize)" },
       "passes": [
         {
           "id": "main",

build/webgpu/metadata.json CHANGED Viewed

@@ -1,18 +1,18 @@
 {
   "name": "ai.onnx.Loop",
-  "id": "_ai_onnx_loop_webgpu_801da86",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "ZSbqJEih0tA29HPs9gNfjrVWJeylslE0PXPC0uKM82s=",
-      "loop-add-step.wgsl.jinja": "Sk4601cYmNLKCMIFX47m3mKo2M5oh/fekvOwLOUN1D8=",
-      "manifest.json": "H3mdjtpFSSHh+Osunjn3M82kXBVm5ga/DK+8Ha0Wwoc=",
       "test.json": "Rp36C1l3Leiipq280ww8mCj4exFK53UlSD0thVu/2P0="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
-  "webgpu": { "manifestSpec": "2.0", "variants": { "lowered_add_step": ["loop-add-step.wgsl.jinja"] } }
 }

 {
   "name": "ai.onnx.Loop",
+  "id": "_ai_onnx_loop_webgpu_6b860c0",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "gjMN5zadXNCeu4axKIR/IDNYDMcA9DgBZi6Dk0UVGLc=",
+      "loop-add-step.wgsl.jinja": "53Jj1FiRGGzYrUleplg+3QHim5zdxOQY2oZB2dESxr0=",
+      "manifest.json": "popStqex+yDicCRjJQwtzATv+7fwuipB98BQBGJncGg=",
       "test.json": "Rp36C1l3Leiipq280ww8mCj4exFK53UlSD0thVu/2P0="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
+  "webgpu": { "manifestSpec": "2.1", "variants": { "lowered_add_step": ["loop-add-step.wgsl.jinja"] } }
 }