sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +12 -8
build/webgpu/bench.json +0 -1
build/webgpu/compare-broadcast-vec4.wgsl.jinja +20 -24
build/webgpu/compare-broadcast.wgsl.jinja +15 -11
build/webgpu/compare-vec4.wgsl.jinja +35 -6
build/webgpu/manifest.json +136 -169
build/webgpu/metadata.json +19 -9
build/webgpu/test.json +42 -16

README.md CHANGED Viewed

@@ -18,16 +18,16 @@ See the [ONNX `Less` spec](https://onnx.ai/onnx/operators/onnx__Less.html) for t
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `A` | `a` | `T` | — | — | First input operand for the less-than comparison. | required |
-| `B` | `b` | `T` | — | — | Second input operand for the less-than comparison. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `C` | `c` | `B` | derived | broadcast result of `A` and `B` | Boolean result tensor; true where A < B. | required |
 ## Type constraints
@@ -38,7 +38,7 @@ See the [ONNX `Less` spec](https://onnx.ai/onnx/operators/onnx__Less.html) for t
 ## 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
@@ -48,10 +48,14 @@ See the [ONNX `Less` spec](https://onnx.ai/onnx/operators/onnx__Less.html) for t
 ## Use with `@huggingface/kernels`
-The loader derives every required output's shape and logical dtype from the manifest contract and this call.
-It then allocates the result tensors automatically.
 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 | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `a` | `A` | `T` | — | — | First input operand for the less-than comparison. | required |
+| `b` | `B` | `T` | — | — | Second input operand for the less-than comparison. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `c` | `C` | `B` | derived | broadcast result of `a` and `b` | Boolean result tensor; true where A < B. | 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
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 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.Less",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

build/webgpu/compare-broadcast-vec4.wgsl.jinja CHANGED Viewed

@@ -42,33 +42,29 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
 {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
 {%- endmacro %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-// Vec4 broadcast comparison op (Equal/Greater/GreaterOrEqual/Less/LessOrEqual),
-// vectorized over the innermost output axis. Each thread writes
-// a vec4<u32> of {0,1} results. The guard requires the output's innermost
 // (stride-1) axis to be a multiple of 4, so a vec4 group of 4 consecutive
 // outputs never crosses that axis. Per the broadcast rules, for the innermost
 // axis each operand is one of: scalar (whole operand is one element), splat
 // (innermost axis 1), or contiguous (innermost axis matches output). Operands
-// stay scalar-typed bindings (4 coalesced scalar loads == one vec4 of bandwidth)
-// and compare in their native type (no f32 widening — exact for int operands).
 // Grid-strides a clamped dispatch so outputs beyond one dispatch dimension are covered.
 {% set a_numel = namespace(value=1) %}
-{% for d in source.aShape %}{% set a_numel.value = a_numel.value * d %}{% endfor %}
 {% set b_numel = namespace(value=1) %}
-{% for d in source.bShape %}{% set b_numel.value = b_numel.value * d %}{% endfor %}
 {% set c_numel = namespace(value=1) %}
-{% for d in source.cShape %}{% set c_numel.value = c_numel.value * d %}{% endfor %}
-{% set a_same = namespace(value=(source.aRank == source.cRank)) %}
-{% if a_same.value %}{% for axis in range(source.cRank) %}{% if source.aShape[axis] != source.cShape[axis] %}{% set a_same.value = false %}{% endif %}{% endfor %}{% endif %}
-{% set b_same = namespace(value=(source.bRank == source.cRank)) %}
-{% if b_same.value %}{% for axis in range(source.cRank) %}{% if source.bShape[axis] != source.cShape[axis] %}{% set b_same.value = false %}{% endif %}{% endfor %}{% endif %}
-{% set a_inner = source.aShape[source.aRank - 1] if source.aRank >= 1 else 1 %}
-{% set b_inner = source.bShape[source.bRank - 1] if source.bRank >= 1 else 1 %}
 {% if a_numel.value == 1 %}{% set a_mode = "scalar" %}
 {% elif a_inner == 1 %}{% set a_mode = "splat" %}
 {% else %}{% set a_mode = "contig" %}{% endif %}
@@ -77,14 +73,14 @@ enable f16;
 {% else %}{% set b_mode = "contig" %}{% endif %}
 {% if a_mode != "scalar" %}
-{{ offset_fn("a_offset", source.aShape, source.aRank, a_same.value, a_numel.value, source.cShape, source.cRank, c_numel.value) }}
 {% endif %}
 {% if b_mode != "scalar" %}
-{{ offset_fn("b_offset", source.bShape, source.bRank, b_same.value, b_numel.value, source.cShape, source.cRank, c_numel.value) }}
 {% endif %}
-{% set OP = {"equal": "==", "greater": ">", "greaterOrEqual": ">=", "less": "<", "lessOrEqual": "<="}[source.op] %}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
@@ -98,17 +94,17 @@ fn main(
 {% if a_mode == "scalar" %}
     let av = vec4<{{ scalar }}>(a[0]);
 {% elif a_mode == "splat" %}
-    let av = vec4<{{ scalar }}>(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }}]);
 {% else %}
-    let ao = {{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }};
     let av = vec4<{{ scalar }}>(a[ao], a[ao + 1u], a[ao + 2u], a[ao + 3u]);
 {% endif %}
 {% if b_mode == "scalar" %}
     let bv = vec4<{{ scalar }}>(b[0]);
 {% elif b_mode == "splat" %}
-    let bv = vec4<{{ scalar }}>(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }}]);
 {% else %}
