sync 2e7068faf55e

Browse files

Files changed (9) hide show

README.md +67 -0
build/webgpu/bench.json +82 -0
build/webgpu/hardmax-axis-tree.wgsl.jinja +60 -0
build/webgpu/hardmax-last-axis-subgroup.wgsl.jinja +119 -0
build/webgpu/hardmax-last-axis-vec4.wgsl.jinja +83 -0
build/webgpu/hardmax.wgsl.jinja +50 -0
build/webgpu/manifest.json +211 -0
build/webgpu/metadata.json +21 -0
build/webgpu/test.json +613 -0

README.md CHANGED Viewed

@@ -1,3 +1,70 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Hardmax
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Computes the hardmax of the input along a single axis: sets the position of the first maximum value along `axis` to 1 and all other positions to 0. The output has the same shape as the input.
+See the [ONNX `Hardmax` spec](https://onnx.ai/onnx/operators/onnx__Hardmax.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `x` | `T` | — | — | Input tensor with rank at least 1. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `y` | `T` | same as `input` | same as `input` | The output tensor with the same shape as the input, containing hardmax values. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axis` | `-1` | The dimension along which hardmax is computed. Negative values count from the back; accepted range is `[-r, r-1]` where `r` is the rank of the input. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## 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
+- [`hardmax-axis-tree.wgsl.jinja`](build/webgpu/hardmax-axis-tree.wgsl.jinja)
+- [`hardmax-last-axis-subgroup.wgsl.jinja`](build/webgpu/hardmax-last-axis-subgroup.wgsl.jinja)
+- [`hardmax-last-axis-vec4.wgsl.jinja`](build/webgpu/hardmax-last-axis-vec4.wgsl.jinja)
+- [`hardmax.wgsl.jinja`](build/webgpu/hardmax.wgsl.jinja)
+## 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.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Hardmax", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [1, 3] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,82 @@

+{
+  "op": "ai.onnx.Hardmax",
+  "cases": [
+    {
+      "name": "rows_4096_cols_128",
+      "vars": { "rows": 4096, "cols": 128, "dtype": "float32" },
+      "inputs": { "x": { "dtype": "float32", "shape": [4096, 128] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4096, 128] } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }]
+      },
+      "preset": "smoke"
+    },
+    {
+      "name": "rows_4096_cols_512",
+      "vars": { "rows": 4096, "cols": 512, "dtype": "float32" },
+      "inputs": { "x": { "dtype": "float32", "shape": [4096, 512] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4096, 512] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "rows_4096_cols_128_f16",
+      "vars": { "rows": 4096, "cols": 128, "dtype": "float16" },
+      "inputs": { "x": { "dtype": "float16", "shape": [4096, 128] } },
+      "outputs": { "y": { "dtype": "float16", "shape": [4096, 128] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "rows_1024_cols_1025",
+      "vars": { "rows": 1024, "cols": 1025, "dtype": "float32" },
+      "inputs": { "x": { "dtype": "float32", "shape": [1024, 1025] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024, 1025] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "rank3_axis1_scalar_fallback",
+      "vars": { "outer": 1024, "axisDim": 1024, "inner": 8, "dtype": "float32" },
+      "inputs": { "x": { "dtype": "float32", "shape": [1024, 1024, 8] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024, 1024, 8] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.outer * args.axisDim * args.inner * dtypeBytes(args.dtype) * 2" }
+        ]
+      }
+    },
+    {
+      "name": "unaligned_last_axis_127_scalar_cliff",
+      "vars": { "rows": 16384, "cols": 127, "dtype": "float32" },
+      "inputs": { "x": { "dtype": "float32", "shape": [16384, 127] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16384, 127] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "rows_1024_cols_1025_f16_subgroup",
+      "vars": { "rows": 1024, "cols": 1025, "dtype": "float16" },
+      "inputs": { "x": { "dtype": "float16", "shape": [1024, 1025] } },
+      "outputs": { "y": { "dtype": "float16", "shape": [1024, 1025] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "axis0_huge_axis_few_threads_launch_poor",
+      "preset": "stress",
+      "vars": { "outer": 1, "axisDim": 65536, "inner": 64, "dtype": "float32" },
+      "attrs": { "axis": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [65536, 64], "dist": "normal", "seed": 1301, "scale": 2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [65536, 64], "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.axisDim * args.inner * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "unaligned_last_axis_1023_scalar_row_cliff",
+      "preset": "edge",
+      "vars": { "rows": 16384, "cols": 1023, "dtype": "float32" },
+      "attrs": { "axis": -1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [16384, 1023], "dist": "normal", "seed": 2207, "scale": 2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16384, 1023], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/hardmax-axis-tree.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,60 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const AXIS_DIM: u32 = {{ axisDim }}u;
+const INNER: u32 = {{ inner }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const NEG_INF: f32 = -3.4028234663852886e38;
+var<workgroup> values: array<f32, WG>;
+var<workgroup> indices: array<u32, WG>;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let row = wg.x + wg.y * params.rowStride;
+  if (row >= params.rows) { return; }
+  let tid = lid.x;
+  let outer = row / INNER;
+  let inner_index = row % INNER;
+  let base = outer * AXIS_DIM * INNER + inner_index;
+  let first = f32(x[base]);
+  var best_value = NEG_INF;
+  var best_axis = 0u;
+  for (var a = tid; a < AXIS_DIM; a = a + WG) {
+    let value = f32(x[base + a * INNER]);
+    if (!is_nan_f32(value) && (value > best_value || (value == best_value && a < best_axis))) {
+      best_value = value;
+      best_axis = a;
+    }
+  }
+  values[tid] = best_value;
+  indices[tid] = best_axis;
+  workgroupBarrier();
+  for (var stride = WG >> 1u; stride > 0u; stride = stride >> 1u) {
+    if (tid < stride) {
+      let rv = values[tid + stride];
+      let ri = indices[tid + stride];
+      if (rv > values[tid] || (rv == values[tid] && ri < indices[tid])) {
+        values[tid] = rv;
+        indices[tid] = ri;
+      }
+    }
+    workgroupBarrier();
+  }
+  // The serial definition seeds from element zero: a NaN there wins forever;
+  // later NaNs never satisfy `value > best_value` and are ignored. It then
+  // marks where the value EQUALS the winner, which a NaN winner never does, so
+  // the reduced row is all zeros — an out-of-range index says exactly that.
+  let winner = select(indices[0], AXIS_DIM, is_nan_f32(first));
+  for (var a = tid; a < AXIS_DIM; a = a + WG) {
+    y[base + a * INNER] = {{ scalar }}(select(0.0, 1.0, a == winner));
+  }
+}

build/webgpu/hardmax-last-axis-subgroup.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,119 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{% if useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const SENTINEL_IDX: u32 = 4294967295u;
+fn neg_inf_f32() -> f32 {
+  var bits = 0xff800000u;
+  return bitcast<f32>(bits);
+}
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+var<workgroup> wgVal: array<f32, WG>;
+var<workgroup> wgIdx: array<u32, WG>;
+var<workgroup> rowBestIdx: u32;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>{% if useSubgroups %},
+        @builtin(subgroup_invocation_id) sgLid: u32,
+        @builtin(subgroup_size) sgSize: u32{% endif %}) {
+  let row = wg.x + wg.y * params.rowStride;
+  if (row >= params.rows) {
+    return;
+  }
+  let tid = lid.x;
+  let base = row * params.cols;
+  let firstValue = f32(x[base]);
+  var bestVal = neg_inf_f32();
+  var bestIdx = SENTINEL_IDX;
+  for (var c = tid; c < params.cols; c = c + WG) {
+    let v = f32(x[base + c]);
+    if (!is_nan_f32(v) && (v > bestVal || (v == bestVal && c < bestIdx))) {
+      bestVal = v;
+      bestIdx = c;
+    }
+  }
+{% if useSubgroups %}
+  let m = subgroupMax(bestVal);
+  let cand = select(SENTINEL_IDX, bestIdx, bestVal == m);
+  let sgIdx = subgroupMin(cand);
+  let safeSg = max(sgSize, 1u);
+  let slotCount = max(1u, WG / safeSg);
+  if (sgLid == 0u) {
+    let slot = min(tid / safeSg, WG - 1u);
+    wgVal[slot] = m;
+    wgIdx[slot] = sgIdx;
+  }
+  workgroupBarrier();
+  if (tid == 0u) {
+    var outVal = wgVal[0];
+    var outIdx = wgIdx[0];
+    for (var i = 1u; i < slotCount; i = i + 1u) {
+      let v = wgVal[i];
+      let vi = wgIdx[i];
+      if (vi != SENTINEL_IDX && (v > outVal || (v == outVal && vi < outIdx))) {
+        outVal = v;
+        outIdx = vi;
+      }
+    }
+    // Match the scalar shader: if x[row, 0] is NaN no later comparison can
+    // replace it, and an all-NaN row selects nothing. The operator writes 1 only
+    // where the value equals the winner, so both cases emit an all-zero row; an index
+    // past params.cols is how that is expressed here.
+    if (is_nan_f32(firstValue) || outIdx == SENTINEL_IDX) {
+      outIdx = params.cols;
+    }
+    rowBestIdx = outIdx;
+  }
+  workgroupBarrier();
+{% else %}
+  // No-subgroup tier: workgroup barrier tree-reduction over the per-thread
+  // (bestVal, bestIdx) pairs with the same lowest-index tie rule. WG is a power
+  // of two (the specialized pow2ceil value); SENTINEL_IDX lanes never win.
+  wgVal[tid] = bestVal;
+  wgIdx[tid] = bestIdx;
+  workgroupBarrier();
+  for (var stride = WG / 2u; stride > 0u; stride = stride >> 1u) {
+    if (tid < stride) {
+      let v = wgVal[tid + stride];
+      let vi = wgIdx[tid + stride];
+      let curIdx = wgIdx[tid];
+      if (vi != SENTINEL_IDX && (curIdx == SENTINEL_IDX || v > wgVal[tid] || (v == wgVal[tid] && vi < curIdx))) {
+        wgVal[tid] = v;
+        wgIdx[tid] = vi;
+      }
+    }
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    var outIdx = wgIdx[0];
+    if (is_nan_f32(firstValue) || outIdx == SENTINEL_IDX) {
+      outIdx = params.cols;
+    }
+    rowBestIdx = outIdx;
+  }
+  workgroupBarrier();
+{% endif %}
+  let outIdx = rowBestIdx;
+  for (var c = tid; c < params.cols; c = c + WG) {
+    y[base + c] = {{ scalar }}(select(0.0, 1.0, c == outIdx));
+  }
+}

build/webgpu/hardmax-last-axis-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,83 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const COLS_VEC: u32 = {{ colsVec }}u;
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+@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 past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+  let row = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (row >= params.rows) {
+    return;
+  }
+  let base = row * COLS_VEC;
+  let first4 = x[base];
+  let firstValue = f32(first4.x);
+  var bestValue = firstValue;
+  var bestAxis = 0u;
+  if (!is_nan_f32(firstValue)) {
+    for (var c4 = 0u; c4 < COLS_VEC; c4 = c4 + 1u) {
+      let v4 = x[base + c4];
+      let axisBase = c4 * 4u;
+      {
+        let value = f32(v4.x);
+        if (!is_nan_f32(value) && value > bestValue) {
+          bestValue = value;
+          bestAxis = axisBase;
+        }
+      }
+      {
+        let value = f32(v4.y);
+        if (!is_nan_f32(value) && value > bestValue) {
+          bestValue = value;
+          bestAxis = axisBase + 1u;
+        }
+      }
+      {
+        let value = f32(v4.z);
+        if (!is_nan_f32(value) && value > bestValue) {
+          bestValue = value;
+          bestAxis = axisBase + 2u;
+        }
+      }
+      {
+        let value = f32(v4.w);
+        if (!is_nan_f32(value) && value > bestValue) {
+          bestValue = value;
+          bestAxis = axisBase + 3u;
+        }
+      }
+    }
+  }
+  // A NaN at element zero wins the ordered scan and then matches nothing on the
+  // equality write, so the row comes out all zeros. COLS_VEC * 4 is past every
+  // position written below.
+  if (is_nan_f32(firstValue)) {
+    bestAxis = COLS_VEC * 4u;
+  }
+  let zero = {{ scalar }}(0.0);
+  let one = {{ scalar }}(1.0);
+  for (var c4 = 0u; c4 < COLS_VEC; c4 = c4 + 1u) {
+    let axisBase = c4 * 4u;
+    y[base + c4] = vec4<{{ scalar }}>(
+      select(zero, one, axisBase == bestAxis),
+      select(zero, one, axisBase + 1u == bestAxis),
+      select(zero, one, axisBase + 2u == bestAxis),
+      select(zero, one, axisBase + 3u == bestAxis)
+    );
+  }
+}

build/webgpu/hardmax.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,50 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+const AXIS_DIM: u32 = {{ axisDim }}u;
+const INNER: u32 = {{ inner }}u;
+@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 past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+  let row = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (row >= params.rows) {
+    return;
+  }
+  let outer = row / INNER;
+  let inner_index = row % INNER;
+  let base = outer * AXIS_DIM * INNER + inner_index;
+  var best_axis = 0u;
+  var best_value = f32(x[base]);
+  for (var a = 1u; a < AXIS_DIM; a = a + 1u) {
+    let value = f32(x[base + a * INNER]);
+    if (value > best_value) {
+      best_value = value;
+      best_axis = a;
+    }
+  }
+  // The operator seeds the running best at element zero and only replaces it
+  // on a strict `>`, so a NaN there wins forever — and it then writes 1 where
+  // the value EQUALS the winner, which NaN never does. The row is therefore all
+  // zeros, not a 1 in slot 0. Steering the winning index out of range is how
+  // that falls out here.
+  // Use a bitcast test because no-NaN optimization can fold `v != v` to false.
+  if (is_nan_f32(best_value)) {
+    best_axis = AXIS_DIM;
+  }
+  for (var a = 0u; a < AXIS_DIM; a = a + 1u) {
+    y[base + a * INNER] = {{ scalar }}(select(0.0, 1.0, a == best_axis));
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,211 @@

+{
+  "domain": "ai.onnx",
+  "name": "Hardmax",
+  "sinceVersion": 13,
+  "description": "Computes the hardmax of the input along a single axis: sets the position of the first maximum value along `axis` to 1 and all other positions to 0. The output has the same shape as the input.",
+  "inputs": [{ "role": "input", "dtype": "T", "description": "Input tensor with rank at least 1." }],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": "ranks.input",
+      "description": "The output tensor with the same shape as the input, containing hardmax values.",
+      "shape": "shapes.input"
+    }
+  ],
+  "attributes": { "axis": -1 },
+  "attributeDescriptions": {
+    "axis": "The dimension along which hardmax is computed. Negative values count from the back; accepted range is `[-r, r-1]` where `r` is the rank of the input."
+  },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "bindingSets": {
+    "subgroupRow": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "numel(shapes.x) / dim(shapes.x, ranks.x - 1)" },
+            { "name": "cols", "type": "u32", "value": "dim(shapes.x, ranks.x - 1)" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(numel(shapes.x) / dim(shapes.x, ranks.x - 1), device.limits.maxComputeWorkgroupsPerDimension))"
+            }
+          ]
+        }
+      }
+    ]
+  },
+  "constants": { "usesF16": "dtypes.T == \"f16\"", "scalar": "dtypes.T" },
+  "variants": [
+    {
+      "id": "last_axis_row",
+      "priority": 20,
+      "when": ["ranks.x >= 1", "ranks.y == ranks.x", "numel(shapes.x) == numel(shapes.y)", "(attrs.axis == -1 or attrs.axis == ranks.x - 1)", "dim(shapes.x, ranks.x - 1) >= 1024", "f16Ok(dtypes.T)"],
+      "constants": {
+        "workgroupSize": "min(tunables.WORKGROUP_SIZE, max(32, pow2ceil(dim(shapes.x, ranks.x - 1))))",
+        "useSubgroups": "device.features.has(\"subgroups\")"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Hardmax.LastAxisRow",
+          "shader": "hardmax-last-axis-subgroup.wgsl.jinja",
+          "bindings": "subgroupRow",
+          "dispatch": { "workgroups": "numel(shapes.x) / dim(shapes.x, ranks.x - 1)" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_vec4",
+      "priority": 30,
+      "when": ["ranks.x >= 1", "ranks.y == ranks.x", "numel(shapes.x) == numel(shapes.y)", "(attrs.axis == -1 or attrs.axis == ranks.x - 1)", "dim(shapes.x, ranks.x - 1) >= 4", "dim(shapes.x, ranks.x - 1) % 4 == 0", "f16Ok(dtypes.T)"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"", "colsVec": "dim(shapes.x, ranks.x - 1) / 4" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Hardmax.LastAxisVec4",
+          "shader": "hardmax-last-axis-vec4.wgsl.jinja",
+          "bindings": [
+            {
+              "name": "x",
+              "arg": "x",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "y",
+              "arg": "y",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "rows", "type": "u32", "value": "numel(shapes.x) / dim(shapes.x, ranks.x - 1)" }]
+              }
+            }
+          ],
+          "dispatch": {
+            "threads": "numel(shapes.x) / dim(shapes.x, ranks.x - 1)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "axis_tree",
+      "priority": 21,
+      "demoteWhen": ["device.features.has(\"subgroups\") and (attrs.axis == -1 or attrs.axis == ranks.x - 1)"],
+      "when": ["ranks.x >= 1", "ranks.y == ranks.x", "sameShape(shapes.y, shapes.x)", "attrs.axis + ranks.x >= 0", "attrs.axis < ranks.x", "dim(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x) >= 1024", "rows(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x) <= device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension", "f16Ok(dtypes.T)"],
+      "constants": {
+        "axisDim": "dim(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)",
+        "inner": "inner(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Hardmax.AxisTree",
+          "shader": "hardmax-axis-tree.wgsl.jinja",
+          "bindings": [
+            {
+              "name": "x",
+              "arg": "x",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "y", "arg": "y", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  {
+                    "name": "rows",
+                    "type": "u32",
+                    "value": "rows(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)"
+                  },
+                  {
+                    "name": "rowStride",
+                    "type": "u32",
+                    "value": "max(1, min(rows(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x), device.limits.maxComputeWorkgroupsPerDimension))"
+                  }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "workgroups": "rows(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)" }
+        }
+      ]
+    },
+    {
+      "id": "axis",
+      "when": ["ranks.x >= 1", "ranks.y == ranks.x", "numel(shapes.x) == numel(shapes.y)", "attrs.axis + ranks.x >= 0", "attrs.axis < ranks.x", "f16Ok(dtypes.T)"],
+      "constants": {
+        "axis": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x",
+        "axisDim": "dim(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)",
+        "inner": "inner(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Hardmax",
+          "shader": "hardmax.wgsl.jinja",
+          "bindings": [
+            {
+              "name": "x",
+              "arg": "x",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "y", "arg": "y", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  {
+                    "name": "rows",
+                    "type": "u32",
+                    "value": "rows(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)"
+                  }
+                ]
+              }
+            }
+          ],
+          "dispatch": {
+            "threads": "rows(shapes.x, attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.x)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,21 @@

+{
+  "name": "ai.onnx.Hardmax",
+  "id": "_ai_onnx_hardmax_webgpu_0d3207a",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "bZefItlvlTUgPlW5qCO0jHoXmOmOqQo8vZMsx7mgh9A=",
+      "hardmax-axis-tree.wgsl.jinja": "1Qd9u20cZKcxIcoVUn8UpI27dye5BPfXCtyLSIpmNq8=",
+      "hardmax-last-axis-subgroup.wgsl.jinja": "YSX63Lf95J690RQT40ZWuwbBXDLOWNZp87ngY9T0UWE=",
+      "hardmax-last-axis-vec4.wgsl.jinja": "4S/FaIFjlrXr1LuWMV5X+3MgBQ6FA4JMZioHRoCjVfQ=",
+      "hardmax.wgsl.jinja": "xRSmORDzxQisVCQ78ZE+txGsYp+E/TBp8et510dk1mI=",
+      "manifest.json": "bbiRLG6GXQQCiUUgEJEgl943Ec7p4oPPgR5QiXLmT2c=",
+      "test.json": "CZ8WfjN4jnNK9CE1kbRfFuJeCyntQ40ql17EsKp87JI="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Hardmax" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,613 @@

+{
+  "op": "ai.onnx.Hardmax",
+  "fixtureArrays": {
+    "ort_three_dims_input_x": [1.0856307, 0.99734545, 0.2829785, 1.5062947, 0.5786002, 1.6514366, 2.4266791, 0.42891264, 1.2659363, 0.8667404, 0.6788862, 0.09470897, 1.4913896, 0.638902, 0.44398195, 0.43435127, 2.20593, 2.1867862, 1.004054, 0.3861864, 0.7373686, 1.4907321, 0.9358339, 1.175829, 1.2538806, 0.6377515, 0.9071052, 1.4286807, 0.14006872, 0.8617549, 0.25561938, 2.798589, 1.7715331, 0.69987726, 0.92746246, 0.17363568, 0.002845916, 0.6882227, 0.87953633, 0.28362733, 0.8053665, 1.7276695, 0.3908998, 0.57380587, 0.33858904, 0.011830495, 2.3923652, 0.41291216, 0.978736, 2.2381434, 1.2940853, 1.0387882, 1.7437122, 0.79806274, 0.02968323, 1.0693159, 0.8907064, 1.7548862, 1.4956441, 1.0693927],
+    "onnx_backend_hardmax_input_x": [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],
+    "onnx_backend_example_4x4_input_x": [3, 0, 1, 2, 2, 5, 1, 0, 0, 1, 3, 2, 0, 1, 2, 3]
+  },
+  "cases": [
+    {
+      "name": "f32_positive_subnormal_beats_zero_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; bit-exact subnormal preservation is unattainable on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.Simple",
+        "notes": "A positive subnormal is strictly greater than zero; Hardmax should put the one-hot at that element rather than tie-breaking on a flushed zero."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_positive_subnormal_beats_zero_axis0_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; bit-exact subnormal preservation is unattainable on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.Simple",
+        "notes": "Axis-0 companion for finite subnormal ordering: a positive subnormal is strictly greater than zero and should receive the one-hot."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_positive_subnormal_beats_zero_last_axis_vec4_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; bit-exact subnormal preservation is unattainable on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.Simple",
+        "notes": "Vec4 last-axis companion: a positive subnormal is strictly greater than zero and should win over zero-valued lanes."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 4], "data": { "kind": "values", "values": [0.0, 1e-40, 0.0, -1e-40] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 1.0, 0.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "f32_positive_subnormal_beats_zero_last_axis_subgroup_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; bit-exact subnormal preservation is unattainable on GPU."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.LargeNumber",
+        "notes": "Subgroup last-axis companion: width 1025 bypasses the vec4 specialization, and a positive subnormal must still beat zero-valued lanes."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1025],
+          "data": { "kind": "cycle", "values": [0.0, 1e-40, 0.0, -1e-40] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1025], "tolerance": 0 } }
+    },
+    {
+      "name": "dispatch_cliff_axis_last_dim1",
+      "attrs": { "axis": -1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [16776961, 1], "data": { "kind": "constant", "value": 1.0 } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16776961, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "simple_last_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.Simple"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3] } }
+    },
+    {
+      "name": "onnx_backend_example_4x4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_example",
+        "test": "test_hardmax_example"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_example_4x4_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_one_hot_tie_first",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_one_hot",
+        "test": "test_hardmax_one_hot"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 4], "data": { "kind": "values", "values": [3.0, 3.0, 3.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "large_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.LargeNumber"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 10000.0, 10001.0, 10002.0, 10003.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4] } }
+    },
+    {
+      "name": "ort_compact_axis0_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.ThreeDimsAxis0",
+        "notes": "Compact axis-0 variant covering the same non-last-axis behavior without the full ORT random tensor."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 5.0, 9.0, 4.0, 7.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "axis1_rank3",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [1.0, 9.0, 7.0, 4.0, 5.0, 6.0, 3.0, 2.0, 8.0, 1.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2] } }
+    },
+    {
+      "name": "negative_axis",
+      "attrs": { "axis": -2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [1.0, 9.0, 7.0, 4.0, 5.0, 6.0, 3.0, 2.0, 8.0, 1.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2] } }
+    },
+    {
+      "name": "ties_choose_first",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, 3.0, 3.0, 2.0, 5.0, 5.0, 1.0, 5.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4] } }
+    },
+    {
+      "name": "ort_axis1_nan_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.Simple",
+        "notes": "Compatibility NaN behavior on the last axis: a leading NaN or all-NaN row produces all zeros, while a NaN after a finite incumbent does not replace the selected maximum."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "data": {
+            "kind": "values",
+            "values": ["NaN", 1.0, 2.0, 0.0, 1.0, "NaN", 2.0, 0.0, 1.0, 2.0, "NaN", 0.0, "NaN", "NaN", "NaN", "NaN"]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_rank3_axis1_nan_columns",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.ThreeDimsAxis1_opset13",
+        "notes": "Compatibility NaN behavior on a non-last axis: leading-NaN and all-NaN reduced columns produce all zeros, while later NaNs after a finite value are ignored."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "data": {
+            "kind": "values",
+            "values": ["NaN", 1.0, 1.0, "NaN", 2.0, 2.0, 1.0, "NaN", 2.0, "NaN", "NaN", "NaN"]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "float16_axis0",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 5.0, 4.0, 2.0, 3.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [3, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "rank4_axis2_ties_choose_first",
+      "attrs": { "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 2],
+          "data": { "kind": "values", "values": [1.0, 9.0, 3.0, 9.0, 3.0, 8.0, 5.0, 0.0, 5.0, 1.0, 4.0, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "float16_ties_choose_first",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, 3.0, 3.0, 2.0, 5.0, 5.0, 1.0, 5.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "last_axis_64_tie_first_f32",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 64],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 100.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0, 33.0, 34.0, 35.0, 36.0, 100.0, 38.0, 39.0, 40.0, 41.0, 42.0, 43.0, 44.0, 45.0, 46.0, 47.0, 48.0, 49.0, 50.0, 51.0, 52.0, 53.0, 54.0, 55.0, 56.0, 57.0, 58.0, 59.0, 60.0, 61.0, 62.0, 63.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64], "tolerance": 0 } }
+    },
+    {
+      "name": "last_axis_128_f16",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 128], "tolerance": 0 } }
+    },
+    {
+      "name": "last_axis_1025_f32",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1025],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1025], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_three_dims_axis1_opset13",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.ThreeDimsAxis1_opset13"
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_three_dims_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_three_dims_default_axis_opset13",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.ThreeDimsDefaultAxis_opset13"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_three_dims_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_three_dims_axis2_opset13",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.ThreeDimsAxis2_opset13"
+      },
+      "attrs": { "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_three_dims_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_three_dims_negative_axis_last",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/hardmax_test.cc",
+        "test": "HardmaxOperator.ThreeDimsNegAxis2",
+        "notes": "Axis=-1 maps to the last dimension; this matches opset-13 last-axis semantics."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_three_dims_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_hardmax_axis_0",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_axis_0" },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_hardmax_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_hardmax_axis_1",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_axis_1" },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_hardmax_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_hardmax_axis_2",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_axis_2" },
+      "attrs": { "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_hardmax_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_hardmax_default_axis",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_default_axis" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_hardmax_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_hardmax_negative_axis",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_negative_axis" },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_hardmax_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_hardmax_example",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_example" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_example_4x4_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 4], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_hardmax_one_hot",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hardmax_one_hot" },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 4], "data": { "kind": "values", "values": [3.0, 3.0, 3.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "inputs": { "x": { "dtype": "float32", "shape": [0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "rank7_last_axis",
+      "provenance": {
+        "source": "ONNX spec: Hardmax permits arbitrary rank; onnxruntime CPU coerces dims [0,axis) to rows and [axis,end) to cols with no rank cap.",
+        "notes": "Rank-7 last-axis coverage for the generalized last_axis_vec4 path. ORT and WebGPU both compute the normal one-hot result."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 1, 1, 1, 4],
+          "data": { "kind": "values", "values": [3.0, 1.0, 4.0, 2.0, 5.0, 9.0, 1.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 1, 1, 1, 1, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank7_axis0",
+      "provenance": {
+        "source": "ONNX spec: Hardmax permits arbitrary rank; onnxruntime CPU reduces over the requested axis at any rank.",
+        "notes": "Rank-7 non-last-axis coverage for the generic axis path. ORT and WebGPU reduce over axis 0 per inner column."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 1, 1, 1, 1, 1, 2],
+          "data": { "kind": "values", "values": [1.0, 5.0, 9.0, 4.0, 7.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 1, 1, 1, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_last_axis_subgroup_1025",
+      "provenance": {
+        "source": "Clean coverage: the adaptive f16 last-axis row path (dim(last) >= 1024, f16Ok) was correctness-untested (only f16 vec4 cols=128 and f16 axis-0 rank2 existed).",
+        "notes": "Last dim 1025 (>=1024, not %4==0) skips last_axis_vec4 and selects last_axis_row with f16. linspace 0..1025 gives exact distinct f16 integers (representable to 2048) so argmax is unambiguous at the final lane."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 1025], "data": { "kind": "linspace", "start": 0.0, "end": 1025.0 } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1025], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_axis1_rank3",
+      "provenance": {
+        "source": "Clean coverage: f16 'axis' scalar fallback on a non-last axis (existing f16 fixtures only cover axis-0 rank2 and last-axis rank2).",
+        "notes": "Rank-3 axis=1 f16 selects the generic 'axis' variant with usesF16. Distinct per-column values avoid ties."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [1.0, 9.0, 7.0, 4.0, 5.0, 6.0, 3.0, 2.0, 8.0, 1.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_lone_finite_among_neg_inf",
+      "attrs": { "axis": -1 },
+      "provenance": {
+        "source": "Coverage gap: vec4 last-axis path with -Infinity seed. Existing vec4 tests use finite fillFloat32/values; a row of all -Infinity except one finite element is untested.",
+        "notes": "cols=8 selects last_axis_vec4. Only index 6 is finite (3.0); every other lane is -Infinity, so the one-hot must land at index 6."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8],
+          "data": {
+            "kind": "values",
+            "values": ["-Infinity", "-Infinity", "-Infinity", "-Infinity", "-Infinity", "-Infinity", 3.0, "-Infinity"]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "subgroup_cross_slot_lowest_index_tie_1025",
+      "attrs": { "axis": -1 },
+      "provenance": {
+        "source": "Coverage gap: subgroup last-axis lowest-index tie reduction. Existing 1025-wide subgroup tests use unique-max linspace, so cross-subgroup/slot ties are untested.",
+        "notes": "cols=1025 (>=1024, %4!=0) selects last_axis_row, whose capability-adaptive reduction uses subgroup or portable execution. Alternating 0/1 makes 1.0 the max at every odd index; first occurrence is index 1, so the one-hot must land at index 1 despite ties spanning multiple reduction slots."
+      },
+      "inputs": { "x": { "dtype": "float32", "shape": [1, 1025], "data": { "kind": "cycle", "values": [0.0, 1.0] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1025], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_subgroup_all_equal_tie_index0_1025",
+      "attrs": { "axis": -1 },
+      "provenance": {
+        "source": "Coverage gap: the adaptive f16 last-axis row path under a full tie. The existing f16_last_axis_subgroup_1025 fixture uses distinct linspace values, so an all-equal f16 row is untested.",
+        "notes": "cols=1025 (>=1024, %4!=0) with f16 selects the adaptive last_axis_row path. All values equal (1.0) => first max is index 0, so the one-hot must land at index 0."
+      },
+      "inputs": { "x": { "dtype": "float16", "shape": [1, 1025], "data": { "kind": "constant", "value": 1.0 } } },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1025], "tolerance": 0 } }
+    },
+    {
+      "name": "axis_tree_axis0_1024x2",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1024, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 2.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank8_last_axis",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 1, 1, 1, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [3.0, 1.0, 4.0, 2.0, 5.0, 9.0, 1.0, 6.0, 2.0, 7.0, 1.0, 8.0, 0.5, 0.25, 3.5, 1.25]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 1, 1, 1, 1, 2, 4], "tolerance": 0 } }
+    }
+  ]
+}