sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +86 -0
build/webgpu/bench.json +76 -0
build/webgpu/manifest.json +730 -0
build/webgpu/metadata.json +20 -0
build/webgpu/nms-groups-compact.wgsl.jinja +49 -0
build/webgpu/nms-groups-select.wgsl.jinja +223 -0
build/webgpu/non-max-suppression.wgsl.jinja +126 -0
build/webgpu/test.json +1179 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.NonMaxSuppression
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 11
+## Description
+Filters boxes whose intersection over union (IoU) with a higher-scoring selected box exceeds `iou_threshold`. Processing is independent per batch and class, and returns logical int64 `[batch_index, class_index, box_index]` triples with lossless uint32 WebGPU storage. The int64 selection limit uses a semantics-preserving saturating projection to uint32. Because the result length is data-dependent, callers provide its exact shape. The current backend implements float32 boxes, scores, and thresholds.
+See the [ONNX `NonMaxSuppression` spec](https://onnx.ai/onnx/operators/onnx__NonMaxSuppression.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `boxes` | `boxes` | `T` | same as logical dtype | `3` | — | Bounding box coordinates with shape `[num_batches, spatial_dimension, 4]`; box format is controlled by `center_point_box`. | required |
+| `scores` | `scores` | `T` | same as logical dtype | `3` | — | Per-class confidence scores with shape `[num_batches, num_classes, spatial_dimension]`. | required |
+| `max_output_boxes_per_class` | `max_output_boxes_per_class` | `M` | `uint32` | `0` | — | Optional logical int64 scalar limiting boxes selected per batch and class. When omitted, the ONNX default is zero and the output is empty. WebGPU maps non-positive values to zero and values above uint32 range to `UINT32_MAX`; this preserves results because no group can select more boxes than its finite input. | optional |
+| `iou_threshold` | `iou_threshold` | `T` | same as logical dtype | `0` | — | Optional scalar IoU threshold in `[0, 1]`; boxes whose IoU is strictly greater are suppressed. Defaults to zero. | optional |
+| `score_threshold` | `score_threshold` | `T` | same as logical dtype | `0` | — | Optional scalar score threshold. When present, only boxes whose score is strictly greater than the threshold are considered; when omitted, scores are not filtered. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `selected_indices` | `selected_indices` | `I` | `uint32` | `2` | — | Logical int64 selected-box indices with shape `[num_selected_indices, 3]`, each row containing `[batch_index, class_index, box_index]`. WebGPU stores the bounded indices as uint32. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `center_point_box` | `0` | Box coordinate format: 0 for corner format `[y1, x1, y2, x2]`, 1 for center format `[x_center, y_center, width, height]`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32` |
+| `M` | `int64` |
+| `I` | `int64` |
+## 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
+- [`nms-groups-compact.wgsl.jinja`](build/webgpu/nms-groups-compact.wgsl.jinja)
+- [`nms-groups-select.wgsl.jinja`](build/webgpu/nms-groups-select.wgsl.jinja)
+- [`non-max-suppression.wgsl.jinja`](build/webgpu/non-max-suppression.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
+The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
+- `selected_indices`
+Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.NonMaxSuppression", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { selected_indices } = await kernel({
+  boxes: { data: boxesData, shape: [1, 1, 4] },
+  scores: { data: scoresData, shape: [1, 1, 1] },
+  max_output_boxes_per_class: { data: max_output_boxes_per_classData, shape: [] },
+  iou_threshold: { data: iou_thresholdData, shape: [] },
+  score_threshold: { data: score_thresholdData, shape: [] },
+}, {
+  outputs: { selected_indices: { shape: [1, 3], dtype: "int64" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,76 @@

+{
+  "op": "ai.onnx.NonMaxSuppression",
+  "cases": [
+    {
+      "name": "small_64_boxes",
+      "inputs": {
+        "boxes": { "dtype": "float32", "shape": [1, 64, 4], "data": { "kind": "constant", "value": 0.0 } },
+        "scores": { "dtype": "float32", "shape": [1, 1, 64], "data": { "kind": "constant", "value": 1.0 } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "constant", "value": 10 } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.5 } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.1 } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [10, 3] } }
+    },
+    {
+      "name": "nms-1class-2048-boxes",
+      "preset": "smoke",
+      "inputs": {
+        "boxes": { "dtype": "float32", "shape": [1, 2048, 4], "data": { "kind": "constant", "value": 0.0 } },
+        "scores": { "dtype": "float32", "shape": [1, 1, 2048], "data": { "kind": "constant", "value": 1.0 } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "constant", "value": 64 } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.5 } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [64, 3] } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.boxes) + numel(shapes.scores))" }]
+      }
+    },
+    {
+      "name": "nms-8class-1024-boxes",
+      "preset": "smoke",
+      "inputs": {
+        "boxes": { "dtype": "float32", "shape": [1, 1024, 4], "data": { "kind": "constant", "value": 0.0 } },
+        "scores": { "dtype": "float32", "shape": [1, 8, 1024], "data": { "kind": "constant", "value": 1.0 } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "constant", "value": 32 } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.5 } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [256, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.boxes) + numel(shapes.scores))" }] }
+    },
+    {
+      "name": "nms_single_class_4096_boxes_max_keep_wg_parallel",
+      "preset": "stress",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": { "dtype": "float32", "shape": [1, 4096, 4], "data": { "kind": "constant", "value": 0.0 } },
+        "scores": { "dtype": "float32", "shape": [1, 1, 4096], "data": { "kind": "constant", "value": 1.0 } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "constant", "value": 4096 } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.01 } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [4096, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.boxes) + numel(shapes.scores))" }] }
+    },
+    {
+      "name": "yolo-80class-8400-boxes-top100",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Worst-case exact-output workload: zero-area boxes never overlap and unit scores pass the threshold, so every class reaches max_output_boxes_per_class."
+      },
+      "attrs": { "center_point_box": 1 },
+      "inputs": {
+        "boxes": { "dtype": "float32", "shape": [1, 8400, 4], "data": { "kind": "constant", "value": 0.0 } },
+        "scores": { "dtype": "float32", "shape": [1, 80, 8400], "data": { "kind": "constant", "value": 1.0 } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "constant", "value": 100 } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.45 } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "constant", "value": 0.25 } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [8000, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.boxes) + numel(shapes.scores))" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,730 @@

+{
+  "domain": "ai.onnx",
+  "name": "NonMaxSuppression",
+  "sinceVersion": 11,
+  "description": "Filters boxes whose intersection over union (IoU) with a higher-scoring selected box exceeds `iou_threshold`. Processing is independent per batch and class, and returns logical int64 `[batch_index, class_index, box_index]` triples with lossless uint32 WebGPU storage. The int64 selection limit uses a semantics-preserving saturating projection to uint32. Because the result length is data-dependent, callers provide its exact shape. The current backend implements float32 boxes, scores, and thresholds.",
+  "inputs": [
+    {
+      "role": "boxes",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Bounding box coordinates with shape `[num_batches, spatial_dimension, 4]`; box format is controlled by `center_point_box`."
+    },
+    {
+      "role": "scores",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Per-class confidence scores with shape `[num_batches, num_classes, spatial_dimension]`."
+    },
+    {
+      "role": "max_output_boxes_per_class",
+      "dtype": "M",
+      "rank": 0,
+      "optional": true,
+      "description": "Optional logical int64 scalar limiting boxes selected per batch and class. When omitted, the ONNX default is zero and the output is empty. WebGPU maps non-positive values to zero and values above uint32 range to `UINT32_MAX`; this preserves results because no group can select more boxes than its finite input."
+    },
+    {
+      "role": "iou_threshold",
+      "dtype": "T",
+      "rank": 0,
+      "optional": true,
+      "description": "Optional scalar IoU threshold in `[0, 1]`; boxes whose IoU is strictly greater are suppressed. Defaults to zero."
+    },
+    {
+      "role": "score_threshold",
+      "dtype": "T",
+      "rank": 0,
+      "optional": true,
+      "description": "Optional scalar score threshold. When present, only boxes whose score is strictly greater than the threshold are considered; when omitted, scores are not filtered."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "selected_indices",
+      "dtype": "I",
+      "rank": 2,
+      "description": "Logical int64 selected-box indices with shape `[num_selected_indices, 3]`, each row containing `[batch_index, class_index, box_index]`. WebGPU stores the bounded indices as uint32."
+    }
+  ],
+  "attributes": { "center_point_box": 0 },
+  "attributeDescriptions": {
+    "center_point_box": "Box coordinate format: 0 for corner format `[y1, x1, y2, x2]`, 1 for center format `[x_center, y_center, width, height]`."
+  },
+  "attributeConstraints": { "center_point_box": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32"], "M": ["int64"], "I": ["int64"] },
+  "args": {
+    "boxes": { "kind": "tensor", "semantic": "boxes", "role": "input" },
+    "scores": { "kind": "tensor", "semantic": "scores", "role": "input" },
+    "max_output_boxes_per_class": {
+      "kind": "tensor",
+      "semantic": "max_output_boxes_per_class",
+      "role": "input",
+      "dtype": "uint32",
+      "narrowing": "saturating",
+      "required": false
+    },
+    "iou_threshold": { "kind": "tensor", "semantic": "iou_threshold", "role": "input", "required": false },
+    "score_threshold": { "kind": "tensor", "semantic": "score_threshold", "role": "input", "required": false },
+    "selected_indices": { "kind": "tensor", "semantic": "selected_indices", "role": "output", "dtype": "uint32" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "storageBufferLimit": "min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+    "validLayout": "ranks.boxes == 3 and ranks.scores == 3 and ranks.selected_indices == 2 and dim(shapes.boxes, 0) == dim(shapes.scores, 0) and dim(shapes.boxes, 1) == dim(shapes.scores, 2) and dim(shapes.boxes, 2) == 4 and dim(shapes.selected_indices, 1) == 3 and dim(shapes.selected_indices, 0) <= dim(shapes.boxes, 0) * dim(shapes.scores, 1) * dim(shapes.boxes, 1) and (present.max_output_boxes_per_class or dim(shapes.selected_indices, 0) == 0) and (not present.max_output_boxes_per_class or (ranks.max_output_boxes_per_class == 0 and numel(shapes.max_output_boxes_per_class) == 1)) and (not present.iou_threshold or (ranks.iou_threshold == 0 and numel(shapes.iou_threshold) == 1)) and (not present.score_threshold or (ranks.score_threshold == 0 and numel(shapes.score_threshold) == 1))",
+    "numGroups": "dim(shapes.boxes, 0) * dim(shapes.scores, 1)",
+    "maxKeep": "max(1, dim(shapes.boxes, 1))",
+    "parallelDispatchFits": "numGroups <= device.limits.maxComputeWorkgroupsPerDimension",
+    "parallelScratchFits": "numGroups * dim(shapes.boxes, 1) * 4 <= storageBufferLimit and numGroups * maxKeep * 4 <= storageBufferLimit and numGroups * 4 <= storageBufferLimit"
+  },
+  "constants": { "workgroupSize": "tunables.WORKGROUP_SIZE" },
+  "bindingSets": {
+    "parallelSelect": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "max_output_boxes_per_class",
+        "arg": "max_output_boxes_per_class",
+        "semantic": "max_output_boxes_per_class",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      {
+        "name": "iou_threshold",
+        "arg": "iou_threshold",
+        "semantic": "iou_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "score_threshold",
+        "arg": "score_threshold",
+        "semantic": "score_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      { "name": "suppressed", "semantic": "suppressed", "buffer": { "type": "storage" }, "elementType": "u32" },
+      { "name": "groupSel", "semantic": "groupSel", "buffer": { "type": "storage" }, "elementType": "u32" },
+      { "name": "groupCount", "semantic": "groupCount", "buffer": { "type": "storage" }, "elementType": "u32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "parallelCompact": [
+      { "name": "groupSel", "semantic": "groupSel", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      {
+        "name": "groupCount",
+        "semantic": "groupCount",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" }
+          ]
+        }
+      }
+    ],
+    "serialDefaults": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "serialMax": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "max_output_boxes_per_class",
+        "arg": "max_output_boxes_per_class",
+        "semantic": "max_output_boxes_per_class",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "serialIou": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "iou_threshold",
+        "arg": "iou_threshold",
+        "semantic": "iou_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "serialScore": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "score_threshold",
+        "arg": "score_threshold",
+        "semantic": "score_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "serialMaxIou": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "max_output_boxes_per_class",
+        "arg": "max_output_boxes_per_class",
+        "semantic": "max_output_boxes_per_class",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      {
+        "name": "iou_threshold",
+        "arg": "iou_threshold",
+        "semantic": "iou_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "serialMaxScore": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "max_output_boxes_per_class",
+        "arg": "max_output_boxes_per_class",
+        "semantic": "max_output_boxes_per_class",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      {
+        "name": "score_threshold",
+        "arg": "score_threshold",
+        "semantic": "score_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "serialIouScore": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "iou_threshold",
+        "arg": "iou_threshold",
+        "semantic": "iou_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "score_threshold",
+        "arg": "score_threshold",
+        "semantic": "score_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ],
+    "serialAll": [
+      {
+        "name": "boxes",
+        "arg": "boxes",
+        "semantic": "boxes",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scores",
+        "arg": "scores",
+        "semantic": "scores",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "max_output_boxes_per_class",
+        "arg": "max_output_boxes_per_class",
+        "semantic": "max_output_boxes_per_class",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      {
+        "name": "iou_threshold",
+        "arg": "iou_threshold",
+        "semantic": "iou_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "score_threshold",
+        "arg": "score_threshold",
+        "semantic": "score_threshold",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "selected_indices",
+        "arg": "selected_indices",
+        "semantic": "selected_indices",
+        "buffer": { "type": "storage" },
+        "elementType": "u32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batches", "type": "u32", "value": "dim(shapes.boxes, 0)" },
+            { "name": "boxesPerBatch", "type": "u32", "value": "dim(shapes.boxes, 1)" },
+            { "name": "classes", "type": "u32", "value": "dim(shapes.scores, 1)" },
+            { "name": "capacity", "type": "u32", "value": "dim(shapes.selected_indices, 0)" },
+            { "name": "centerPointBox", "type": "u32", "value": "attrs.center_point_box" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "groups_parallel",
+      "priority": 25,
+      "when": ["validLayout", "present.max_output_boxes_per_class", "present.iou_threshold", "present.score_threshold", "dim(shapes.selected_indices, 0) > 0", "dim(shapes.boxes, 1) > 0", "parallelDispatchFits", "parallelScratchFits"],
+      "intermediates": [
+        { "id": "suppressed", "dtype": "uint32", "shape": "[numGroups * dim(shapes.boxes, 1)]" },
+        { "id": "groupSel", "dtype": "uint32", "shape": "[numGroups * maxKeep]" },
+        { "id": "groupCount", "dtype": "uint32", "shape": "[numGroups]" }
+      ],
+      "passes": [
+        {
+          "id": "select",
+          "name": "NonMaxSuppression.GroupSelect",
+          "source": {
+            "shader": "nms-groups-select.wgsl.jinja",
+            "inputs": {
+              "maxKeep": "maxKeep",
+              "useSubgroups": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "parallelSelect",
+          "dispatch": { "x": "numGroups" }
+        },
+        {
+          "id": "compact",
+          "name": "NonMaxSuppression.GroupCompact",
+          "source": { "shader": "nms-groups-compact.wgsl.jinja", "inputs": { "maxKeep": "maxKeep" } },
+          "bindings": "parallelCompact",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_defaults",
+      "when": ["validLayout", "not present.max_output_boxes_per_class", "not present.iou_threshold", "not present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": false, "hasIouThreshold": false, "hasScoreThreshold": false }
+          },
+          "bindings": "serialDefaults",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_max",
+      "when": ["validLayout", "present.max_output_boxes_per_class", "not present.iou_threshold", "not present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": true, "hasIouThreshold": false, "hasScoreThreshold": false }
+          },
+          "bindings": "serialMax",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_iou",
+      "when": ["validLayout", "not present.max_output_boxes_per_class", "present.iou_threshold", "not present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": false, "hasIouThreshold": true, "hasScoreThreshold": false }
+          },
+          "bindings": "serialIou",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_score",
+      "when": ["validLayout", "not present.max_output_boxes_per_class", "not present.iou_threshold", "present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": false, "hasIouThreshold": false, "hasScoreThreshold": true }
+          },
+          "bindings": "serialScore",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_max_iou",
+      "when": ["validLayout", "present.max_output_boxes_per_class", "present.iou_threshold", "not present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": true, "hasIouThreshold": true, "hasScoreThreshold": false }
+          },
+          "bindings": "serialMaxIou",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_max_score",
+      "when": ["validLayout", "present.max_output_boxes_per_class", "not present.iou_threshold", "present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": true, "hasIouThreshold": false, "hasScoreThreshold": true }
+          },
+          "bindings": "serialMaxScore",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_iou_score",
+      "when": ["validLayout", "not present.max_output_boxes_per_class", "present.iou_threshold", "present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": false, "hasIouThreshold": true, "hasScoreThreshold": true }
+          },
+          "bindings": "serialIouScore",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    },
+    {
+      "id": "serial_all_inputs",
+      "when": ["validLayout", "present.max_output_boxes_per_class", "present.iou_threshold", "present.score_threshold"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "NonMaxSuppression",
+          "source": {
+            "shader": "non-max-suppression.wgsl.jinja",
+            "inputs": { "hasMaxOutput": true, "hasIouThreshold": true, "hasScoreThreshold": true }
+          },
+          "bindings": "serialAll",
+          "dispatch": { "x": 1 }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,20 @@

+{
+  "name": "ai.onnx.NonMaxSuppression",
+  "id": "_ai_onnx_nonmaxsuppression_webgpu_d44f6ed",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "QfNxa5siC+dAWXaBHejTokSn3gllrgjOelQ0UNQlxEE=",
+      "manifest.json": "vZuoGvM3nYg4zijLh8n62me7YOGuyRElAeZQelssago=",
+      "nms-groups-compact.wgsl.jinja": "4z2+LMPnLoFRYfgH4bMN9D9eQw/besrivDzeM81JI5c=",
+      "nms-groups-select.wgsl.jinja": "OKea1RdhUe0+cDFuFGIrx32rOxnRbUxgDJA+0FHIHt8=",
+      "non-max-suppression.wgsl.jinja": "D4xiNb1cSvy5W1TjKi3jp4zA/H6besnr0zfsh3OAk+Q=",
+      "test.json": "ehZ2AurQHFMhpSAmDim7zkCfdJL/IdB6VaCVCakwAMY="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.NonMaxSuppression" }
+}

build/webgpu/nms-groups-compact.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,49 @@

+// Pass 2 of group-parallel NonMaxSuppression. One workgroup walks per-group
+// selections in (batch, class)-major order and compacts them into the final
+// [num_selected, 3] tensor. The running offset is the sum of prior group counts,
+// whose exact data-dependent shape is supplied by the caller. Selections within
+// a group are already in score-descending greedy order. Defensive bounds and
+// initialization keep an invalid allocation from producing out-of-bounds writes.
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const MAX_KEEP: u32 = {{ source.maxKeep }}u;
+var<workgroup> wgTotal: u32;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
+  let tid = lid.x;
+  let numGroups = params.batches * params.classes;
+  if (tid == 0u) {
+    var total = 0u;
+    for (var group = 0u; group < numGroups; group = group + 1u) {
+      let batch = group / params.classes;
+      let class_id = group % params.classes;
+      let selBase = group * MAX_KEEP;
+      let n = groupCount[group];
+      for (var k = 0u; k < n; k = k + 1u) {
+        if (total < params.capacity) {
+          let box = groupSel[selBase + k];
+          let row = total * 3u;
+          selected_indices[row] = batch;
+          selected_indices[row + 1u] = class_id;
+          selected_indices[row + 2u] = box;
+        }
+        total = total + 1u;
+      }
+    }
+    wgTotal = total;
+  }
+  let total = workgroupUniformLoad(&wgTotal);
+  // Initialize any defensive tail rows to (0,0,0).
+  let written = min(total, params.capacity);
+  for (var row = written + tid; row < params.capacity; row = row + WG) {
+    let base = row * 3u;
+    selected_indices[base] = 0u;
+    selected_indices[base + 1u] = 0u;
+    selected_indices[base + 2u] = 0u;
+  }
+}

build/webgpu/nms-groups-select.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,223 @@

+// Pass 1 of group-parallel NonMaxSuppression. Independent (batch, class) groups
+// do not interact, so each gets one workgroup. Within a group, the greedy step
+// and each suppression sweep are parallelized across the workgroup. The sweep
+// also finds the next surviving score maximum. Ties choose the lower box index,
+// and the chosen corners are shared for the iteration.
+//
+// Per-group scratch stores sticky suppression flags, score-descending selected
+// indices, and the selection count. Every group reaches its natural limit;
+// pass 2 applies batch/class-major ordering into the exact output allocation.
+{% if source.useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const MAX_KEEP: u32 = {{ source.maxKeep }}u;
+const INVALID: u32 = 0xffffffffu;
+fn box_offset(batch: u32, box: u32) -> u32 {
+  return (batch * params.boxesPerBatch + box) * 4u;
+}
+fn score_at(batch: u32, class_id: u32, box: u32) -> f32 {
+  return scores[(batch * params.classes + class_id) * params.boxesPerBatch + box];
+}
+fn box_corners(batch: u32, box: u32) -> vec4<f32> {
+  let offset = box_offset(batch, box);
+  let v0 = boxes[offset];
+  let v1 = boxes[offset + 1u];
+  let v2 = boxes[offset + 2u];
+  let v3 = boxes[offset + 3u];
+  if (params.centerPointBox != 0u) {
+    let x_center = v0;
+    let y_center = v1;
+    let half_w = v2 * 0.5;
+    let half_h = v3 * 0.5;
+    return vec4<f32>(y_center - half_h, x_center - half_w, y_center + half_h, x_center + half_w);
+  }
+  return vec4<f32>(v0, v1, v2, v3);
+}
+fn iou_with_box(batch: u32, a: u32, b_box: vec4<f32>) -> f32 {
+  let a_box = box_corners(batch, a);
+  let ay1_raw = a_box.x;
+  let ax1_raw = a_box.y;
+  let ay2_raw = a_box.z;
+  let ax2_raw = a_box.w;
+  let by1_raw = b_box.x;
+  let bx1_raw = b_box.y;
+  let by2_raw = b_box.z;
+  let bx2_raw = b_box.w;
+  let ay1 = min(ay1_raw, ay2_raw);
+  let ax1 = min(ax1_raw, ax2_raw);
+  let ay2 = max(ay1_raw, ay2_raw);
+  let ax2 = max(ax1_raw, ax2_raw);
+  let by1 = min(by1_raw, by2_raw);
+  let bx1 = min(bx1_raw, bx2_raw);
+  let by2 = max(by1_raw, by2_raw);
+  let bx2 = max(bx1_raw, bx2_raw);
+  let inter_y1 = max(ay1, by1);
+  let inter_x1 = max(ax1, bx1);
+  let inter_y2 = min(ay2, by2);
+  let inter_x2 = min(ax2, bx2);
+  let inter = max(0.0, inter_y2 - inter_y1) * max(0.0, inter_x2 - inter_x1);
+  let area_a = max(0.0, ay2 - ay1) * max(0.0, ax2 - ax1);
+  let area_b = max(0.0, by2 - by1) * max(0.0, bx2 - bx1);
+  let denom = area_a + area_b - inter;
+  if (denom <= 0.0) { return 0.0; }
+  return inter / denom;
+}
+var<workgroup> wgScore: array<f32, WG>;
+var<workgroup> wgBox: array<u32, WG>;
+var<workgroup> wgChosen: u32;
+var<workgroup> wgChosenCorners: vec4<f32>;
+fn reduce_best(
+  tid: u32,
+{% if source.useSubgroups %}
+  sg_lane: u32,
+  sg_id: u32,
+  num_sg: u32,
+  sg_size: u32,
+{% endif %}
+  local_score: f32,
+  local_box: u32
+) {
+{% if source.useSubgroups %}
+  // Reduce within each hardware subgroup, then reduce the subgroup winners.
+  // Each lane folds a strided set of winners first, so this is independent of
+  // subgroup width and also covers devices where num_sg > sg_size.
+  let lane_score = select(-3.4028234663852886e38, local_score, local_box != INVALID);
+  let sg_score = subgroupMax(lane_score);
+  let lane_box = select(INVALID, local_box, local_box != INVALID && local_score == sg_score);
+  let sg_box = subgroupMin(lane_box);
+  if (sg_lane == 0u) {
+    wgScore[sg_id] = sg_score;
+    wgBox[sg_id] = sg_box;
+  }
+  workgroupBarrier();
+  var block_score = -3.4028234663852886e38;
+  var block_box = INVALID;
+  for (var i = sg_lane; i < num_sg; i += sg_size) {
+    let candidate_score = wgScore[i];
+    let candidate_box = wgBox[i];
+    if (candidate_box != INVALID &&
+        (block_box == INVALID || candidate_score > block_score ||
+         (candidate_score == block_score && candidate_box < block_box))) {
+      block_score = candidate_score;
+      block_box = candidate_box;
+    }
+  }
+  let final_score = subgroupMax(block_score);
+  let final_box = subgroupMin(select(INVALID, block_box, block_score == final_score));
+  if (tid == 0u) {
+    wgChosen = final_box;
+  }
+  workgroupBarrier();
+{% else %}
+  wgScore[tid] = local_score;
+  wgBox[tid] = local_box;
+  var stride = WG / 2u;
+  while (stride > 0u) {
+    workgroupBarrier();
+    if (tid < stride) {
+      let o_score = wgScore[tid + stride];
+      let o_box = wgBox[tid + stride];
+      let m_score = wgScore[tid];
+      let m_box = wgBox[tid];
+      if (o_box != INVALID && (m_box == INVALID || o_score > m_score || (o_score == m_score && o_box < m_box))) {
+        wgScore[tid] = o_score;
+        wgBox[tid] = o_box;
+      }
+    }
+    stride = stride >> 1u;
+  }
+  workgroupBarrier();
+  if (tid == 0u) {
+    wgChosen = wgBox[0];
+  }
+  workgroupBarrier();
+{% endif %}
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>{% if source.useSubgroups %},
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_id) sg_id: u32,
+  @builtin(num_subgroups) num_sg: u32,
+  @builtin(subgroup_size) sg_size: u32{% endif %}
+) {
+  let tid = lid.x;
+  let group = wid.x;
+  // group index decodes to (batch, class) in (batch, class)-major order, exactly
+  // matching the serial iteration order.
+  let batch = group / params.classes;
+  let class_id = group % params.classes;
+  let suppBase = group * params.boxesPerBatch;
+  let selBase = group * MAX_KEEP;
+  let max_per_class = min(max_output_boxes_per_class[0], MAX_KEEP);
+  let iou_threshold_value = iou_threshold[0];
+  let score_threshold_value = score_threshold[0];
+  var best_box = INVALID;
+  var best_score = 0.0;
+  for (var box = tid; box < params.boxesPerBatch; box = box + WG) {
+    suppressed[suppBase + box] = 0u;
+    let score = score_at(batch, class_id, box);
+    if (score <= score_threshold_value) { continue; }
+    if (best_box == INVALID || score > best_score || (score == best_score && box < best_box)) {
+      best_score = score;
+      best_box = box;
+    }
+  }
+  var selected_for_class = 0u;
+  reduce_best(tid, {% if source.useSubgroups %}sg_lane, sg_id, num_sg, sg_size, {% endif %}best_score, best_box);
+  var chosen = workgroupUniformLoad(&wgChosen);
+  while (chosen != INVALID && selected_for_class < max_per_class) {
+    if (tid == 0u) {
+      groupSel[selBase + selected_for_class] = chosen;
+    }
+    selected_for_class = selected_for_class + 1u;
+    // The last retained box cannot affect the requested output; avoid one
+    // complete IoU sweep and reduction when maxKeep is reached.
+    if (selected_for_class >= max_per_class) { break; }
+    if (tid == 0u) {
+      wgChosenCorners = box_corners(batch, chosen);
+    }
+    let chosen_corners = workgroupUniformLoad(&wgChosenCorners);
+    best_box = INVALID;
+    best_score = score_threshold_value;
+    for (var box = tid; box < params.boxesPerBatch; box = box + WG) {
+      if (box == chosen) {
+        suppressed[suppBase + box] = 1u;
+        continue;
+      }
+      if (suppressed[suppBase + box] != 0u) { continue; }
+      if (iou_with_box(batch, box, chosen_corners) > iou_threshold_value) {
+        suppressed[suppBase + box] = 1u;
+        continue;
+      }
+      let score = score_at(batch, class_id, box);
+      if (score <= score_threshold_value) { continue; }
+      if (best_box == INVALID || score > best_score || (score == best_score && box < best_box)) {
+        best_score = score;
+        best_box = box;
+      }
+    }
+    reduce_best(tid, {% if source.useSubgroups %}sg_lane, sg_id, num_sg, sg_size, {% endif %}best_score, best_box);
+    chosen = workgroupUniformLoad(&wgChosen);
+  }
+  if (tid == 0u) {
+    groupCount[group] = selected_for_class;
+  }
+}

build/webgpu/non-max-suppression.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,126 @@

+{{ env.wgsl.resourceDeclarations }}
+fn box_offset(batch: u32, box: u32) -> u32 {
+  return (batch * params.boxesPerBatch + box) * 4u;
+}
+fn score_at(batch: u32, class_id: u32, box: u32) -> f32 {
+  return scores[(batch * params.classes + class_id) * params.boxesPerBatch + box];
+}
+fn box_corners(batch: u32, box: u32) -> vec4<f32> {
+  let offset = box_offset(batch, box);
+  let v0 = boxes[offset];
+  let v1 = boxes[offset + 1u];
+  let v2 = boxes[offset + 2u];
+  let v3 = boxes[offset + 3u];
+  if (params.centerPointBox != 0u) {
+    let x_center = v0;
+    let y_center = v1;
+    let half_w = v2 * 0.5;
+    let half_h = v3 * 0.5;
+    return vec4<f32>(y_center - half_h, x_center - half_w, y_center + half_h, x_center + half_w);
+  }
+  return vec4<f32>(v0, v1, v2, v3);
+}
+fn iou_with_box(batch: u32, a: u32, b_box: vec4<f32>) -> f32 {
+  let a_box = box_corners(batch, a);
+  let ay1_raw = a_box.x;
+  let ax1_raw = a_box.y;
+  let ay2_raw = a_box.z;
+  let ax2_raw = a_box.w;
+  let by1_raw = b_box.x;
+  let bx1_raw = b_box.y;
+  let by2_raw = b_box.z;
+  let bx2_raw = b_box.w;
+  let ay1 = min(ay1_raw, ay2_raw);
+  let ax1 = min(ax1_raw, ax2_raw);
+  let ay2 = max(ay1_raw, ay2_raw);
+  let ax2 = max(ax1_raw, ax2_raw);
+  let by1 = min(by1_raw, by2_raw);
+  let bx1 = min(bx1_raw, bx2_raw);
+  let by2 = max(by1_raw, by2_raw);
+  let bx2 = max(bx1_raw, bx2_raw);
+  let inter_y1 = max(ay1, by1);
+  let inter_x1 = max(ax1, bx1);
+  let inter_y2 = min(ay2, by2);
+  let inter_x2 = min(ax2, bx2);
+  let inter = max(0.0, inter_y2 - inter_y1) * max(0.0, inter_x2 - inter_x1);
+  let area_a = max(0.0, ay2 - ay1) * max(0.0, ax2 - ax1);
+  let area_b = max(0.0, by2 - by1) * max(0.0, bx2 - bx1);
+  let denom = area_a + area_b - inter;
+  if (denom <= 0.0) { return 0.0; }
+  return inter / denom;
+}
+fn iou(batch: u32, a: u32, b: u32) -> f32 {
+  return iou_with_box(batch, a, box_corners(batch, b));
+}
+fn selected_box(row: u32) -> u32 {
+  return selected_indices[row * 3u + 2u];
+}
+fn write_selected(row: u32, batch: u32, class_id: u32, box: u32) {
+  let base = row * 3u;
+  selected_indices[base] = batch;
+  selected_indices[base + 1u] = class_id;
+  selected_indices[base + 2u] = box;
+}
+@compute @workgroup_size(1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  if (gid.x != 0u) { return; }
+{% if source.hasMaxOutput %}
+  let max_per_class = max_output_boxes_per_class[0];
+{% else %}
+  let max_per_class = 0u;
+{% endif %}
+{% if source.hasIouThreshold %}
+  let iou_threshold_value = iou_threshold[0];
+{% else %}
+  let iou_threshold_value = 0.0;
+{% endif %}
+  var total = 0u;
+  for (var batch = 0u; batch < params.batches; batch = batch + 1u) {
+    for (var class_id = 0u; class_id < params.classes; class_id = class_id + 1u) {
+      let class_start = total;
+      var selected_for_class = 0u;
+      loop {
+        if (selected_for_class >= max_per_class) { break; }
+        var best_box = 0xffffffffu;
+        var best_score = 0.0;
+        for (var box = 0u; box < params.boxesPerBatch; box = box + 1u) {
+          let score = score_at(batch, class_id, box);
+{% if source.hasScoreThreshold %}
+          if (score <= score_threshold[0]) { continue; }
+{% endif %}
+          var suppressed = false;
+          for (var prev = 0u; prev < selected_for_class; prev = prev + 1u) {
+            let sel = selected_box(class_start + prev);
+            // An already-selected box must never be re-selected. For normal boxes
+            // iou(box, box) == 1 > threshold handles this, but a zero-area box has
+            // iou 0/0 -> 0 (the denom<=0 guard), so it would be re-picked every
+            // iteration; skip it explicitly by index (no-op for normal boxes).
+            if (box == sel || iou(batch, box, sel) > iou_threshold_value) {
+              suppressed = true;
+            }
+          }
+          if (!suppressed && (best_box == 0xffffffffu || score > best_score || (score == best_score && box < best_box))) {
+            best_score = score;
+            best_box = box;
+          }
+        }
+        if (best_box == 0xffffffffu) { break; }
+        write_selected(total, batch, class_id, best_box);
+        total = total + 1u;
+        selected_for_class = selected_for_class + 1u;
+      }
+    }
+  }
+  for (var row = total; row < params.capacity; row = row + 1u) {
+    write_selected(row, 0u, 0u, 0u);
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1179 @@

+{
+  "op": "ai.onnx.NonMaxSuppression",
+  "fixtureArrays": {
+    "nms_six_boxes": [0, 0, 1, 1, 0, 0.1, 1, 1.1, 0, -0.1, 1, 0.9, 0, 10, 1, 11, 0, 10.1, 1, 11.1, 0, 100, 1, 101],
+    "identical_boxes_tie_lower_index_ort_input_boxes": [0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1],
+    "two_batches_one_class_ort_input_boxes": [0, 0, 1, 1, 0, 0.1, 1, 1.1, 0, -0.1, 1, 0.9, 0, 10, 1, 11, 0, 10.1, 1, 11.1, 0, 100, 1, 101, 0, 0, 1, 1, 0, 0.1, 1, 1.1, 0, -0.1, 1, 0.9, 0, 10, 1, 11, 0, 10.1, 1, 11.1, 0, 100, 1, 101],
+    "two_batches_two_classes_ort_input_boxes": [0, 0, 0.3, 0.3, 0, 0, 0.4, 0.4, 0, 0, 0.5, 0.5, 0.5, 0.5, 0.9, 0.9, 0.5, 0.5, 1, 1, 0, 0, 0.3, 0.3, 0, 0, 0.4, 0.4, 0.5, 0.5, 0.95, 0.95, 0.5, 0.5, 0.96, 0.96, 0.5, 0.5, 1, 1],
+    "two_batches_two_classes_ort_input_scores": [0.1, 0.2, 0.6, 0.3, 0.9, 0.1, 0.2, 0.6, 0.3, 0.9, 0.1, 0.2, 0.6, 0.3, 0.9, 0.1, 0.2, 0.6, 0.3, 0.9],
+    "flipped_coordinates_ort_input_boxes": [1, 1, 0, 0, 0, 0.1, 1, 1.1, 0, 0.9, 1, -0.1, 0, 10, 1, 11, 1, 10.1, 0, 11.1, 1, 101, 0, 100],
+    "onnx_backend_nonmaxsuppression_limit_output_size_input_boxes": [0, 0, 1, 1, 0, 0.10000000149011612, 1, 1.100000023841858, 0, -0.10000000149011612, 1, 0.8999999761581421, 0, 10, 1, 11, 0, 10.100000381469727, 1, 11.100000381469727, 0, 100, 1, 101]
+  },
+  "cases": [
+    {
+      "name": "single_class_suppresses_overlap",
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 10.0, 10.0, 1.0, 1.0, 11.0, 11.0, 20.0, 20.0, 30.0, 30.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [0.9, 0.8, 0.7] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } }
+    },
+    {
+      "name": "f32_subnormal_score_above_zero_threshold_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: Metal flushes denormals to zero in the ALU; the subnormal score is flushed when loaded for the > score_threshold comparison and the ordered best-score selection. Permanent FTZ limitation."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.WithScoreThreshold",
+        "notes": "Score threshold filtering is strict; a finite positive subnormal score is greater than a zero threshold and should produce one selected index."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0, 2.0, 2.0, 3.0, 3.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [0.0, 1e-40] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": {
+        "selected_indices": {
+          "dtype": "uint32",
+          "shape": [1, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 1] }
+        }
+      }
+    },
+    {
+      "name": "two_classes_threshold",
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0, 0.0, 0.0, 2.0, 2.0, 4.0, 4.0, 5.0, 5.0] }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [0.4, 0.9, 0.7, 0.6, 0.2, 0.8] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.3] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } }
+    },
+    {
+      "name": "ort_two_classes_full_capacity",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.TwoClasses",
+        "notes": "The fixture supplies the exact data-dependent ONNX output shape."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 2, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3, 0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [6] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [6, 3] } }
+    },
+    {
+      "name": "ort_with_score_threshold",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.WithScoreThreshold",
+        "notes": "The fixture supplies the exact data-dependent ONNX output shape."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.4] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } }
+    },
+    {
+      "name": "single_box_capacity_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.SelectSingleBox"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 1], "data": { "kind": "values", "values": [0.9] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [1, 3] } }
+    },
+    {
+      "name": "identical_boxes_tie_lower_index_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.SelectFromIdenticalBoxes"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 10, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/identical_boxes_tie_lower_index_ort_input_boxes" }
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 10],
+          "data": { "kind": "values", "values": [0.9, 0.9, 0.9, 0.9, 0.9, 0.9, 0.9, 0.9, 0.9, 0.9] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [1, 3] } }
+    },
+    {
+      "name": "select_thirty_capacity_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.SelectThirty"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [30] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3] } }
+    },
+    {
+      "name": "with_iou_threshold_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.WithIOUThreshold"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3] } }
+    },
+    {
+      "name": "without_score_threshold_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.WithoutScoreThreshold"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3] } }
+    },
+    {
+      "name": "select_two_capacity_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.SelectTwo"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } }
+    },
+    {
+      "name": "two_batches_one_class_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.TwoBatches_OneClass"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [2, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/two_batches_one_class_ort_input_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [2, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3, 0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [4, 3] } }
+    },
+    {
+      "name": "two_batches_two_classes_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.TwoBatches_TwoClasses"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [2, 5, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/two_batches_two_classes_ort_input_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [2, 2, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/two_batches_two_classes_ort_input_scores" } }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.8] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [8, 3] } }
+    },
+    {
+      "name": "score_threshold_zero_scores_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.WithScoreThresholdZeroScores"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.1, 0.0, 0.0, 0.3, 0.2, -5.0] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [6] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-3.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } }
+    },
+    {
+      "name": "flipped_coordinates_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.FlippedCoordinates"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/flipped_coordinates_ort_input_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3] } }
+    },
+    {
+      "name": "empty_input_capacity_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.EmptyInput"
+      },
+      "inputs": {
+        "boxes": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } },
+        "scores": { "dtype": "float32", "shape": [1, 1, 0], "data": { "kind": "values", "values": [] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [30] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [0, 3] } }
+    },
+    {
+      "name": "zero_max_output_per_class_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.ZeroMaxOutputPerClass"
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.4] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [0, 3] } }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_center_point_box_format",
+      "attrs": { "center_point_box": 1 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.5, 0.5, 1.0, 1.0, 0.5, 0.6000000238418579, 1.0, 1.0, 0.5, 0.4000000059604645, 1.0, 1.0, 0.5, 10.5, 1.0, 1.0, 0.5, 10.600000381469727, 1.0, 1.0, 0.5, 100.5, 1.0, 1.0]
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_center_point_box_format",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_flipped_coordinates",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 1.0, 0.0, 0.0, 0.0, 0.10000000149011612, 1.0, 1.100000023841858, 0.0, 0.8999999761581421, 1.0, -0.10000000149011612, 0.0, 10.0, 1.0, 11.0, 1.0, 10.100000381469727, 0.0, 11.100000381469727, 1.0, 101.0, 0.0, 100.0]
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_flipped_coordinates",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_identical_boxes",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 10, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/identical_boxes_tie_lower_index_ort_input_boxes" }
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 10],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.8999999761581421, 0.8999999761581421, 0.8999999761581421, 0.8999999761581421, 0.8999999761581421, 0.8999999761581421, 0.8999999761581421, 0.8999999761581421, 0.8999999761581421]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [1, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_identical_boxes",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_iou_threshold_boundary",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0, 0.5, 0.5, 1.5, 1.5] }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 2],
+          "data": { "kind": "values", "values": [0.8999999761581421, 0.800000011920929] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "values", "values": [0.1428571492433548] }
+        },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_iou_threshold_boundary",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_limit_output_size",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_nonmaxsuppression_limit_output_size_input_boxes" }
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_limit_output_size",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_single_box",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0] }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 1],
+          "data": { "kind": "values", "values": [0.8999999761581421] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [1, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_single_box",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_suppress_by_IOU",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_nonmaxsuppression_limit_output_size_input_boxes" }
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_suppress_by_IOU",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_suppress_by_IOU_and_scores",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_nonmaxsuppression_limit_output_size_input_boxes" }
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "values", "values": [0.4000000059604645] }
+        }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_suppress_by_IOU_and_scores",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_two_batches",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [2, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.0, 1.0, 1.0, 0.0, 0.10000000149011612, 1.0, 1.100000023841858, 0.0, -0.10000000149011612, 1.0, 0.8999999761581421, 0.0, 10.0, 1.0, 11.0, 0.0, 10.100000381469727, 1.0, 11.100000381469727, 0.0, 100.0, 1.0, 101.0, 0.0, 0.0, 1.0, 1.0, 0.0, 0.10000000149011612, 1.0, 1.100000023841858, 0.0, -0.10000000149011612, 1.0, 0.8999999761581421, 0.0, 10.0, 1.0, 11.0, 0.0, 10.100000381469727, 1.0, 11.100000381469727, 0.0, 100.0, 1.0, 101.0]
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [2, 1, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896, 0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [4, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_two_batches",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_nonmaxsuppression_two_classes",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_nonmaxsuppression_limit_output_size_input_boxes" }
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 2, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896, 0.8999999761581421, 0.75, 0.6000000238418579, 0.949999988079071, 0.5, 0.30000001192092896]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [4, 3] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_nonmaxsuppression_two_classes",
+        "notes": "The ONNX optional scalar inputs are supplied as rank-0 tensors, the output uses its exact data-dependent shape, and ONNX int64 indices are represented as uint32 where representable."
+      }
+    },
+    {
+      "name": "ort_with_iou_threshold",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.WithIOUThreshold"
+      },
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_center_point_box_format",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.CenterPointBoxFormat"
+      },
+      "attrs": { "center_point_box": 1 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.5, 0.5, 1.0, 1.0, 0.5, 0.6, 1.0, 1.0, 0.5, 0.4, 1.0, 1.0, 0.5, 10.5, 1.0, 1.0, 0.5, 10.6, 1.0, 1.0, 0.5, 100.5, 1.0, 1.0]
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_two_batches_one_class",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.TwoBatches_OneClass"
+      },
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [2, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/two_batches_one_class_ort_input_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [2, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3, 0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [4, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_two_batches_two_classes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.TwoBatches_TwoClasses"
+      },
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [2, 5, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/two_batches_two_classes_ort_input_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [2, 2, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/two_batches_two_classes_ort_input_scores" } }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.8] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [8, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_score_threshold_zero_scores",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.WithScoreThresholdZeroScores"
+      },
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.1, 0.0, 0.0, 0.3, 0.2, -5.0] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [6] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-3.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_flipped_coordinates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.FlippedCoordinates"
+      },
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/flipped_coordinates_ort_input_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_zero_max_output_per_class",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.ZeroMaxOutputPerClass"
+      },
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.4] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "groups_parallel_single_class_512",
+      "provenance": {
+        "notes": "512 boxes on an integer grid (13-value cycle gives 13 repeating geometries -> heavy score/index tie-breaking). Grid coordinates keep IoU ratios far from the threshold so float comparisons are unambiguous. Exercises the group-parallel fused suppression/next-argmax path."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 512, 4],
+          "data": { "kind": "cycle", "values": [0.0, 0.0, 4.0, 4.0, 1.0, 1.0, 5.0, 5.0, 2.0, 2.0, 6.0, 6.0, 3.0] }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 512],
+          "data": {
+            "kind": "cycle",
+            "values": [0.9, 0.1, 0.55, 0.7, 0.3, 0.85, 0.2, 0.65, 0.4, 0.75, 0.05, 0.95, 0.5, 0.6, 0.15, 0.8, 0.35]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [16] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [4, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "groups_parallel_two_batches_two_classes",
+      "provenance": {
+        "notes": "2 batches x 2 classes x 320 boxes; 19-value box cycle yields varied grid geometries. Per-group suppression state must remain isolated across batches and classes."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [2, 320, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [0.0, 0.0, 2.0, 2.0, 1.0, 0.0, 3.0, 2.0, 2.0, 0.0, 4.0, 2.0, 0.0, 1.0, 2.0, 3.0, 1.0, 1.0, 3.0]
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [2, 2, 320],
+          "data": {
+            "kind": "cycle",
+            "values": [0.31, 0.97, 0.12, 0.58, 0.74, 0.05, 0.66, 0.83, 0.22, 0.49, 0.91, 0.38, 0.7, 0.15, 0.62, 0.88, 0.27, 0.53, 0.79, 0.08, 0.44, 0.95, 0.35]
+          }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [8] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.4] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.3] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [30, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "groups_parallel_center_point_box_256",
+      "provenance": {
+        "notes": "center_point_box=1 decode through the group-parallel kernel; 17-value positive cycle keeps widths/heights valid and corners on a half-integer grid."
+      },
+      "attrs": { "center_point_box": 1 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 256, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [2.0, 2.0, 2.0, 2.0, 5.0, 3.0, 4.0, 2.0, 3.0, 6.0, 2.0, 4.0, 7.0, 7.0, 6.0, 6.0, 1.0]
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "cycle", "values": [0.45, 0.9, 0.1, 0.7, 0.3, 0.85, 0.55, 0.25, 0.65, 0.4, 0.75] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.2] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [10, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_big_int_max_output_boxes_per_class",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/object_detection/non_max_suppression_test.cc",
+        "test": "NonMaxSuppressionOpTest.BigIntMaxOutputBoxesPerClass",
+        "notes": "ORT source coverage uses INT64_MAX; this C-long-safe large positive scalar has the same no-practical-limit behavior while the fixture supplies the exact result shape."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/nms_six_boxes" } }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 1, 6],
+          "data": { "kind": "values", "values": [0.9, 0.75, 0.6, 0.95, 0.5, 0.3] }
+        },
+        "max_output_boxes_per_class": {
+          "dtype": "uint32",
+          "shape": [],
+          "data": { "kind": "values", "values": [2147483647] }
+        },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.4] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "score_exactly_at_threshold_excluded",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0, 5.0, 5.0, 6.0, 6.0, 10.0, 10.0, 11.0, 11.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [0.8, 0.5, 0.3] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": {
+        "selected_indices": {
+          "dtype": "uint32",
+          "shape": [1, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "iou_exactly_at_threshold_kept",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 5.0, 4.0, 3.0, 0.0, 8.0, 4.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [0.9, 0.8] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": {
+        "selected_indices": {
+          "dtype": "uint32",
+          "shape": [2, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 1] }
+        }
+      }
+    },
+    {
+      "name": "groups_parallel_multi_class_suppression_isolation",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.0, 2.0, 2.0, 0.5, 0.5, 2.5, 2.5, 10.0, 10.0, 12.0, 12.0, 20.0, 20.0, 22.0, 22.0]
+          }
+        },
+        "scores": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.9, 0.8, 0.7, 0.6, 0.2, 0.95, 0.85, 0.15] }
+        },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.1] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [6, 3] } }
+    },
+    {
+      "name": "zero_area_boxes_never_suppress_each_other",
+      "attrs": { "center_point_box": 0 },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [5.0, 5.0, 5.0, 5.0, 5.0, 5.0, 5.0, 5.0, 3.0, 3.0, 3.0, 3.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [0.9, 0.8, 0.7] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } },
+      "provenance": {
+        "notes": "Reference-generated regression for zero-area boxes: explicit selected-box tracking prevents a zero self-IoU from selecting the same box repeatedly."
+      }
+    },
+    {
+      "name": "all_optional_inputs_omitted_defaults_to_empty",
+      "provenance": {
+        "source": "ONNX NonMaxSuppression-11",
+        "notes": "Covers all three optional scalar inputs omitted: max_output_boxes_per_class defaults to zero, so the exact result is empty even though scores include negative and positive values."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0, 2.0, 2.0, 3.0, 3.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [-1.0, 1.0] } }
+      },
+      "outputs": {
+        "selected_indices": {
+          "dtype": "uint32",
+          "shape": [0, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [] }
+        }
+      }
+    },
+    {
+      "name": "only_max_output_input_present",
+      "provenance": {
+        "source": "ONNX NonMaxSuppression-11",
+        "notes": "With score_threshold omitted, negative scores remain eligible; with iou_threshold omitted, its zero default keeps disjoint boxes."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0, 2.0, 2.0, 3.0, 3.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [-1.0, 1.0] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "selected_indices": {
+          "dtype": "uint32",
+          "shape": [2, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 1, 0, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "only_iou_threshold_input_present",
+      "provenance": {
+        "source": "ONNX NonMaxSuppression-11",
+        "notes": "An IOU input does not change the omitted max_output_boxes_per_class default of zero."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 1], "data": { "kind": "values", "values": [1.0] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "only_score_threshold_input_present",
+      "provenance": {
+        "source": "ONNX NonMaxSuppression-11",
+        "notes": "A score input does not change the omitted max_output_boxes_per_class default of zero."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 1], "data": { "kind": "values", "values": [1.0] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "max_and_score_inputs_iou_default",
+      "provenance": {
+        "source": "ONNX NonMaxSuppression-11",
+        "notes": "Covers the optional-input hole before score_threshold: iou_threshold is omitted and defaults to zero."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0, 2.0, 2.0, 3.0, 3.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [-1.0, 1.0] } },
+        "max_output_boxes_per_class": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": {
+        "selected_indices": {
+          "dtype": "uint32",
+          "shape": [1, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 1] }
+        }
+      }
+    },
+    {
+      "name": "iou_and_score_inputs_max_default",
+      "provenance": {
+        "source": "ONNX NonMaxSuppression-11",
+        "notes": "Covers both threshold inputs with max_output_boxes_per_class omitted; its zero default still produces an empty result."
+      },
+      "inputs": {
+        "boxes": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1.0, 1.0] }
+        },
+        "scores": { "dtype": "float32", "shape": [1, 1, 1], "data": { "kind": "values", "values": [1.0] } },
+        "iou_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "score_threshold": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "selected_indices": { "dtype": "uint32", "shape": [0, 3], "tolerance": 0 } }
+    }
+  ]
+}