-    let bo = {{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }};
     let bv = vec4<{{ scalar }}>(b[bo], b[bo + 1u], b[bo + 2u], b[bo + 3u]);
 {% endif %}
     c[i4] = select(vec4<u32>(0u), vec4<u32>(1u), av {{ OP }} bv);

 {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
 {%- endmacro %}
 {{ env.wgsl.resourceDeclarations }}
+// Broadcast comparison vectorized over the innermost output axis. Each thread
+// writes a vec4<u32> of {0,1} results. This specialization requires the output's innermost
 // (stride-1) axis to be a multiple of 4, so a vec4 group of 4 consecutive
 // outputs never crosses that axis. Per the broadcast rules, for the innermost
 // axis each operand is one of: scalar (whole operand is one element), splat
 // (innermost axis 1), or contiguous (innermost axis matches output). Operands
+// stay scalar-typed bindings; contiguous operands are loaded from four
+// consecutive addresses. Comparisons use the native type with no f32 widening.
 // Grid-strides a clamped dispatch so outputs beyond one dispatch dimension are covered.
 {% set a_numel = namespace(value=1) %}
+{% for d in aShape %}{% set a_numel.value = a_numel.value * d %}{% endfor %}
 {% set b_numel = namespace(value=1) %}
+{% for d in bShape %}{% set b_numel.value = b_numel.value * d %}{% endfor %}
 {% set c_numel = namespace(value=1) %}
+{% for d in cShape %}{% set c_numel.value = c_numel.value * d %}{% endfor %}
+{% set a_same = namespace(value=(aRank == cRank)) %}
+{% if a_same.value %}{% for axis in range(cRank) %}{% if aShape[axis] != cShape[axis] %}{% set a_same.value = false %}{% endif %}{% endfor %}{% endif %}
+{% set b_same = namespace(value=(bRank == cRank)) %}
+{% if b_same.value %}{% for axis in range(cRank) %}{% if bShape[axis] != cShape[axis] %}{% set b_same.value = false %}{% endif %}{% endfor %}{% endif %}
+{% set a_inner = aShape[aRank - 1] if aRank >= 1 else 1 %}
+{% set b_inner = bShape[bRank - 1] if bRank >= 1 else 1 %}
 {% if a_numel.value == 1 %}{% set a_mode = "scalar" %}
 {% elif a_inner == 1 %}{% set a_mode = "splat" %}
 {% else %}{% set a_mode = "contig" %}{% endif %}
 {% else %}{% set b_mode = "contig" %}{% endif %}
 {% if a_mode != "scalar" %}
+{{ offset_fn("a_offset", aShape, aRank, a_same.value, a_numel.value, cShape, cRank, c_numel.value) }}
 {% endif %}
 {% if b_mode != "scalar" %}
+{{ offset_fn("b_offset", bShape, bRank, b_same.value, b_numel.value, cShape, cRank, c_numel.value) }}
 {% endif %}
+{% set OP = {"equal": "==", "greater": ">", "greaterOrEqual": ">=", "less": "<", "lessOrEqual": "<="}[op] %}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 {% if a_mode == "scalar" %}
     let av = vec4<{{ scalar }}>(a[0]);
 {% elif a_mode == "splat" %}
+    let av = vec4<{{ scalar }}>(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base") }}]);
 {% else %}
+    let ao = {{ broadcast_offset_call("a_offset", aShape, cShape, "base") }};
     let av = vec4<{{ scalar }}>(a[ao], a[ao + 1u], a[ao + 2u], a[ao + 3u]);
 {% endif %}
 {% if b_mode == "scalar" %}
     let bv = vec4<{{ scalar }}>(b[0]);
 {% elif b_mode == "splat" %}
+    let bv = vec4<{{ scalar }}>(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base") }}]);
 {% else %}
+    let bo = {{ broadcast_offset_call("b_offset", bShape, cShape, "base") }};
     let bv = vec4<{{ scalar }}>(b[bo], b[bo + 1u], b[bo + 2u], b[bo + 3u]);
 {% endif %}
     c[i4] = select(vec4<u32>(0u), vec4<u32>(1u), av {{ OP }} bv);

build/webgpu/compare-broadcast.wgsl.jinja CHANGED Viewed

@@ -1,13 +1,14 @@
 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
-  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
-  let begin = invocation * {{ source.itemsPerInvocation }}u;
-  let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
@@ -75,20 +76,23 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
 {% endif %}
 {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
 {%- endmacro %}{% macro binary_broadcast_offsets() %}
-{{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.cShape, source.cRank) }}
-{{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.cShape, source.cRank) }}
 {%- endmacro %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {{ binary_broadcast_offsets() }}
 {{ flat_tail_open() }}
-  c[i] = select(0u, 1u, a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] < b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}]);
 {{ flat_tail_close() -}}
 }

 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
