sync 91d990483a17

Browse files

Files changed (6) hide show

README.md +16 -13
build/webgpu/bench.json +0 -1
build/webgpu/if-select.wgsl.jinja +15 -15
build/webgpu/manifest.json +13 -54
build/webgpu/metadata.json +7 -7
build/webgpu/test.json +8 -9

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [standard ONNX `If` spec](https://onnx.ai/onnx/operators/onnx__If.html)
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `cond` | `cond` | `B` | — | — | Scalar boolean condition that selects which value tensor to output. | required |
-| `then_value` | `then_value` | `T` | — | — | Values to output when `cond` is true. | required |
-| `else_value` | `else_value` | `T` | — | — | Values to output when `cond` is false. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `y` | `y` | `T` | — | — | Output tensor with the same number of elements as `then_value` and `else_value`. | required |
 ## Type constraints
@@ -39,7 +39,7 @@ See the [standard ONNX `If` spec](https://onnx.ai/onnx/operators/onnx__If.html)
 ## Files
-- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, 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
@@ -47,15 +47,18 @@ See the [standard ONNX `If` spec](https://onnx.ai/onnx/operators/onnx__If.html)
 ## Use with `@huggingface/kernels`
-The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
-The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
-- `y`
-Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Inputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `cond` | `B` | — | — | Scalar boolean condition that selects which value tensor to output. | required |
+| `then_value` | `T` | — | — | Values to output when `cond` is true. | required |
+| `else_value` | `T` | — | — | Values to output when `cond` is false. | required |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `y` | `T` | — | — | Output tensor with the same number of elements as `then_value` and `else_value`. | required |
 ## Type constraints
 ## Files
+- [`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
 ## 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.
+This example supplies explicit metadata for:
+- `y`
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.If",
   "cases": [
     {
       "name": "lowered_select_1m",

 {
   "cases": [
     {
       "name": "lowered_select_1m",

build/webgpu/if-select.wgsl.jinja CHANGED Viewed

@@ -1,30 +1,30 @@
 {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
 {% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
 {% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at the device's per-axis workgroup
-  // limit; gid.y carries the high portion of the output index.
 {% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
 {% elif note == "element-limit" %}
   // 2D-folded flat element index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "dispatch" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit.
 {% endif %}
 {% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
@@ -34,7 +34,7 @@
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
 {{ flat_index_2d(guardInline=true) }}
   y[i] = select(else_value[i], then_value[i], cond[0] != 0u);
 }

 {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
 {% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; gid.y carries the high portion of the output index.
 {% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (the dispatch caps x and spills into y).
 {% elif note == "element-limit" %}
   // 2D-folded flat element index: gid.y carries the high bits past the
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d(guardInline=true) }}
   y[i] = select(else_value[i], then_value[i], cond[0] != 0u);
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -3,31 +3,10 @@
   "name": "If",
   "conformance": "internal-lowering",
   "sinceVersion": 25,
-  "description": "Support status: the standard ONNX `If` control-flow operator is not implemented because standalone kernel packages cannot carry or execute its `then_branch` and `else_branch` graph attributes. This internal lowering only selects elementwise between two pre-evaluated, equal-sized tensors from a scalar condition and must not be treated as ONNX `If`.",
-  "inputs": [
-    {
-      "role": "cond",
-      "dtype": "B",
-      "description": "Scalar boolean condition that selects which value tensor to output."
-    },
-    { "role": "then_value", "dtype": "T", "description": "Values to output when `cond` is true." },
-    { "role": "else_value", "dtype": "T", "description": "Values to output when `cond` is false." }
-  ],
-  "outputs": [
-    {
-      "role": "y",
-      "dtype": "T",
-      "description": "Output tensor with the same number of elements as `then_value` and `else_value`."
-    }
-  ],
   "typeConstraints": { "T": ["float32"], "B": ["uint32", "bool"] },
-  "args": {
-    "cond": { "kind": "tensor", "semantic": "cond", "role": "input" },
-    "then_value": { "kind": "tensor", "semantic": "then_value", "role": "input" },
-    "else_value": { "kind": "tensor", "semantic": "else_value", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "variants": [
     {
       "id": "lowered_select",
@@ -38,37 +17,17 @@
           "name": "If",
           "shader": "if-select.wgsl.jinja",
           "bindings": [
-            {
-              "name": "cond",
-              "arg": "cond",
-              "semantic": "cond",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "u32",
-              "length": 1
-            },
-            {
-              "name": "then_value",
-              "arg": "then_value",
-              "semantic": "then_value",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            {
-              "name": "else_value",
-              "arg": "else_value",
-              "semantic": "else_value",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "f32" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
-            }
           ],
-          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "name": "If",
   "conformance": "internal-lowering",
   "sinceVersion": 25,
+  "inputs": { "cond": { "dtype": "B" }, "then_value": { "dtype": "T" }, "else_value": { "dtype": "T" } },
+  "outputs": { "y": { "dtype": "T" } },
   "typeConstraints": { "T": ["float32"], "B": ["uint32", "bool"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "variants": [
     {
       "id": "lowered_select",
           "name": "If",
           "shader": "if-select.wgsl.jinja",
           "bindings": [
+            { "arg": "cond", "elementType": "u32", "length": 1 },
+            { "arg": "then_value", "elementType": "f32" },
+            { "arg": "else_value", "elementType": "f32" },
+            { "arg": "y", "elementType": "f32" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
           ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,18 +1,18 @@
 {
   "name": "ai.onnx.If",
-  "id": "_ai_onnx_if_webgpu_2b93a6b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "aVzy6O2BXjbtp7QD6jSgMGQnCFr9rv3ErzC9HuwA6gE=",
-      "if-select.wgsl.jinja": "SOCO3HJ0GVfXYEOn7wNTEDBq94rDoQ5sGuijBo/4u6Y=",
-      "manifest.json": "qnXbI/F6Tz6YY42S8aJWiA5scMh+fR/ZkPTGfwzYYcc=",
-      "test.json": "wPMNj7icFakVJ6lHXUZxozdTSnpyTakK1is0VuwCvzE="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.If" }
 }

 {
   "name": "ai.onnx.If",
+  "id": "_ai_onnx_if_webgpu_fc82927",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "pKYDQ5YSgA8pWav7C7GCSLohOpQ9qIZlwV+NNK8NSic=",
+      "if-select.wgsl.jinja": "ywcv63EICuDTOALNqhfHgflu7LBg2gMSAG9rVDWXZNE=",
+      "manifest.json": "O5az6AZ2qds85ksKnrtYq+QUiS3/b7AbAZG0hzJ83q4=",
+      "test.json": "S8aT5V28HOrUh8vxW+Is0BVgTE78h+8AUypNHxDJdUA="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": { "manifestSpec": "2.0", "variants": { "lowered_select": ["if-select.wgsl.jinja"] } }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.If",
   "cases": [
     {
       "name": "lowered_dispatch_cliff_select_then",
@@ -41,7 +40,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ShapeInMainGraph_NoShapeInSubgraph_True",
-        "notes": "Projection onto the framework's lowered select variant: ORT's then branch computes split_out_0 + if_graph_input_0 = 2."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
@@ -55,7 +54,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ShapeInMainGraph_NoShapeInSubgraph_False",
-        "notes": "Projection onto the framework's lowered select variant: ORT's else branch computes split_out_1 + if_graph_input_0 = 11."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
@@ -69,7 +68,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ThenBranchExecution",
-        "notes": "Projection onto the framework's lowered select variant: branch subgraphs are represented by precomputed branch tensors."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
@@ -83,7 +82,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ElseBranchExecution",
-        "notes": "Projection onto the framework's lowered select variant: branch subgraphs are represented by precomputed branch tensors."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
@@ -97,7 +96,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.Opset11ThenAndElseBranchesProduceDifferentOutputShapes",
-        "notes": "Projection onto the framework's lowered select variant: branch tensors use different ranks with equal storage size, and the selected then branch fixes the output shape."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
@@ -119,7 +118,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.Opset11ThenAndElseBranchesProduceDifferentOutputShapes",
-        "notes": "Projection onto the framework's lowered select variant: branch tensors use different ranks with equal storage size, and the selected else branch fixes the output shape."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
@@ -150,7 +149,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ThenBranchExecution",
-        "notes": "ONNX If condition is a scalar bool. The project-lowered select form should accept scalar logical uint32 conditions as well as [1] conditions."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [1] } },
@@ -196,7 +195,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ElseBranchExecution",
-        "notes": "Length-1 logical bool condition companion for the lowered select form; false must select the else tensor exactly."
       },
       "inputs": {
         "cond": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [0] } },

 {
   "cases": [
     {
       "name": "lowered_dispatch_cliff_select_then",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ShapeInMainGraph_NoShapeInSubgraph_True",
+        "notes": "The internal lowered-select variant receives precomputed branch values; the selected then value is 2."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ShapeInMainGraph_NoShapeInSubgraph_False",
+        "notes": "The internal lowered-select variant receives precomputed branch values; the selected else value is 11."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ThenBranchExecution",
+        "notes": "The internal lowered-select variant receives precomputed branch tensors and selects the then value."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ElseBranchExecution",
+        "notes": "The internal lowered-select variant receives precomputed branch tensors and selects the else value."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.Opset11ThenAndElseBranchesProduceDifferentOutputShapes",
+        "notes": "Precomputed branch tensors have different ranks but equal storage size; selecting the then tensor determines the output shape."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.Opset11ThenAndElseBranchesProduceDifferentOutputShapes",
+        "notes": "Precomputed branch tensors have different ranks but equal storage size; selecting the else tensor determines the output shape."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ThenBranchExecution",
+        "notes": "The fixed-branch select lowering accepts a scalar logical uint32 condition and chooses the complete two-element then value."
       },
       "inputs": {
         "cond": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [1] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/controlflow/if_test.cc",
         "test": "If.ConditionalBranchesOnlyContainConstantNodes_ElseBranchExecution",
+        "notes": "A length-one logical bool condition set to false must select the else tensor exactly in the lowered-select form."
       },
       "inputs": {
         "cond": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [0] } },