+{% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
+  let begin = invocation * {{ itemsPerInvocation }}u;
+  let end = min(begin + {{ itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
 {% endif %}
 {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
 {%- endmacro %}{% macro binary_broadcast_offsets() %}
+{% set aShape = aShape | default([]) %}
+{% set aRank = aRank | default(0) %}
+{% set bShape = bShape | default([]) %}
+{% set bRank = bRank | default(0) %}
+{% set cShape = cShape | default([]) %}
+{% set cRank = cRank | default(0) %}
+{{ broadcast_offset_fn("a_offset", aShape, aRank, cShape, cRank) }}
+{{ broadcast_offset_fn("b_offset", bShape, bRank, cShape, cRank) }}
 {%- endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 {{ binary_broadcast_offsets() }}
 {{ flat_tail_open() }}
+  c[i] = select(0u, 1u, a[{{ broadcast_offset_call("a_offset", aShape, cShape, "i") }}] < b[{{ broadcast_offset_call("b_offset", bShape, cShape, "i") }}]);
 {{ flat_tail_close() -}}
 }

build/webgpu/compare-vec4.wgsl.jinja CHANGED Viewed

@@ -1,17 +1,46 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
-  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
   let av = a[i];
   let bv = b[i];
   c[i] = select(vec4<u32>(0u), vec4<u32>(1u), av < bv);
 }

 {{ env.wgsl.resourceDeclarations }}
+{% set vec4PerThread = vec4PerThread %}
+{% if vec4PerThread > 1 %}
+const ITEMS: u32 = {{ vec4PerThread }}u;
+{% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
+{% if vec4PerThread > 1 %}
+  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
+  // access consecutive words on every step, while each lane can keep several
+  // independent loads in flight.
+  let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  let span = (params.count + ITEMS - 1u) / ITEMS;
+  for (var j = 0u; j < ITEMS; j = j + 1u) {
+    let i = tid + j * span;
+    if (i >= params.count) {
+      break;
+    }
+{% else %}
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
+{% endif %}
+{% set scalarOperand = scalarOperand if scalarOperand is defined else "" %}
+{% if scalarOperand == "a" %}
+  // One-element operand: read once and splat across the vector.
+  let av = {{ vectorScalar }}(a[0]);
+{% else %}
   let av = a[i];
+{% endif %}
+{% if scalarOperand == "b" %}
+  let bv = {{ vectorScalar }}(b[0]);
+{% else %}
   let bv = b[i];
+{% endif %}
   c[i] = select(vec4<u32>(0u), vec4<u32>(1u), av < bv);
+{% if vec4PerThread > 1 %}
+  }
+{% endif %}
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -2,215 +2,182 @@
   "domain": "ai.onnx",
   "name": "Less",
   "sinceVersion": 13,
-  "description": "Performs an elementwise `less-than` comparison between tensors `A` and `B` with NumPy-style multidirectional broadcasting, producing a boolean result tensor `C`. Each output element is `true` where the corresponding element of `A` is strictly less than that of `B`.",
-  "inputs": [
-    { "role": "A", "dtype": "T", "description": "First input operand for the less-than comparison." },
-    { "role": "B", "dtype": "T", "description": "Second input operand for the less-than comparison." }
-  ],
-  "outputs": [
-    {
-      "role": "C",
-      "dtype": "B",
-      "rank": "max(ranks.A, ranks.B)",
-      "description": "Boolean result tensor; true where A < B.",
-      "shape": "broadcastShape(shapes.A, shapes.B)"
-    }
-  ],
-  "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "uint32", "int8", "uint8"], "B": ["bool"] },
-  "args": {
-    "a": { "kind": "tensor", "semantic": "A", "role": "input" },
-    "b": { "kind": "tensor", "semantic": "B", "role": "input" },
-    "c": { "kind": "tensor", "semantic": "C", "role": "output" }
   },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "variants": [
     {
       "id": "same_shape_vec4",
       "priority": 20,
-      "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 == 0", "f16Ok(dtypes.T)"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Less.vec4",
-          "source": { "shader": "compare-vec4.wgsl.jinja", "inputs": { "op": "\"less\"" } },
-          "bindings": [
-            {
-              "name": "a",
-              "arg": "a",
-              "semantic": "A",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "b",
-              "arg": "b",
-              "semantic": "B",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$vectorScalar"
-            },
-            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "vec4<u32>" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
-              }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "broadcast_vec4",
-      "when": ["ranks.A <= ranks.C", "ranks.B <= ranks.C", "ranks.C >= 1", "dim(shapes.C, ranks.C - 1) % 4 == 0", "numel(shapes.C) % 4 == 0", "numel(shapes.C) >= 4", "f16Ok(dtypes.T)"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Less",
-          "source": {
-            "shader": "compare-broadcast-vec4.wgsl.jinja",
-            "inputs": {
-              "aShape": "shapes.A",
-              "bShape": "shapes.B",
-              "cShape": "shapes.C",
-              "aRank": "ranks.A",
-              "bRank": "ranks.B",
-              "cRank": "ranks.C",
-              "op": "\"less\""
-            }
           },
-          "bindings": [
-            {
-              "name": "a",
-              "arg": "a",
-              "semantic": "A",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "b",
-              "arg": "b",
-              "semantic": "B",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "vec4<u32>" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
-              }
-            }
-          ],
-          "dispatch": { "gridStride": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
-      ],
-      "priority": 10
     },
     {
       "id": "broadcast",
-      "when": ["ranks.A <= ranks.C", "ranks.B <= ranks.C", "f16Ok(dtypes.T)"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Less",
-          "source": {
-            "shader": "compare-broadcast.wgsl.jinja",
-            "inputs": {
-              "aShape": "shapes.A",
-              "bShape": "shapes.B",
-              "cShape": "shapes.C",
-              "aRank": "ranks.A",
-              "bRank": "ranks.B",
-              "cRank": "ranks.C",
-              "op": "\"less\"",
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": [
-            {
-              "name": "a",
-              "arg": "a",
-              "semantic": "A",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "b",
-              "arg": "b",
-              "semantic": "B",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "u32" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
-            }
-          ],
-          "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "same_shape_scalar_x4",
       "priority": 15,
-      "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 != 0", "ranks.A <= ranks.C", "ranks.B <= ranks.C", "f16Ok(dtypes.T)"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Less",
-          "source": {
-            "shader": "compare-broadcast.wgsl.jinja",
-            "inputs": {
-              "aShape": "shapes.A",
-              "bShape": "shapes.B",
-              "cShape": "shapes.C",
-              "aRank": "ranks.A",
-              "bRank": "ranks.B",
-              "cRank": "ranks.C",
-              "op": "\"less\"",
-              "itemsPerInvocation": 4
-            }
           },
-          "bindings": [
-            {
-              "name": "a",
-              "arg": "a",
-              "semantic": "A",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "b",
-              "arg": "b",
-              "semantic": "B",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "u32" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
-            }
-          ],
-          "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }
-  ]
 }

   "domain": "ai.onnx",
   "name": "Less",
   "sinceVersion": 13,
+  "inputs": { "a": { "onnx": "A", "dtype": "T" }, "b": { "onnx": "B", "dtype": "T" } },
+  "outputs": {
+    "c": { "onnx": "C", "dtype": "B", "rank": "max(ranks.a, ranks.b)", "shape": "broadcastShape(shapes.a, shapes.b)" }
   },
+  "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "uint32", "int8", "uint8"], "B": ["bool"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "variants": [
     {
       "id": "same_shape_vec4",
       "priority": 20,
+      "when": ["sameShape(shapes.a, shapes.c)", "sameShape(shapes.b, shapes.c)", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Less.vec4",
+          "shader": "compare-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"less\"",
+            "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
+          },
+          "bindings": ["a", "b", "c", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "scalar_b_vec4",
+      "priority": 18,
+      "when": ["sameShape(shapes.a, shapes.c)", "numel(shapes.b) == 1", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"", "scalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Less.scalarBVec4",
+          "shader": "compare-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"less\"",
+            "scalarOperand": "\"b\"",
+            "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
+          },
+          "bindings": ["a", "b_2", "c", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "scalar_a_vec4",
+      "priority": 18,
+      "when": ["sameShape(shapes.b, shapes.c)", "numel(shapes.a) == 1", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"", "scalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Less.scalarAVec4",
+          "shader": "compare-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"less\"",
+            "scalarOperand": "\"a\"",
+            "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
+          },
+          "bindings": ["a_2", "b", "c", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "broadcast_vec4",
+      "priority": 10,
+      "when": ["ranks.a <= ranks.c", "ranks.b <= ranks.c", "ranks.c >= 1", "dim(shapes.c, ranks.c - 1) % 4 == 0", "numel(shapes.c) % 4 == 0", "numel(shapes.c) >= 4", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Less",
+          "shader": "compare-broadcast-vec4.wgsl.jinja",
+          "derive": {
+            "aShape": "shapes.a",
+            "bShape": "shapes.b",
+            "cShape": "shapes.c",
+            "aRank": "ranks.a",
+            "bRank": "ranks.b",
+            "cRank": "ranks.c",
+            "op": "\"less\""
           },
+          "bindings": ["a_2", "b_2", "c", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.c) / 4), (tunables.WORKGROUP_SIZE)), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": 1,
+            "z": 1
+          }
         }
+      ]
     },
     {
       "id": "broadcast",
+      "when": ["ranks.a <= ranks.c", "ranks.b <= ranks.c", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Less",
+          "shader": "compare-broadcast.wgsl.jinja",
+          "derive": {
+            "aShape": "shapes.a",
+            "bShape": "shapes.b",
+            "cShape": "shapes.c",
+            "aRank": "ranks.a",
+            "bRank": "ranks.b",
+            "cRank": "ranks.c",
+            "op": "\"less\"",
+            "itemsPerInvocation": 4
           },
+          "bindings": ["a_2", "b_2", "c_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "same_shape_scalar_x4",
       "priority": 15,
+      "when": ["sameShape(shapes.a, shapes.c)", "sameShape(shapes.b, shapes.c)", "numel(shapes.c) > 0", "numel(shapes.c) % 4 != 0", "ranks.a <= ranks.c", "ranks.b <= ranks.c", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Less",
+          "shader": "compare-broadcast.wgsl.jinja",
+          "derive": {
+            "aShape": "shapes.a",
+            "bShape": "shapes.b",
+            "cShape": "shapes.c",
+            "aRank": "ranks.a",
+            "bRank": "ranks.b",
+            "cRank": "ranks.c",
+            "op": "\"less\"",
+            "itemsPerInvocation": 4
           },
+          "bindings": ["a_2", "b_2", "c_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }
+  ],
+  "bindings": {
+    "a": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "b": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "c": { "buffer": "storage", "elementType": "vec4<u32>" },
+    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c) / 4" }] },
+    "b_2": { "buffer": "read-only-storage", "name": "b", "elementType": "$scalar" },
+    "a_2": { "buffer": "read-only-storage", "name": "a", "elementType": "$scalar" },
+    "c_2": { "buffer": "storage", "name": "c", "elementType": "u32" },
+    "params_2": {
+      "buffer": "uniform",
+      "name": "params",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c)" }]
+    }
+  }
 }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,20 +1,30 @@
 {
   "name": "ai.onnx.Less",
-  "id": "_ai_onnx_less_webgpu_622f298",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "bi459h1j97iJXf9cdTyyrHoisesrs12uXVRMJmdV3jk=",
-      "compare-broadcast-vec4.wgsl.jinja": "tSx+Yubq+i8IwWTdoOppEsdRyIkK2mENofLE+InFA0w=",
-      "compare-broadcast.wgsl.jinja": "8JaFKuwhZhkyz/5Z52Y12lDykGVTeSupcVNarIVl/3k=",
-      "compare-vec4.wgsl.jinja": "cXSOC920mmdSmQNgxaMiUZDh4jmFywSvCjPxMr8SAjI=",
-      "manifest.json": "wAE9MngLrMCvMj2SOAHfFV2Z06CFHidS/BltK6nFkr4=",
-      "test.json": "1ZrEwckIRspwkQp9fLSrgGzAcBEUXKJKNXIU81H+DYE="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Less" }
 }

 {
   "name": "ai.onnx.Less",
+  "id": "_ai_onnx_less_webgpu_cd5aeae",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "DYOpGhl+2huEW8ES9/klafDCuls+BKZycGYBsG3xuXc=",
+      "compare-broadcast-vec4.wgsl.jinja": "e9IjvIN8zaZuy3Oxyo0917Z43virxHPxIKYg4GD5DDo=",
+      "compare-broadcast.wgsl.jinja": "149Y9wl0draPoa3sCNPbGVSvEk3Z4G0Rt1lifzBbCFM=",
+      "compare-vec4.wgsl.jinja": "HHor4TZ32hagFcdLFOoBRtFnH6CyrYjaIcviPAUhSWM=",
+      "manifest.json": "IhvdJxWOC3OjvHtb6mjf/7GCfUAHPm1gIjVpfWwLKNk=",
+      "test.json": "BkFQK6dgyaPaA0W1QSVZO1MlUw5afjqO9YwzV1lymrg="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "same_shape_vec4": ["compare-vec4.wgsl.jinja"],
+      "scalar_b_vec4": ["compare-vec4.wgsl.jinja"],
+      "scalar_a_vec4": ["compare-vec4.wgsl.jinja"],
+      "broadcast_vec4": ["compare-broadcast-vec4.wgsl.jinja"],
+      "broadcast": ["compare-broadcast.wgsl.jinja"],
+      "same_shape_scalar_x4": ["compare-broadcast.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,8 +1,8 @@
 {
-  "op": "ai.onnx.Less",
   "fixtureArrays": {
     "onnx_backend_input_a": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902],
-    "broadcast_vec4_lhs_splat_lower_rank_input_b": [9, 10, 11, 10, 19, 20, 21, 20, 29, 30, 31, 30, 30, 20, 10, 0, 10, 20, 30, 40, 31, 29, 30, 10]
   },
   "cases": [
     {
@@ -31,7 +31,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Less",
-        "notes": "Scalar-path companion: negative subnormals are strictly less than zero and must keep mask lanes true."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } },
@@ -167,11 +167,7 @@
           "shape": [4],
           "data": { "kind": "values", "values": ["NaN", 1.0, "Infinity", "-Infinity"] }
         },
-        "b": {
-          "dtype": "float32",
-          "shape": [4],
-          "data": { "kind": "values", "values": [0.0, "NaN", "Infinity", "-Infinity"] }
-        }
       },
       "outputs": { "c": { "dtype": "bool", "shape": [4], "tolerance": 0 } }
     },
@@ -435,18 +431,18 @@
     {
       "name": "f32_broadcast_odd_inner_dim_4x17",
       "provenance": {
-        "notes": "Compact f32 sibling for the odd-inner broadcast benchmark; preserves the [rows, odd] x [odd] shape family without benchmark-scale rows."
       },
       "inputs": {
         "a": {
           "dtype": "float32",
           "shape": [4, 17],
-          "data": { "kind": "fillFloat32", "sinStep": 0.037, "cosStep": 0.019, "scale": 0.5 }
         },
         "b": {
           "dtype": "float32",
           "shape": [17],
-          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.043, "scale": 0.5 }
         }
       },
       "outputs": { "c": { "dtype": "bool", "shape": [4, 17], "tolerance": 0 } }
@@ -478,9 +474,9 @@
     {
       "name": "broadcast_vec4_lhs_splat_lower_rank",
       "provenance": {
-        "source": "authored for render coverage",
         "test": "compare-broadcast-vec4 lhs splat arm",
-        "notes": "A is rank 2 against a rank 3 output and its innermost axis is 1, so compare-broadcast-vec4 renders `a_same` false straight from the rank test and takes the a_mode == \"splat\" branch, which reads one A element per vec4 group instead of four contiguous ones. Every other vec4-eligible case has A either the same rank as C (contiguous) or a single element (scalar), so neither the lower-rank A arm nor the A-splat arm had ever rendered. Each row of A holds a different value and B differs at every position, so a wrong a_offset changes the result."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [10.0, 20.0, 30.0] } },
@@ -498,7 +494,7 @@
     {
       "name": "broadcast_vec4_rhs_splat_lower_rank",
       "provenance": {
-        "source": "authored for render coverage",
         "test": "compare-broadcast-vec4 rhs splat arm",
         "notes": "B is rank 2 with innermost axis 1 against a rank-3 output, so each vec4 output group broadcasts one B value while A remains contiguous. This mirrors the corresponding left-hand splat case."
       },
@@ -518,9 +514,9 @@
     {
       "name": "broadcast_vec4_true_scalar_lhs",
       "provenance": {
-        "source": "authored for render coverage",
         "test": "compare-broadcast-vec4 rank-0 lhs arm",
-        "notes": "A is a true rank-0 scalar, so `source.aRank >= 1` is false and compare-broadcast-vec4 takes the left-hand scalar arm. C's innermost dimension is 4 and numel is 8, selecting broadcast_vec4 while distinguishing A-scalar from B-scalar indexing."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [20.0] } },
@@ -531,6 +527,36 @@
         }
       },
       "outputs": { "c": { "dtype": "bool", "shape": [2, 4], "tolerance": 0 } }
     }
   ]
 }

 {
   "fixtureArrays": {
     "onnx_backend_input_a": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902],
+    "broadcast_vec4_lhs_splat_lower_rank_input_b": [9, 10, 11, 10, 19, 20, 21, 20, 29, 30, 31, 30, 30, 20, 10, 0, 10, 20, 30, 40, 31, 29, 30, 10],
+    "scalarOperandRouteValues": [0.5, -1, 2, 0.25, 3, -2, 0.25, 1, 0.75, 0.25, -0.5, 4, 0.125, 0.25, -3, 2.5]
   },
   "cases": [
     {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
         "test": "MathOpTest.Less",
+        "notes": "On an unaligned scalar path, negative subnormals are strictly less than zero and must produce true mask lanes."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } },
           "shape": [4],
           "data": { "kind": "values", "values": ["NaN", 1.0, "Infinity", "-Infinity"] }
         },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, "NaN", "Infinity", 0.0] } }
       },
       "outputs": { "c": { "dtype": "bool", "shape": [4], "tolerance": 0 } }
     },
     {
       "name": "f32_broadcast_odd_inner_dim_4x17",
       "provenance": {
+        "notes": "A compact float32 [rows, odd] by [odd] broadcast exercises an odd inner dimension without benchmark-scale rows."
       },
       "inputs": {
         "a": {
           "dtype": "float32",
           "shape": [4, 17],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.53, "scale": 0.5 }
         },
         "b": {
           "dtype": "float32",
           "shape": [17],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.71, "scale": 0.5 }
         }
       },
       "outputs": { "c": { "dtype": "bool", "shape": [4, 17], "tolerance": 0 } }
     {
       "name": "broadcast_vec4_lhs_splat_lower_rank",
       "provenance": {
+        "source": "constructed fixture",
         "test": "compare-broadcast-vec4 lhs splat arm",
+        "notes": "A has rank 2 against a rank-3 output and an innermost dimension of 1, so each A row is splatted across one output row. Distinct values in every A row and B position make an incorrect broadcast offset observable."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [10.0, 20.0, 30.0] } },
     {
       "name": "broadcast_vec4_rhs_splat_lower_rank",
       "provenance": {
+        "source": "constructed fixture",
         "test": "compare-broadcast-vec4 rhs splat arm",
         "notes": "B is rank 2 with innermost axis 1 against a rank-3 output, so each vec4 output group broadcasts one B value while A remains contiguous. This mirrors the corresponding left-hand splat case."
       },
     {
       "name": "broadcast_vec4_true_scalar_lhs",
       "provenance": {
+        "source": "constructed fixture",
         "test": "compare-broadcast-vec4 rank-0 lhs arm",
+        "notes": "A is a rank-0 scalar that broadcasts to every output position. An innermost dimension of four and eight total outputs exercise vectorized comparison, while distinct B values make scalar-A indexing observable."
       },
       "inputs": {
         "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [20.0] } },
         }
       },
       "outputs": { "c": { "dtype": "bool", "shape": [2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "scalar_b_vec4_route",
+      "provenance": {
+        "notes": "Route lock for the one-element-operand vec4 kernel: the other operand matches the output and the vector count is a multiple of four, so the scalar reads once and splats. Values straddle and hit the scalar exactly, so the compare emits both results."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scalarOperandRouteValues" } }
+        },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": { "c": { "dtype": "bool", "shape": [2, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "scalar_a_vec4_route",
+      "provenance": {
+        "notes": "Route lock for the one-element-operand vec4 kernel: the other operand matches the output and the vector count is a multiple of four, so the scalar reads once and splats. Values straddle and hit the scalar exactly, so the compare emits both results."
+      },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/scalarOperandRouteValues" } }
+        }
+      },
+      "outputs": { "c": { "dtype": "bool", "shape": [2, 8], "tolerance": 0 } }
     }
   ]
 }