Xenova HF Staff commited on 4 days ago

Commit

63dd0ba

verified ·

1 Parent(s): 1f2c04e

sync 91d990483a17

Browse files

Files changed (29) hide show

README.md +72 -14
build/webgpu/bench.json +1214 -5
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +104 -18
build/webgpu/scan-block-prefix-u32.wgsl.jinja +44 -0
build/webgpu/test.json +0 -0
build/webgpu/unique-axis-compact-sort.wgsl.jinja +31 -21
build/webgpu/unique-axis-dedup.wgsl.jinja +22 -20
build/webgpu/unique-axis-hash.wgsl.jinja +32 -7
build/webgpu/unique-axis-scalar-inverse.wgsl.jinja +109 -0
build/webgpu/unique-axis-scalar-ranks.wgsl.jinja +7 -0
build/webgpu/unique-axis-scatter.wgsl.jinja +7 -10
build/webgpu/unique-axis.wgsl.jinja +38 -54
build/webgpu/unique-compact-sort.wgsl.jinja +82 -26
build/webgpu/unique-flag-block-scan.wgsl.jinja +20 -0
build/webgpu/unique-flat-metadata.wgsl.jinja +129 -0
build/webgpu/unique-global-sort-exchange.wgsl.jinja +36 -0
build/webgpu/unique-global-sort-output.wgsl.jinja +9 -0
build/webgpu/unique-global-sort-shared.wgsl.jinja +69 -0
build/webgpu/unique-hash-build.wgsl.jinja +49 -4
build/webgpu/unique-hash-collect.wgsl.jinja +7 -7
build/webgpu/unique-hash-init.wgsl.jinja +2 -2
build/webgpu/unique-hash-mark.wgsl.jinja +52 -3
build/webgpu/unique-hash-sort-collected-key-only.wgsl.jinja +2 -2
build/webgpu/unique-scalar-compact.wgsl.jinja +60 -0
build/webgpu/unique-scalar-metadata.wgsl.jinja +122 -0
build/webgpu/unique-scalar-output.wgsl.jinja +20 -0
build/webgpu/unique-single-class.wgsl.jinja +32 -0
build/webgpu/unique.wgsl.jinja +14 -8

README.md CHANGED Viewed

@@ -12,24 +12,24 @@ tags:
 ## Description
-Finds unique values or subtensors along an optional `axis`. Without an axis, `X` is flattened; results are sorted or retain first-occurrence order. Sub-32-bit integers and booleans use lossless widened 32-bit storage. Metadata outputs remain logical int64 but use lossless uint32 storage because all values are bounded by an addressable tensor extent. Callers supply exact data-dependent output shapes. ONNX-permitted uint16, 64-bit, string, and complex inputs remain unsupported because the runtime lacks matching WebGPU storage.
 See the [ONNX `Unique` spec](https://onnx.ai/onnx/operators/onnx__Unique.html) for the reference semantics.
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `X` | `x` | `T` | — | — | The N-D input tensor from which unique values or subtensors are extracted. When `axis` is omitted, tensors of any rank are flattened in row-major order. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | derived | — | Tensor containing all unique values or subtensors of X, sorted or in first-occurrence order. | required |
-| `indices` | `indices` | `I` | `uint32` | `1` | — | Optional logical int64 indices of each `Y` value or slice's first occurrence in `X`; stored as bounded uint32 values by WebGPU. | optional |
-| `inverse_indices` | `inverse_indices` | `I` | `uint32` | `1` | — | Optional logical int64 mapping from each flattened input value, or each input-axis slice, to its corresponding index in `Y`; stored as bounded uint32 values by WebGPU. | optional |
-| `counts` | `counts` | `I` | `uint32` | `1` | — | Optional logical int64 occurrence count for each unique value or slice in `Y`; stored as bounded uint32 values by WebGPU. | optional |
 ## Attributes
@@ -37,8 +37,8 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
-| `sorted` | `1` | Whether to sort unique elements in ascending order before output; 1 (default) sorts, 0 retains first-occurrence order. |
 | `axis` | — | Optional axis along which unique subtensors are identified. Negative values count from the back; when omitted, the input is flattened. |
 ## Type constraints
@@ -47,37 +47,95 @@ Attributes and default values (overridable per request):
 | `T` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool` |
 | `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
 - [`unique-axis-compact-sort.wgsl.jinja`](build/webgpu/unique-axis-compact-sort.wgsl.jinja)
 - [`unique-axis-dedup.wgsl.jinja`](build/webgpu/unique-axis-dedup.wgsl.jinja)
 - [`unique-axis-hash.wgsl.jinja`](build/webgpu/unique-axis-hash.wgsl.jinja)
 - [`unique-axis-scatter.wgsl.jinja`](build/webgpu/unique-axis-scatter.wgsl.jinja)
 - [`unique-axis.wgsl.jinja`](build/webgpu/unique-axis.wgsl.jinja)
 - [`unique-compact-sort.wgsl.jinja`](build/webgpu/unique-compact-sort.wgsl.jinja)
 - [`unique-dedup.wgsl.jinja`](build/webgpu/unique-dedup.wgsl.jinja)
 - [`unique-hash-build.wgsl.jinja`](build/webgpu/unique-hash-build.wgsl.jinja)
 - [`unique-hash-collect.wgsl.jinja`](build/webgpu/unique-hash-collect.wgsl.jinja)
 - [`unique-hash-init.wgsl.jinja`](build/webgpu/unique-hash-init.wgsl.jinja)
 - [`unique-hash-mark.wgsl.jinja`](build/webgpu/unique-hash-mark.wgsl.jinja)
 - [`unique-hash-sort-collected-key-only.wgsl.jinja`](build/webgpu/unique-hash-sort-collected-key-only.wgsl.jinja)
 - [`unique.wgsl.jinja`](build/webgpu/unique.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:
-- `y`
-Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Description
+Finds unique values or subtensors along an optional `axis`. Without an axis, `X` is flattened; results are sorted or retain first-occurrence order. Sub-32-bit integers and booleans use lossless widened 32-bit storage. Metadata outputs remain logical int64 but use lossless uint32 storage because all values are bounded by an addressable tensor extent. Exact data-dependent output shapes must be supplied. ONNX-permitted uint16, 64-bit, string, and complex inputs are unsupported by this package.
 See the [ONNX `Unique` spec](https://onnx.ai/onnx/operators/onnx__Unique.html) for the reference semantics.
 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `x` | `X` | `T` | — | — | The N-D input tensor from which unique values or subtensors are extracted. When `axis` is omitted, tensors of any rank are flattened in row-major order. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | derived | — | Tensor containing all unique values or subtensors of X, sorted or in first-occurrence order. | required |
+| `indices` | — | `I` | `uint32` | `1` | — | Optional logical int64 indices of each `Y` value or slice's first occurrence in `X`; stored as bounded uint32 values by WebGPU. | optional |
+| `inverse_indices` | — | `I` | `uint32` | `1` | — | Optional logical int64 mapping from each flattened input value, or each input-axis slice, to its corresponding index in `Y`; stored as bounded uint32 values by WebGPU. | optional |
+| `counts` | — | `I` | `uint32` | `1` | — | Optional logical int64 occurrence count for each unique value or slice in `Y`; stored as bounded uint32 values by WebGPU. | optional |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `axis` | — | Optional axis along which unique subtensors are identified. Negative values count from the back; when omitted, the input is flattened. |
+| `sorted` | `1` | Whether to sort unique elements in ascending order before output; 1 (default) sorts, 0 retains first-occurrence order. |
 ## Type constraints
 | `T` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool` |
 | `I` | `int64` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `single_class_y` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_y_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_y_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `single_class_indices` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_indices_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_indices_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `single_class_inverse` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_inverse_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_inverse_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `single_class_indices_inverse` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_indices_inverse_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_indices_inverse_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `single_class_counts` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_counts_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_counts_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `single_class_indices_counts` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_indices_counts_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_indices_counts_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `single_class_inverse_counts` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_inverse_counts_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_inverse_counts_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `single_class_indices_inverse_counts` — Copy the first representative and materialize requested metadata when the exact output contract proves there is one distinct class.
+- `scalar_hash_parallel_indices_inverse_counts_unsorted` — Hash first occurrences, scan flag blocks in parallel, and scatter scalar representatives in input order with only the requested metadata.
+- `scalar_hash_parallel_indices_inverse_counts_sorted` — Hash and compact scalar representatives in parallel, sort their order with shared and global bitonic stages, and materialize only the requested metadata.
+- `flat_parallel_metadata_indices` — Builds the unique values with the parallel dedup and compaction passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_parallel_metadata_inverse` — Builds the unique values with the parallel dedup and compaction passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_parallel_metadata_indices_inverse` — Builds the unique values with the parallel dedup and compaction passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_parallel_metadata_counts` — Builds the unique values with the parallel dedup and compaction passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_parallel_metadata_indices_counts` — Builds the unique values with the parallel dedup and compaction passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_parallel_metadata_inverse_counts` — Builds the unique values with the parallel dedup and compaction passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_parallel_metadata_all` — Builds the unique values with the parallel dedup and compaction passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_hash_metadata_indices` — Builds the unique values with the hash-set dedup passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_hash_metadata_inverse` — Builds the unique values with the hash-set dedup passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_hash_metadata_indices_inverse` — Builds the unique values with the hash-set dedup passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_hash_metadata_counts` — Builds the unique values with the hash-set dedup passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_hash_metadata_indices_counts` — Builds the unique values with the hash-set dedup passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_hash_metadata_inverse_counts` — Builds the unique values with the hash-set dedup passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
+- `flat_hash_metadata_all` — Builds the unique values with the hash-set dedup passes, then resolves every metadata output from the finished result: one thread per input element for the bucket and the first-occurrence store, one thread per unique value for the tally.
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`scan-block-prefix-u32.wgsl.jinja`](build/webgpu/scan-block-prefix-u32.wgsl.jinja)
 - [`unique-axis-compact-sort.wgsl.jinja`](build/webgpu/unique-axis-compact-sort.wgsl.jinja)
 - [`unique-axis-dedup.wgsl.jinja`](build/webgpu/unique-axis-dedup.wgsl.jinja)
 - [`unique-axis-hash.wgsl.jinja`](build/webgpu/unique-axis-hash.wgsl.jinja)
+- [`unique-axis-scalar-inverse.wgsl.jinja`](build/webgpu/unique-axis-scalar-inverse.wgsl.jinja)
+- [`unique-axis-scalar-ranks.wgsl.jinja`](build/webgpu/unique-axis-scalar-ranks.wgsl.jinja)
 - [`unique-axis-scatter.wgsl.jinja`](build/webgpu/unique-axis-scatter.wgsl.jinja)
 - [`unique-axis.wgsl.jinja`](build/webgpu/unique-axis.wgsl.jinja)
 - [`unique-compact-sort.wgsl.jinja`](build/webgpu/unique-compact-sort.wgsl.jinja)
 - [`unique-dedup.wgsl.jinja`](build/webgpu/unique-dedup.wgsl.jinja)
+- [`unique-flag-block-scan.wgsl.jinja`](build/webgpu/unique-flag-block-scan.wgsl.jinja)
+- [`unique-flat-metadata.wgsl.jinja`](build/webgpu/unique-flat-metadata.wgsl.jinja)
+- [`unique-global-sort-exchange.wgsl.jinja`](build/webgpu/unique-global-sort-exchange.wgsl.jinja)
+- [`unique-global-sort-output.wgsl.jinja`](build/webgpu/unique-global-sort-output.wgsl.jinja)
+- [`unique-global-sort-shared.wgsl.jinja`](build/webgpu/unique-global-sort-shared.wgsl.jinja)
 - [`unique-hash-build.wgsl.jinja`](build/webgpu/unique-hash-build.wgsl.jinja)
 - [`unique-hash-collect.wgsl.jinja`](build/webgpu/unique-hash-collect.wgsl.jinja)
 - [`unique-hash-init.wgsl.jinja`](build/webgpu/unique-hash-init.wgsl.jinja)
 - [`unique-hash-mark.wgsl.jinja`](build/webgpu/unique-hash-mark.wgsl.jinja)
 - [`unique-hash-sort-collected-key-only.wgsl.jinja`](build/webgpu/unique-hash-sort-collected-key-only.wgsl.jinja)
+- [`unique-scalar-compact.wgsl.jinja`](build/webgpu/unique-scalar-compact.wgsl.jinja)
+- [`unique-scalar-metadata.wgsl.jinja`](build/webgpu/unique-scalar-metadata.wgsl.jinja)
+- [`unique-scalar-output.wgsl.jinja`](build/webgpu/unique-scalar-output.wgsl.jinja)
+- [`unique-single-class.wgsl.jinja`](build/webgpu/unique-single-class.wgsl.jinja)
 - [`unique.wgsl.jinja`](build/webgpu/unique.wgsl.jinja)
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
+This example supplies explicit metadata for:
+- `y`
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Unique",
   "cases": [
     {
       "name": "int32-64k-moderate-distinct-sorted",
@@ -116,9 +115,12 @@
       "preset": "smoke",
       "attrs": { "axis": 0, "sorted": 1 },
       "inputs": {
-        "x": { "dtype": "float32", "shape": [4096, 256], "dist": "uniform", "seed": 1531, "scale": 8, "offset": 0 }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [4096, 256] } },
       "bench": {
         "primary": true,
         "metrics": [
@@ -130,6 +132,27 @@
         ]
       }
     },
     {
       "name": "axis_hash_split_scatter_unsorted_zero_fills_tail",
       "preset": "smoke",
@@ -206,7 +229,7 @@
       "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + dim(shapes.y, 0))" }] }
     },
     {
-      "name": "axis0-f32-70000x1-above-hash-capacity-serial",
       "preset": "stress",
       "attrs": { "axis": 0, "sorted": 1 },
       "inputs": {
@@ -217,7 +240,10 @@
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [70000, 1] } },
-      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + dim(shapes.y, 0))" }] }
     },
     {
       "name": "int32-262144-Y4096-token-vocab-sorted",
@@ -242,6 +268,1189 @@
       "inputs": { "x": { "dtype": "int32", "shape": [8192], "dist": "linearMod", "mod": 8192 } },
       "outputs": { "y": { "dtype": "int32", "shape": [8192] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
     }
-  ]
 }

 {
   "cases": [
     {
       "name": "int32-64k-moderate-distinct-sorted",
       "preset": "smoke",
       "attrs": { "axis": 0, "sorted": 1 },
       "inputs": {
+        "x": { "dtype": "float32", "shape": [4096, 256], "data": { "kind": "linspace", "start": 8.0, "end": -8.0 } }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [4096, 256] } },
+      "provenance": {
+        "notes": "Descending values give 4096 distinct 256-element rows. Lexicographic sorting reverses row order while retaining each row's column order."
+      },
       "bench": {
         "primary": true,
         "metrics": [
         ]
       }
     },
+    {
+      "name": "axis0-f32-4096x256-1024-distinct-random-rows",
+      "preset": "smoke",
+      "attrs": { "axis": 0, "sorted": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4096, 256], "dist": "uniform", "seed": 1531, "scale": 8, "offset": 0 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024, 256] } },
+      "provenance": {
+        "notes": "Input contains four copies of each of 1024 distinct 256-element rows. Output contains the distinct rows in lexicographic order."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      }
+    },
     {
       "name": "axis_hash_split_scatter_unsorted_zero_fills_tail",
       "preset": "smoke",
       "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + dim(shapes.y, 0))" }] }
     },
     {
+      "name": "axis0-f32-70000x1-above-old-hash-capacity",
       "preset": "stress",
       "attrs": { "axis": 0, "sorted": 1 },
       "inputs": {
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [70000, 1] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + dim(shapes.y, 0))" }] },
+      "provenance": {
+        "notes": "A 70,000-row float tensor with every row distinct exercises large-axis hash deduplication and sorted emission."
+      }
     },
     {
       "name": "int32-262144-Y4096-token-vocab-sorted",
       "inputs": { "x": { "dtype": "int32", "shape": [8192], "dist": "linearMod", "mod": 8192 } },
       "outputs": { "y": { "dtype": "int32", "shape": [8192] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "int32-4096-128-distinct-inverse-metadata",
+      "provenance": {
+        "notes": "A 4,096-token stream with 128 distinct values exercises parallel local deduplication while producing an inverse map for vocabulary compaction."
+      },
+      "preset": "smoke",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [4096], "dist": "linearMod", "mod": 128 } },
+      "outputs": {
+        "y": { "dtype": "int32", "shape": [128] },
+        "inverse_indices": { "dtype": "uint32", "shape": [4096] }
+      },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "12 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "int32-65536-512-distinct-inverse-metadata-hash",
+      "provenance": {
+        "notes": "A 65,536-token stream with 512 distinct values exercises hash-set deduplication while producing an inverse map for vocabulary compaction."
+      },
+      "preset": "smoke",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [65536], "dist": "linearMod", "mod": 512 } },
+      "outputs": {
+        "y": { "dtype": "int32", "shape": [512] },
+        "inverse_indices": { "dtype": "uint32", "shape": [65536] }
+      },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "12 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "axis0-int32-512x16-record-table-shared-markers",
+      "provenance": {
+        "notes": "A fixed-width record table whose shared marker fields sit at the four positions an evenly-spaced four-sample prefilter reads, and whose identifier does not. Every pair of rows therefore agrees on the prefilter and only the exact comparator separates them."
+      },
+      "attrs": { "axis": 0, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [512, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/axis_record_table_marker_collision" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [512, 16], "dist": "empty" } }
+    },
+    {
+      "name": "flat-i32-70000-distinct-above-old-capacity",
+      "preset": "stress",
+      "provenance": {
+        "notes": "A flat 70,000-element float tensor with every value distinct exercises large-output hash deduplication and sorted emission."
+      },
+      "attrs": { "sorted": 1 },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [70000], "data": { "kind": "linspace", "start": -35000, "end": 34999 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [70000], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "axis0-f32-16384x1-inverse-metadata",
+      "preset": "stress",
+      "provenance": {
+        "notes": "A float axis request with an inverse map and more distinct rows than the ordered-map head cache holds, so the comparator reads its keys back from the order buffer instead."
+      },
+      "attrs": { "axis": 0, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [16384, 1],
+          "data": { "kind": "linspace", "start": -8192.0, "end": 8191.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [16384, 1], "dist": "empty" },
+        "inverse_indices": { "dtype": "uint32", "shape": [16384], "dist": "empty" }
+      }
+    },
+    {
+      "name": "flat-i32-1049600-1024-classes-label-raster",
+      "preset": "stress",
+      "provenance": {
+        "notes": "A 1,049,600-element class-label raster with 1,024 classes exercises large-input flat deduplication in first-occurrence order."
+      },
+      "attrs": { "sorted": 0 },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [1025, 1024], "data": { "kind": "linspace", "start": 0, "end": 1023 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1024], "dist": "empty" } }
+    },
+    {
+      "name": "flat-i32-70000-distinct-unsorted-first-order",
+      "preset": "stress",
+      "provenance": {
+        "source": "synthetic",
+        "notes": "Descending distinct integers measure large-output hash deduplication while preserving first-occurrence order for sorted=0."
+      },
+      "attrs": { "sorted": 0 },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [70000], "data": { "kind": "linspace", "start": 34999, "end": -35000 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [70000], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "axis-scalar-f32-inverse-descending",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Scalar-axis inverse map: descending. Checks the hash/rank route with non-identity ordering, duplicate-rich or history-dependent NaN inputs."
+      },
+      "attrs": { "axis": 0, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [16384, 1],
+          "data": { "kind": "linspace", "start": 8191.0, "end": -8192.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [16384, 1], "dist": "empty" },
+        "inverse_indices": { "dtype": "uint32", "shape": [16384], "dist": "empty" }
+      }
+    },
+    {
+      "name": "axis-scalar-f32-inverse-duplicate64",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Scalar-axis inverse map: duplicate64. Checks the hash/rank route with non-identity ordering, duplicate-rich or history-dependent NaN inputs."
+      },
+      "attrs": { "axis": 0, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4096, 1],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/axis-scalar-f32-inverse-duplicate64_input_x" }
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [64, 1], "dist": "empty" },
+        "inverse_indices": { "dtype": "uint32", "shape": [4096], "dist": "empty" }
+      }
+    },
+    {
+      "name": "axis-scalar-f32-inverse-duplicate64-unsorted",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Scalar-axis inverse map: duplicate64-unsorted. Checks the hash/rank route with non-identity ordering, duplicate-rich or history-dependent NaN inputs."
+      },
+      "attrs": { "axis": 0, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4096, 1],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/axis-scalar-f32-inverse-duplicate64_input_x" }
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [64, 1], "dist": "empty" },
+        "inverse_indices": { "dtype": "uint32", "shape": [4096], "dist": "empty" }
+      }
+    },
+    {
+      "name": "axis-scalar-f32-inverse-nan-prefix",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Scalar-axis inverse map: nan-prefix. Checks the hash/rank route with non-identity ordering, duplicate-rich or history-dependent NaN inputs."
+      },
+      "attrs": { "axis": 0, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8192, 1],
+          "data": { "kind": "cycle", "values": [5.0, "NaN", 1.0, "NaN", -2.0, "NaN", 3.0, 5.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [4, 1], "dist": "empty" },
+        "inverse_indices": { "dtype": "uint32", "shape": [8192], "dist": "empty" }
+      }
+    },
+    {
+      "name": "axis-scalar-f32-inverse-nan-prefix-unsorted",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Scalar-axis inverse map: nan-prefix-unsorted. Checks the hash/rank route with non-identity ordering, duplicate-rich or history-dependent NaN inputs."
+      },
+      "attrs": { "axis": 0, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8192, 1],
+          "data": { "kind": "cycle", "values": [5.0, "NaN", 1.0, "NaN", -2.0, "NaN", 3.0, 5.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [4, 1], "dist": "empty" },
+        "inverse_indices": { "dtype": "uint32", "shape": [8192], "dist": "empty" }
+      }
+    },
+    {
+      "name": "axis-scalar-f32-inverse-hybrid-floor",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Scalar-axis inverse map: hybrid-floor. Checks the hash/rank route with non-identity ordering, duplicate-rich or history-dependent NaN inputs."
+      },
+      "attrs": { "axis": 0, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4096, 1],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/axis-scalar-f32-inverse-hybrid-floor_input_x" }
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2049, 1], "dist": "empty" },
+        "inverse_indices": { "dtype": "uint32", "shape": [4096], "dist": "empty" }
+      }
+    },
+    {
+      "name": "r4_scalar_scan_n32767_u13_sorted0",
+      "attrs": { "axis": -2, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32767, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "r4_scalar_scan_n32767_u13_sorted1",
+      "attrs": { "axis": -2, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32767, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "r4_scalar_scan_n32768_u13_sorted0",
+      "attrs": { "axis": -2, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32768, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "r4_scalar_scan_n32768_u13_sorted1",
+      "attrs": { "axis": -2, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32768, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "r4_scalar_scan_n32769_u13_sorted0",
+      "attrs": { "axis": -2, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32769, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "r4_scalar_scan_n32769_u13_sorted1",
+      "attrs": { "axis": -2, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32769, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "r4_scalar_scan_n65537_u13_sorted0",
+      "attrs": { "axis": -2, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 65537, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "r4_scalar_scan_n65537_u13_sorted1",
+      "attrs": { "axis": -2, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 65537, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } },
+      "provenance": {
+        "notes": "Neighboring natural scan sizes at dispatch boundaries with bounded distinct-count oracle."
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "4 * (numel(shapes.x) + numel(shapes.y))"
+          }
+        ]
+      },
+      "preset": "stress"
+    },
+    {
+      "name": "scalar-f32-metadata-mask0-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } }
+    },
+    {
+      "name": "scalar-f32-metadata-mask0-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] } }
+    },
+    {
+      "name": "scalar-f32-metadata-mask1-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] }, "indices": { "dtype": "uint32", "shape": [13] } }
+    },
+    {
+      "name": "scalar-f32-metadata-mask1-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] }, "indices": { "dtype": "uint32", "shape": [13] } }
+    },
+    {
+      "name": "scalar-f32-metadata-mask2-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask2-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask3-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "indices": { "dtype": "uint32", "shape": [13] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask3-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "indices": { "dtype": "uint32", "shape": [13] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask4-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] }, "counts": { "dtype": "uint32", "shape": [13] } }
+    },
+    {
+      "name": "scalar-f32-metadata-mask4-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 13, 1] }, "counts": { "dtype": "uint32", "shape": [13] } }
+    },
+    {
+      "name": "scalar-f32-metadata-mask5-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "indices": { "dtype": "uint32", "shape": [13] },
+        "counts": { "dtype": "uint32", "shape": [13] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask5-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "indices": { "dtype": "uint32", "shape": [13] },
+        "counts": { "dtype": "uint32", "shape": [13] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask6-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] },
+        "counts": { "dtype": "uint32", "shape": [13] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask6-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] },
+        "counts": { "dtype": "uint32", "shape": [13] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask7-sorted0-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "indices": { "dtype": "uint32", "shape": [13] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] },
+        "counts": { "dtype": "uint32", "shape": [13] }
+      }
+    },
+    {
+      "name": "scalar-f32-metadata-mask7-sorted1-n16385-u13",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16385, 1],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0, 11.0, 0.0, 4.0, 6.0, 8.0, 9.0, 10.0, 12.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 13, 1] },
+        "indices": { "dtype": "uint32", "shape": [13] },
+        "inverse_indices": { "dtype": "uint32", "shape": [16385] },
+        "counts": { "dtype": "uint32", "shape": [13] }
+      }
+    },
+    {
+      "name": "single-class-f32-n65537-mask0",
+      "preset": "stress",
+      "inputs": { "x": { "dtype": "float32", "shape": [65537], "data": { "kind": "constant", "value": 3.0 } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1] } }
+    },
+    {
+      "name": "single-class-f32-n65537-mask2",
+      "preset": "stress",
+      "inputs": { "x": { "dtype": "float32", "shape": [65537], "data": { "kind": "constant", "value": 3.0 } } },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1] },
+        "inverse_indices": { "dtype": "uint32", "shape": [65537] }
+      }
+    },
+    {
+      "name": "single-class-f32-n65537-mask4",
+      "preset": "stress",
+      "inputs": { "x": { "dtype": "float32", "shape": [65537], "data": { "kind": "constant", "value": 3.0 } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1] }, "counts": { "dtype": "uint32", "shape": [1] } }
+    },
+    {
+      "name": "single-class-f32-n65537-mask7",
+      "preset": "stress",
+      "inputs": { "x": { "dtype": "float32", "shape": [65537], "data": { "kind": "constant", "value": 3.0 } } },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1] },
+        "indices": { "dtype": "uint32", "shape": [1] },
+        "inverse_indices": { "dtype": "uint32", "shape": [65537] },
+        "counts": { "dtype": "uint32", "shape": [1] }
+      }
+    },
+    {
+      "name": "single-class-axis1-f32-outer3-n4097-inner5-all-metadata",
+      "preset": "stress",
+      "attrs": { "axis": 1, "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4097, 5],
+          "data": { "kind": "cycle", "values": [5.0, 2.0, 7.0, 1.0, 3.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3, 1, 5] },
+        "indices": { "dtype": "uint32", "shape": [1] },
+        "inverse_indices": { "dtype": "uint32", "shape": [4097] },
+        "counts": { "dtype": "uint32", "shape": [1] }
+      }
+    },
+    {
+      "name": "flat-global-crossover-i32-n2048-u2048-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [2048], "dist": "linearMod", "mod": 2048 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [2048] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-i32-n2048-u2048-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "int32", "shape": [2048], "dist": "linearMod", "mod": 2048 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [2048] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-i32-n2049-u2049-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [2049], "dist": "linearMod", "mod": 2049 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [2049] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-i32-n2049-u2049-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "int32", "shape": [2049], "dist": "linearMod", "mod": 2049 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [2049] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-i32-n4097-u4097-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [4097], "dist": "linearMod", "mod": 4097 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [4097] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-i32-n4097-u4097-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "int32", "shape": [4097], "dist": "linearMod", "mod": 4097 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [4097] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-i32-n8193-u8193-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [8193], "dist": "linearMod", "mod": 8193 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [8193] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-i32-n8193-u8193-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "int32", "shape": [8193], "dist": "linearMod", "mod": 8193 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [8193] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "flat-global-crossover-int32-n4097-u2049-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [4097],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/axis-scalar-f32-inverse-hybrid-floor_input_x" }
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2049] } }
+    },
+    {
+      "name": "flat-global-crossover-int32-n4097-u2049-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [4097],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/axis-scalar-f32-inverse-hybrid-floor_input_x" }
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2049] } }
+    },
+    {
+      "name": "flat-global-crossover-uint32-n8193-u2049-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [8193],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/axis-scalar-f32-inverse-hybrid-floor_input_x" }
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [2049] } }
+    },
+    {
+      "name": "flat-global-crossover-uint32-n8193-u2049-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [8193],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/axis-scalar-f32-inverse-hybrid-floor_input_x" }
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [2049] } }
+    },
+    {
+      "name": "output-span-int32-n257-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [257], "dist": "linearMod", "mod": 257 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [257] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "output-span-int32-n257-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "int32", "shape": [257], "dist": "linearMod", "mod": 257 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [257] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "output-span-float32-n257-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [257], "dist": "linearMod", "mod": 257 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [257] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "output-span-float32-n257-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [257], "dist": "linearMod", "mod": 257 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [257] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "output-span-int32-n1025-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "int32", "shape": [1025], "dist": "linearMod", "mod": 1025 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [1025] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "output-span-int32-n1025-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "int32", "shape": [1025], "dist": "linearMod", "mod": 1025 } },
+      "outputs": { "y": { "dtype": "int32", "shape": [1025] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "output-span-float32-n1025-sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [1025], "dist": "linearMod", "mod": 1025 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1025] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "output-span-float32-n1025-sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [1025], "dist": "linearMod", "mod": 1025 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1025] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (numel(shapes.x) + numel(shapes.y))" }] }
+    },
+    {
+      "name": "float_hash_work_float32_n15_u15_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [15], "dist": "linearMod", "mod": 15 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [15] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n15_u15_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [15], "dist": "linearMod", "mod": 15 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [15] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n16_u16_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [16], "dist": "linearMod", "mod": 16 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n16_u16_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [16], "dist": "linearMod", "mod": 16 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n32_u8_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [32], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n32_u8_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [32], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n64_u64_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [64], "dist": "linearMod", "mod": 64 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [64] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n64_u64_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [64], "dist": "linearMod", "mod": 64 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [64] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n128_u8_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [128], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n128_u8_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [128], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n512_u512_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float32", "shape": [512], "dist": "linearMod", "mod": 512 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [512] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float32_n512_u512_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float32", "shape": [512], "dist": "linearMod", "mod": 512 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [512] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float32" }
+    },
+    {
+      "name": "float_hash_work_float16_n15_u15_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float16", "shape": [15], "dist": "linearMod", "mod": 15 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [15] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n15_u15_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float16", "shape": [15], "dist": "linearMod", "mod": 15 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [15] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n16_u16_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float16", "shape": [16], "dist": "linearMod", "mod": 16 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [16] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n16_u16_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float16", "shape": [16], "dist": "linearMod", "mod": 16 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [16] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n32_u8_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float16", "shape": [32], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n32_u8_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float16", "shape": [32], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n64_u64_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float16", "shape": [64], "dist": "linearMod", "mod": 64 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [64] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n64_u64_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float16", "shape": [64], "dist": "linearMod", "mod": 64 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [64] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n128_u8_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float16", "shape": [128], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n128_u8_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float16", "shape": [128], "dist": "linearMod", "mod": 8 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [8] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n512_u512_sorted0",
+      "preset": "stress",
+      "attrs": { "sorted": 0 },
+      "inputs": { "x": { "dtype": "float16", "shape": [512], "dist": "linearMod", "mod": 512 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [512] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
+    },
+    {
+      "name": "float_hash_work_float16_n512_u512_sorted1",
+      "preset": "stress",
+      "attrs": { "sorted": 1 },
+      "inputs": { "x": { "dtype": "float16", "shape": [512], "dist": "linearMod", "mod": 512 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [512] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "dtypeBytes(args.dtype) * (numel(shapes.x) + numel(shapes.y))" }]
+      },
+      "vars": { "dtype": "float16" }
     }
+  ],
+  "fixtureArrays": {
+    "axis_record_table_marker_collision": [7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 0, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 1, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 2, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 3, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 4, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 5, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 6, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 7, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 8, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 9, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 10, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 11, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 12, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 13, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 14, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 15, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 16, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 17, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 18, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 19, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 20, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 21, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 22, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 23, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 24, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 25, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 26, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 27, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 28, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 29, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 30, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 31, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 32, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 33, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 34, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 35, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 36, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 37, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 38, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 39, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 40, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 41, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 42, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 43, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 44, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 45, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 46, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 47, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 48, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 49, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 50, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 51, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 52, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 53, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 54, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 55, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 56, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 57, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 58, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 59, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 60, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 61, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 62, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 63, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 64, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 65, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 66, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 67, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 68, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 69, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 70, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 71, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 72, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 73, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 74, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 75, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 76, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 77, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 78, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 79, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 80, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 81, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 82, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 83, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 84, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 85, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 86, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 87, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 88, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 89, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 90, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 91, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 92, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 93, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 94, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 95, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 96, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 97, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 98, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 99, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 100, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 101, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 102, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 103, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 104, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 105, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 106, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 107, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 108, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 109, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 110, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 111, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 112, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 113, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 114, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 115, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 116, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 117, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 118, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 119, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 120, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 121, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 122, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 123, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 124, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 125, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 126, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 127, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 128, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 129, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 130, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 131, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 132, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 133, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 134, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 135, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 136, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 137, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 138, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 139, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 140, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 141, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 142, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 143, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 144, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 145, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 146, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 147, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 148, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 149, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 150, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 151, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 152, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 153, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 154, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 155, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 156, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 157, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 158, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 159, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 160, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 161, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 162, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 163, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 164, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 165, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 166, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 167, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 168, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 169, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 170, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 171, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 172, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 173, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 174, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 175, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 176, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 177, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 178, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 179, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 180, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 181, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 182, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 183, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 184, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 185, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 186, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 187, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 188, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 189, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 190, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 191, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 192, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 193, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 194, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 195, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 196, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 197, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 198, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 199, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 200, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 201, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 202, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 203, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 204, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 205, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 206, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 207, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 208, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 209, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 210, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 211, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 212, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 213, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 214, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 215, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 216, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 217, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 218, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 219, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 220, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 221, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 222, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 223, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 224, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 225, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 226, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 227, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 228, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 229, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 230, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 231, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 232, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 233, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 234, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 235, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 236, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 237, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 238, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 239, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 240, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 241, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 242, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 243, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 244, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 245, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 246, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 247, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 248, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 249, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 250, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 251, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 252, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 253, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 254, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 255, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 256, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 257, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 258, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 259, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 260, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 261, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 262, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 263, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 264, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 265, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 266, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 267, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 268, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 269, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 270, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 271, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 272, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 273, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 274, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 275, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 276, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 277, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 278, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 279, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 280, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 281, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 282, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 283, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 284, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 285, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 286, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 287, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 288, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 289, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 290, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 291, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 292, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 293, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 294, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 295, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 296, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 297, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 298, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 299, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 300, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 301, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 302, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 303, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 304, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 305, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 306, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 307, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 308, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 309, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 310, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 311, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 312, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 313, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 314, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 315, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 316, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 317, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 318, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 319, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 320, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 321, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 322, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 323, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 324, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 325, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 326, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 327, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 328, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 329, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 330, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 331, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 332, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 333, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 334, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 335, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 336, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 337, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 338, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 339, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 340, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 341, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 342, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 343, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 344, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 345, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 346, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 347, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 348, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 349, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 350, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 351, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 352, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 353, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 354, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 355, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 356, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 357, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 358, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 359, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 360, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 361, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 362, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 363, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 364, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 365, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 366, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 367, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 368, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 369, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 370, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 371, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 372, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 373, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 374, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 375, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 376, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 377, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 378, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 379, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 380, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 381, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 382, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 383, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 384, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 385, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 386, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 387, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 388, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 389, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 390, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 391, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 392, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 393, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 394, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 395, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 396, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 397, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 398, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 399, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 400, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 401, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 402, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 403, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 404, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 405, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 406, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 407, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 408, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 409, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 410, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 411, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 412, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 413, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 414, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 415, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 416, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 417, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 418, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 419, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 420, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 421, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 422, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 423, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 424, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 425, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 426, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 427, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 428, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 429, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 430, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 431, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 432, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 433, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 434, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 435, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 436, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 437, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 438, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 439, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 440, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 441, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 442, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 443, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 444, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 445, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 446, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 447, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 448, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 449, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 450, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 451, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 452, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 453, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 454, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 455, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 456, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 457, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 458, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 459, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 460, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 461, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 462, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 463, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 464, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 465, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 466, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 467, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 468, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 469, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 470, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 471, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 472, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 473, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 474, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 475, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 476, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 477, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 478, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 479, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 480, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 481, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 482, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 483, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 484, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 485, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 486, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 487, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 488, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 489, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 490, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 491, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 492, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 493, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 494, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 495, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 496, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 497, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 498, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 499, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 500, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 501, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 502, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 503, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 504, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 505, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 506, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 507, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 508, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 509, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 510, 17, 7, 7, 7, 7, 7, 11, 7, 7, 7, 7, 13, 7, 7, 7, 511, 17],
+    "axis-scalar-f32-inverse-duplicate64_input_x": [63, 62, 61, 60, 59, 58, 57, 56, 55, 54, 53, 52, 51, 50, 49, 48, 47, 46, 45, 44, 43, 42, 41, 40, 39, 38, 37, 36, 35, 34, 33, 32, 31, 30, 29, 28, 27, 26, 25, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0],
+    "axis-scalar-f32-inverse-hybrid-floor_input_x": [2048, 2047, 2046, 2045, 2044, 2043, 2042, 2041, 2040, 2039, 2038, 2037, 2036, 2035, 2034, 2033, 2032, 2031, 2030, 2029, 2028, 2027, 2026, 2025, 2024, 2023, 2022, 2021, 2020, 2019, 2018, 2017, 2016, 2015, 2014, 2013, 2012, 2011, 2010, 2009, 2008, 2007, 2006, 2005, 2004, 2003, 2002, 2001, 2000, 1999, 1998, 1997, 1996, 1995, 1994, 1993, 1992, 1991, 1990, 1989, 1988, 1987, 1986, 1985, 1984, 1983, 1982, 1981, 1980, 1979, 1978, 1977, 1976, 1975, 1974, 1973, 1972, 1971, 1970, 1969, 1968, 1967, 1966, 1965, 1964, 1963, 1962, 1961, 1960, 1959, 1958, 1957, 1956, 1955, 1954, 1953, 1952, 1951, 1950, 1949, 1948, 1947, 1946, 1945, 1944, 1943, 1942, 1941, 1940, 1939, 1938, 1937, 1936, 1935, 1934, 1933, 1932, 1931, 1930, 1929, 1928, 1927, 1926, 1925, 1924, 1923, 1922, 1921, 1920, 1919, 1918, 1917, 1916, 1915, 1914, 1913, 1912, 1911, 1910, 1909, 1908, 1907, 1906, 1905, 1904, 1903, 1902, 1901, 1900, 1899, 1898, 1897, 1896, 1895, 1894, 1893, 1892, 1891, 1890, 1889, 1888, 1887, 1886, 1885, 1884, 1883, 1882, 1881, 1880, 1879, 1878, 1877, 1876, 1875, 1874, 1873, 1872, 1871, 1870, 1869, 1868, 1867, 1866, 1865, 1864, 1863, 1862, 1861, 1860, 1859, 1858, 1857, 1856, 1855, 1854, 1853, 1852, 1851, 1850, 1849, 1848, 1847, 1846, 1845, 1844, 1843, 1842, 1841, 1840, 1839, 1838, 1837, 1836, 1835, 1834, 1833, 1832, 1831, 1830, 1829, 1828, 1827, 1826, 1825, 1824, 1823, 1822, 1821, 1820, 1819, 1818, 1817, 1816, 1815, 1814, 1813, 1812, 1811, 1810, 1809, 1808, 1807, 1806, 1805, 1804, 1803, 1802, 1801, 1800, 1799, 1798, 1797, 1796, 1795, 1794, 1793, 1792, 1791, 1790, 1789, 1788, 1787, 1786, 1785, 1784, 1783, 1782, 1781, 1780, 1779, 1778, 1777, 1776, 1775, 1774, 1773, 1772, 1771, 1770, 1769, 1768, 1767, 1766, 1765, 1764, 1763, 1762, 1761, 1760, 1759, 1758, 1757, 1756, 1755, 1754, 1753, 1752, 1751, 1750, 1749, 1748, 1747, 1746, 1745, 1744, 1743, 1742, 1741, 1740, 1739, 1738, 1737, 1736, 1735, 1734, 1733, 1732, 1731, 1730, 1729, 1728, 1727, 1726, 1725, 1724, 1723, 1722, 1721, 1720, 1719, 1718, 1717, 1716, 1715, 1714, 1713, 1712, 1711, 1710, 1709, 1708, 1707, 1706, 1705, 1704, 1703, 1702, 1701, 1700, 1699, 1698, 1697, 1696, 1695, 1694, 1693, 1692, 1691, 1690, 1689, 1688, 1687, 1686, 1685, 1684, 1683, 1682, 1681, 1680, 1679, 1678, 1677, 1676, 1675, 1674, 1673, 1672, 1671, 1670, 1669, 1668, 1667, 1666, 1665, 1664, 1663, 1662, 1661, 1660, 1659, 1658, 1657, 1656, 1655, 1654, 1653, 1652, 1651, 1650, 1649, 1648, 1647, 1646, 1645, 1644, 1643, 1642, 1641, 1640, 1639, 1638, 1637, 1636, 1635, 1634, 1633, 1632, 1631, 1630, 1629, 1628, 1627, 1626, 1625, 1624, 1623, 1622, 1621, 1620, 1619, 1618, 1617, 1616, 1615, 1614, 1613, 1612, 1611, 1610, 1609, 1608, 1607, 1606, 1605, 1604, 1603, 1602, 1601, 1600, 1599, 1598, 1597, 1596, 1595, 1594, 1593, 1592, 1591, 1590, 1589, 1588, 1587, 1586, 1585, 1584, 1583, 1582, 1581, 1580, 1579, 1578, 1577, 1576, 1575, 1574, 1573, 1572, 1571, 1570, 1569, 1568, 1567, 1566, 1565, 1564, 1563, 1562, 1561, 1560, 1559, 1558, 1557, 1556, 1555, 1554, 1553, 1552, 1551, 1550, 1549, 1548, 1547, 1546, 1545, 1544, 1543, 1542, 1541, 1540, 1539, 1538, 1537, 1536, 1535, 1534, 1533, 1532, 1531, 1530, 1529, 1528, 1527, 1526, 1525, 1524, 1523, 1522, 1521, 1520, 1519, 1518, 1517, 1516, 1515, 1514, 1513, 1512, 1511, 1510, 1509, 1508, 1507, 1506, 1505, 1504, 1503, 1502, 1501, 1500, 1499, 1498, 1497, 1496, 1495, 1494, 1493, 1492, 1491, 1490, 1489, 1488, 1487, 1486, 1485, 1484, 1483, 1482, 1481, 1480, 1479, 1478, 1477, 1476, 1475, 1474, 1473, 1472, 1471, 1470, 1469, 1468, 1467, 1466, 1465, 1464, 1463, 1462, 1461, 1460, 1459, 1458, 1457, 1456, 1455, 1454, 1453, 1452, 1451, 1450, 1449, 1448, 1447, 1446, 1445, 1444, 1443, 1442, 1441, 1440, 1439, 1438, 1437, 1436, 1435, 1434, 1433, 1432, 1431, 1430, 1429, 1428, 1427, 1426, 1425, 1424, 1423, 1422, 1421, 1420, 1419, 1418, 1417, 1416, 1415, 1414, 1413, 1412, 1411, 1410, 1409, 1408, 1407, 1406, 1405, 1404, 1403, 1402, 1401, 1400, 1399, 1398, 1397, 1396, 1395, 1394, 1393, 1392, 1391, 1390, 1389, 1388, 1387, 1386, 1385, 1384, 1383, 1382, 1381, 1380, 1379, 1378, 1377, 1376, 1375, 1374, 1373, 1372, 1371, 1370, 1369, 1368, 1367, 1366, 1365, 1364, 1363, 1362, 1361, 1360, 1359, 1358, 1357, 1356, 1355, 1354, 1353, 1352, 1351, 1350, 1349, 1348, 1347, 1346, 1345, 1344, 1343, 1342, 1341, 1340, 1339, 1338, 1337, 1336, 1335, 1334, 1333, 1332, 1331, 1330, 1329, 1328, 1327, 1326, 1325, 1324, 1323, 1322, 1321, 1320, 1319, 1318, 1317, 1316, 1315, 1314, 1313, 1312, 1311, 1310, 1309, 1308, 1307, 1306, 1305, 1304, 1303, 1302, 1301, 1300, 1299, 1298, 1297, 1296, 1295, 1294, 1293, 1292, 1291, 1290, 1289, 1288, 1287, 1286, 1285, 1284, 1283, 1282, 1281, 1280, 1279, 1278, 1277, 1276, 1275, 1274, 1273, 1272, 1271, 1270, 1269, 1268, 1267, 1266, 1265, 1264, 1263, 1262, 1261, 1260, 1259, 1258, 1257, 1256, 1255, 1254, 1253, 1252, 1251, 1250, 1249, 1248, 1247, 1246, 1245, 1244, 1243, 1242, 1241, 1240, 1239, 1238, 1237, 1236, 1235, 1234, 1233, 1232, 1231, 1230, 1229, 1228, 1227, 1226, 1225, 1224, 1223, 1222, 1221, 1220, 1219, 1218, 1217, 1216, 1215, 1214, 1213, 1212, 1211, 1210, 1209, 1208, 1207, 1206, 1205, 1204, 1203, 1202, 1201, 1200, 1199, 1198, 1197, 1196, 1195, 1194, 1193, 1192, 1191, 1190, 1189, 1188, 1187, 1186, 1185, 1184, 1183, 1182, 1181, 1180, 1179, 1178, 1177, 1176, 1175, 1174, 1173, 1172, 1171, 1170, 1169, 1168, 1167, 1166, 1165, 1164, 1163, 1162, 1161, 1160, 1159, 1158, 1157, 1156, 1155, 1154, 1153, 1152, 1151, 1150, 1149, 1148, 1147, 1146, 1145, 1144, 1143, 1142, 1141, 1140, 1139, 1138, 1137, 1136, 1135, 1134, 1133, 1132, 1131, 1130, 1129, 1128, 1127, 1126, 1125, 1124, 1123, 1122, 1121, 1120, 1119, 1118, 1117, 1116, 1115, 1114, 1113, 1112, 1111, 1110, 1109, 1108, 1107, 1106, 1105, 1104, 1103, 1102, 1101, 1100, 1099, 1098, 1097, 1096, 1095, 1094, 1093, 1092, 1091, 1090, 1089, 1088, 1087, 1086, 1085, 1084, 1083, 1082, 1081, 1080, 1079, 1078, 1077, 1076, 1075, 1074, 1073, 1072, 1071, 1070, 1069, 1068, 1067, 1066, 1065, 1064, 1063, 1062, 1061, 1060, 1059, 1058, 1057, 1056, 1055, 1054, 1053, 1052, 1051, 1050, 1049, 1048, 1047, 1046, 1045, 1044, 1043, 1042, 1041, 1040, 1039, 1038, 1037, 1036, 1035, 1034, 1033, 1032, 1031, 1030, 1029, 1028, 1027, 1026, 1025, 1024, 1023, 1022, 1021, 1020, 1019, 1018, 1017, 1016, 1015, 1014, 1013, 1012, 1011, 1010, 1009, 1008, 1007, 1006, 1005, 1004, 1003, 1002, 1001, 1000, 999, 998, 997, 996, 995, 994, 993, 992, 991, 990, 989, 988, 987, 986, 985, 984, 983, 982, 981, 980, 979, 978, 977, 976, 975, 974, 973, 972, 971, 970, 969, 968, 967, 966, 965, 964, 963, 962, 961, 960, 959, 958, 957, 956, 955, 954, 953, 952, 951, 950, 949, 948, 947, 946, 945, 944, 943, 942, 941, 940, 939, 938, 937, 936, 935, 934, 933, 932, 931, 930, 929, 928, 927, 926, 925, 924, 923, 922, 921, 920, 919, 918, 917, 916, 915, 914, 913, 912, 911, 910, 909, 908, 907, 906, 905, 904, 903, 902, 901, 900, 899, 898, 897, 896, 895, 894, 893, 892, 891, 890, 889, 888, 887, 886, 885, 884, 883, 882, 881, 880, 879, 878, 877, 876, 875, 874, 873, 872, 871, 870, 869, 868, 867, 866, 865, 864, 863, 862, 861, 860, 859, 858, 857, 856, 855, 854, 853, 852, 851, 850, 849, 848, 847, 846, 845, 844, 843, 842, 841, 840, 839, 838, 837, 836, 835, 834, 833, 832, 831, 830, 829, 828, 827, 826, 825, 824, 823, 822, 821, 820, 819, 818, 817, 816, 815, 814, 813, 812, 811, 810, 809, 808, 807, 806, 805, 804, 803, 802, 801, 800, 799, 798, 797, 796, 795, 794, 793, 792, 791, 790, 789, 788, 787, 786, 785, 784, 783, 782, 781, 780, 779, 778, 777, 776, 775, 774, 773, 772, 771, 770, 769, 768, 767, 766, 765, 764, 763, 762, 761, 760, 759, 758, 757, 756, 755, 754, 753, 752, 751, 750, 749, 748, 747, 746, 745, 744, 743, 742, 741, 740, 739, 738, 737, 736, 735, 734, 733, 732, 731, 730, 729, 728, 727, 726, 725, 724, 723, 722, 721, 720, 719, 718, 717, 716, 715, 714, 713, 712, 711, 710, 709, 708, 707, 706, 705, 704, 703, 702, 701, 700, 699, 698, 697, 696, 695, 694, 693, 692, 691, 690, 689, 688, 687, 686, 685, 684, 683, 682, 681, 680, 679, 678, 677, 676, 675, 674, 673, 672, 671, 670, 669, 668, 667, 666, 665, 664, 663, 662, 661, 660, 659, 658, 657, 656, 655, 654, 653, 652, 651, 650, 649, 648, 647, 646, 645, 644, 643, 642, 641, 640, 639, 638, 637, 636, 635, 634, 633, 632, 631, 630, 629, 628, 627, 626, 625, 624, 623, 622, 621, 620, 619, 618, 617, 616, 615, 614, 613, 612, 611, 610, 609, 608, 607, 606, 605, 604, 603, 602, 601, 600, 599, 598, 597, 596, 595, 594, 593, 592, 591, 590, 589, 588, 587, 586, 585, 584, 583, 582, 581, 580, 579, 578, 577, 576, 575, 574, 573, 572, 571, 570, 569, 568, 567, 566, 565, 564, 563, 562, 561, 560, 559, 558, 557, 556, 555, 554, 553, 552, 551, 550, 549, 548, 547, 546, 545, 544, 543, 542, 541, 540, 539, 538, 537, 536, 535, 534, 533, 532, 531, 530, 529, 528, 527, 526, 525, 524, 523, 522, 521, 520, 519, 518, 517, 516, 515, 514, 513, 512, 511, 510, 509, 508, 507, 506, 505, 504, 503, 502, 501, 500, 499, 498, 497, 496, 495, 494, 493, 492, 491, 490, 489, 488, 487, 486, 485, 484, 483, 482, 481, 480, 479, 478, 477, 476, 475, 474, 473, 472, 471, 470, 469, 468, 467, 466, 465, 464, 463, 462, 461, 460, 459, 458, 457, 456, 455, 454, 453, 452, 451, 450, 449, 448, 447, 446, 445, 444, 443, 442, 441, 440, 439, 438, 437, 436, 435, 434, 433, 432, 431, 430, 429, 428, 427, 426, 425, 424, 423, 422, 421, 420, 419, 418, 417, 416, 415, 414, 413, 412, 411, 410, 409, 408, 407, 406, 405, 404, 403, 402, 401, 400, 399, 398, 397, 396, 395, 394, 393, 392, 391, 390, 389, 388, 387, 386, 385, 384, 383, 382, 381, 380, 379, 378, 377, 376, 375, 374, 373, 372, 371, 370, 369, 368, 367, 366, 365, 364, 363, 362, 361, 360, 359, 358, 357, 356, 355, 354, 353, 352, 351, 350, 349, 348, 347, 346, 345, 344, 343, 342, 341, 340, 339, 338, 337, 336, 335, 334, 333, 332, 331, 330, 329, 328, 327, 326, 325, 324, 323, 322, 321, 320, 319, 318, 317, 316, 315, 314, 313, 312, 311, 310, 309, 308, 307, 306, 305, 304, 303, 302, 301, 300, 299, 298, 297, 296, 295, 294, 293, 292, 291, 290, 289, 288, 287, 286, 285, 284, 283, 282, 281, 280, 279, 278, 277, 276, 275, 274, 273, 272, 271, 270, 269, 268, 267, 266, 265, 264, 263, 262, 261, 260, 259, 258, 257, 256, 255, 254, 253, 252, 251, 250, 249, 248, 247, 246, 245, 244, 243, 242, 241, 240, 239, 238, 237, 236, 235, 234, 233, 232, 231, 230, 229, 228, 227, 226, 225, 224, 223, 222, 221, 220, 219, 218, 217, 216, 215, 214, 213, 212, 211, 210, 209, 208, 207, 206, 205, 204, 203, 202, 201, 200, 199, 198, 197, 196, 195, 194, 193, 192, 191, 190, 189, 188, 187, 186, 185, 184, 183, 182, 181, 180, 179, 178, 177, 176, 175, 174, 173, 172, 171, 170, 169, 168, 167, 166, 165, 164, 163, 162, 161, 160, 159, 158, 157, 156, 155, 154, 153, 152, 151, 150, 149, 148, 147, 146, 145, 144, 143, 142, 141, 140, 139, 138, 137, 136, 135, 134, 133, 132, 131, 130, 129, 128, 127, 126, 125, 124, 123, 122, 121, 120, 119, 118, 117, 116, 115, 114, 113, 112, 111, 110, 109, 108, 107, 106, 105, 104, 103, 102, 101, 100, 99, 98, 97, 96, 95, 94, 93, 92, 91, 90, 89, 88, 87, 86, 85, 84, 83, 82, 81, 80, 79, 78, 77, 76, 75, 74, 73, 72, 71, 70, 69, 68, 67, 66, 65, 64, 63, 62, 61, 60, 59, 58, 57, 56, 55, 54, 53, 52, 51, 50, 49, 48, 47, 46, 45, 44, 43, 42, 41, 40, 39, 38, 37, 36, 35, 34, 33, 32, 31, 30, 29, 28, 27, 26, 25, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0]
+  }
 }

build/webgpu/manifest.json CHANGED Viewed

The diff for this file is too large to render. See raw diff

build/webgpu/metadata.json CHANGED Viewed

@@ -1,30 +1,116 @@
 {
   "name": "ai.onnx.Unique",
-  "id": "_ai_onnx_unique_webgpu_6e781ae",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "B7+zElIduemQX6xUT7ASYxUA4g57rhMrU3Lt1A761IQ=",
-      "manifest.json": "/vKIMvGvv1gB51KCCioQsOc43XoxZzBZir/AM+oP7qk=",
-      "test.json": "GUvkml01vopw45if9L8HSq5uiIzpmkFLGUH8I3StJVQ=",
-      "unique-axis-compact-sort.wgsl.jinja": "PJRMzhP+QgoHpmF7c9m2pQLOmHUL5aZTX7Ywhxqjrxs=",
-      "unique-axis-dedup.wgsl.jinja": "O1EWypziZGw+vSmIZWyYo7IBihHTAEEYrUtUSdA+sk4=",
-      "unique-axis-hash.wgsl.jinja": "Q8xwZ5T74OaJS9TpQURo3QPHXzpmMCA8ZbBfl/bucQA=",
-      "unique-axis-scatter.wgsl.jinja": "Yjvj7qKzNzVJNKK8NeZ7ooDSetmEl218RIw7ftyt0lk=",
-      "unique-axis.wgsl.jinja": "fjQxtH8MjP43bhdPOcxya+xVerIi30R4cgs7URqc71s=",
-      "unique-compact-sort.wgsl.jinja": "jeh9pkkZIYSv2jQA/StxvAaQTVp/sI0eUsekeYSnkvo=",
       "unique-dedup.wgsl.jinja": "muTd580L/rfaL/p5NjOwMPp2F3ALTR9S7SzVYl9r+wQ=",
-      "unique-hash-build.wgsl.jinja": "HkomLl6SIyeWaOI0dRrXbl2chUwav16ClFwxZ+EDZzY=",
-      "unique-hash-collect.wgsl.jinja": "5+RKxUzJlrcqJ3Z/RfyznHl43WSoVKJSwGlQ/KzwnI4=",
-      "unique-hash-init.wgsl.jinja": "d4tz0KAxmNaRn4ZvRQemr+7wk4GwRZY9tLQguTuIcGo=",
-      "unique-hash-mark.wgsl.jinja": "87mwSXE1BpY/3zHxtaOhLpm4IZAF8hSX6bLvr6GfncE=",
-      "unique-hash-sort-collected-key-only.wgsl.jinja": "Rpuf6QQZwx2XvogNtJoTS5/vsNHbA721YwBqwWt0SLQ=",
-      "unique.wgsl.jinja": "J4xqWcvx4Ca4uW4ZbJ1xG2pfZbvL3cdnYU33fgHg1eo="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Unique" }
 }

 {
   "name": "ai.onnx.Unique",
+  "id": "_ai_onnx_unique_webgpu_b3bffe7",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "6IjCP8Xe2mSSuxPK7xy73JMu42BF5ZgUhgwgS3Y2qcM=",
+      "manifest.json": "LLZErOf8Z+wJ7nP3zzdL38BwDHrMScCWVqLZ/v/bC/I=",
+      "scan-block-prefix-u32.wgsl.jinja": "EDY0dOyl7TX0GjGdfkGRqho+8VaqNTv69XIUAGWwW7w=",
+      "test.json": "4n/yg8WUYERuAKaqDlVQLlpguI6LoD/VkjllQSd/r+A=",
+      "unique-axis-compact-sort.wgsl.jinja": "USY2zwmM0qLzs5Zu+jeWhvsF/kksOxlMdBCRQyzSeEI=",
+      "unique-axis-dedup.wgsl.jinja": "tdp/nGtmgjimm0FRK2Q+nVPJXabpIbsE53/N88b6ZkI=",
+      "unique-axis-hash.wgsl.jinja": "KIRzJC48pNLKin6XkQlYmlOdj7UaMvHlHkORtDJBHbg=",
+      "unique-axis-scalar-inverse.wgsl.jinja": "8XJ/PidrZqxZ7E4exCzsghcDAnDkBWO6XhKoneojGMo=",
+      "unique-axis-scalar-ranks.wgsl.jinja": "MKa09aSgiScgGAke+eaZ8ECSPtZl1DvQRJhNnsGhx9Q=",
+      "unique-axis-scatter.wgsl.jinja": "H4nuOCxlxZHOBVm71ht58g4oFZfEwT+pXLXbx0yhjU0=",
+      "unique-axis.wgsl.jinja": "xGYbso+QeI9d4zpCXmrwWaueY1/dNII5ton7Nl6qAZQ=",
+      "unique-compact-sort.wgsl.jinja": "V1et+Wyn4847Ej1jkpjfPGBfexgbpKk1bk6VIapFU3o=",
       "unique-dedup.wgsl.jinja": "muTd580L/rfaL/p5NjOwMPp2F3ALTR9S7SzVYl9r+wQ=",
+      "unique-flag-block-scan.wgsl.jinja": "UHiW1pUoMetLcSVFtbb27NBiHyrj1BqnpvsmKHR4eUo=",
+      "unique-flat-metadata.wgsl.jinja": "DjjbOABSeqFVcZWMK3tfyBjT+1IxqRaF8IkSb1LbHNI=",
+      "unique-global-sort-exchange.wgsl.jinja": "ycyAh+i1Jf8f827QgsW8JFZsyPuTSq8dTe9eID4FtG8=",
+      "unique-global-sort-output.wgsl.jinja": "OHzPDIUYaZ71Lb0s+dkXqvYTUyQD5j5WOrb4pl2ehcE=",
+      "unique-global-sort-shared.wgsl.jinja": "PrT6/StTV6cAJKH9RsoQ9DL8jMi7xgBiUoxIl5Nhbuo=",
+      "unique-hash-build.wgsl.jinja": "R6FsBjCfL8Gs8Obi+UT7TPXH66BM4PlmA7N5zcf6boY=",
+      "unique-hash-collect.wgsl.jinja": "QjGT7dXhz4D9ka0Ry6EDJyCk+4ZxqEWZ7CQvaLl00lg=",
+      "unique-hash-init.wgsl.jinja": "oijfIva4kwhMMSpHUUWnwAEA3mZZhVSr4wQaXHgqaBI=",
+      "unique-hash-mark.wgsl.jinja": "vNqJd6BsN5s1QyOwoZfurA7FMKeSDA7n8+8t9IzvDaA=",
+      "unique-hash-sort-collected-key-only.wgsl.jinja": "r3/0IAivXcDWz5je9+0kn+EMahlO7IsQAqggcMPG01g=",
+      "unique-scalar-compact.wgsl.jinja": "icYtAk99HoPQMRoKcaDuBKhbiBZQabtx1+EnvM9AXSU=",
+      "unique-scalar-metadata.wgsl.jinja": "9K5i3e7rE0hGfnS9y1knf9M/IZNMRztJ2/mijo1uUvw=",
+      "unique-scalar-output.wgsl.jinja": "AzH6/sqeVpuLzEA8BN8I9u7eOz3cqvoDqLst5Awh31Q=",
+      "unique-single-class.wgsl.jinja": "7eZJTSljWA00bg6rPO5YbebToQuzuhd99IFCXbf96Fs=",
+      "unique.wgsl.jinja": "cwLLFAvO05OQZw6C6xpXeQgJPuwQTymG1PjBtHwwBmM="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "single_class_y": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_y_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_y_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "single_class_indices": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_indices_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_indices_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "single_class_inverse": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_inverse_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_inverse_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "single_class_indices_inverse": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_indices_inverse_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_indices_inverse_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "single_class_counts": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_counts_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_counts_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "single_class_indices_counts": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_indices_counts_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_indices_counts_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "single_class_inverse_counts": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_inverse_counts_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_inverse_counts_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "single_class_indices_inverse_counts": ["unique-single-class.wgsl.jinja"],
+      "scalar_hash_parallel_indices_inverse_counts_unsorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "scalar_hash_parallel_indices_inverse_counts_sorted": ["scan-block-prefix-u32.wgsl.jinja", "unique-flag-block-scan.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja", "unique-scalar-compact.wgsl.jinja", "unique-scalar-metadata.wgsl.jinja", "unique-scalar-output.wgsl.jinja"],
+      "hash_dedup_hybrid_sort_large_y": ["unique-compact-sort.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-output.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "hash_dedup_large_y": ["unique-compact-sort.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "hash_dedup_large_y_unsorted": ["unique-compact-sort.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "axis_scalar_hash_hybrid_metadata_inverse": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-scalar-inverse.wgsl.jinja", "unique-axis-scalar-ranks.wgsl.jinja", "unique-axis-scatter.wgsl.jinja", "unique-global-sort-exchange.wgsl.jinja", "unique-global-sort-shared.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "axis_scalar_hash_metadata_inverse": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-scalar-inverse.wgsl.jinja", "unique-axis-scalar-ranks.wgsl.jinja", "unique-axis-scatter.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "axis_scalar_hash_metadata_inverse_unsorted": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-scalar-inverse.wgsl.jinja", "unique-axis-scalar-ranks.wgsl.jinja", "unique-axis-scatter.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "hash_dedup_sorted_collect": ["unique-hash-build.wgsl.jinja", "unique-hash-collect.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-sort-collected-key-only.wgsl.jinja"],
+      "hash_dedup": ["unique-compact-sort.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "rank1_bounded": ["unique.wgsl.jinja"],
+      "axis_serial": ["unique-axis.wgsl.jinja"],
+      "axis_hash_bounded": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-hash.wgsl.jinja", "unique-axis-scatter.wgsl.jinja"],
+      "axis_hash_bounded_unsorted": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-hash.wgsl.jinja", "unique-axis-scatter.wgsl.jinja"],
+      "axis_hash_bounded_nan_safe": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-hash.wgsl.jinja", "unique-axis-scatter.wgsl.jinja", "unique-axis.wgsl.jinja"],
+      "axis_hash_bounded_unsorted_nan_safe": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-hash.wgsl.jinja", "unique-axis-scatter.wgsl.jinja", "unique-axis.wgsl.jinja"],
+      "parallel_dedup": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja"],
+      "parallel_dedup_large_y": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja"],
+      "parallel_dedup_large_y_unsorted": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja"],
+      "axis_bounded": ["unique-axis-compact-sort.wgsl.jinja", "unique-axis-dedup.wgsl.jinja"],
+      "flat_metadata_indices": ["unique.wgsl.jinja"],
+      "axis_metadata_indices": ["unique-axis.wgsl.jinja"],
+      "flat_metadata_inverse": ["unique.wgsl.jinja"],
+      "axis_metadata_inverse": ["unique-axis.wgsl.jinja"],
+      "flat_metadata_counts": ["unique.wgsl.jinja"],
+      "axis_metadata_counts": ["unique-axis.wgsl.jinja"],
+      "flat_metadata_indices_inverse": ["unique.wgsl.jinja"],
+      "axis_metadata_indices_inverse": ["unique-axis.wgsl.jinja"],
+      "flat_metadata_indices_counts": ["unique.wgsl.jinja"],
+      "axis_metadata_indices_counts": ["unique-axis.wgsl.jinja"],
+      "flat_metadata_inverse_counts": ["unique.wgsl.jinja"],
+      "axis_metadata_inverse_counts": ["unique-axis.wgsl.jinja"],
+      "flat_metadata_all": ["unique.wgsl.jinja"],
+      "axis_metadata_all": ["unique-axis.wgsl.jinja"],
+      "flat_parallel_metadata_indices": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja", "unique-flat-metadata.wgsl.jinja"],
+      "flat_parallel_metadata_inverse": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja", "unique-flat-metadata.wgsl.jinja"],
+      "flat_parallel_metadata_indices_inverse": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja", "unique-flat-metadata.wgsl.jinja"],
+      "flat_parallel_metadata_counts": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja", "unique-flat-metadata.wgsl.jinja"],
+      "flat_parallel_metadata_indices_counts": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja", "unique-flat-metadata.wgsl.jinja"],
+      "flat_parallel_metadata_inverse_counts": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja", "unique-flat-metadata.wgsl.jinja"],
+      "flat_parallel_metadata_all": ["unique-compact-sort.wgsl.jinja", "unique-dedup.wgsl.jinja", "unique-flat-metadata.wgsl.jinja"],
+      "flat_hash_metadata_indices": ["unique-compact-sort.wgsl.jinja", "unique-flat-metadata.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "flat_hash_metadata_inverse": ["unique-compact-sort.wgsl.jinja", "unique-flat-metadata.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "flat_hash_metadata_indices_inverse": ["unique-compact-sort.wgsl.jinja", "unique-flat-metadata.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "flat_hash_metadata_counts": ["unique-compact-sort.wgsl.jinja", "unique-flat-metadata.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "flat_hash_metadata_indices_counts": ["unique-compact-sort.wgsl.jinja", "unique-flat-metadata.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "flat_hash_metadata_inverse_counts": ["unique-compact-sort.wgsl.jinja", "unique-flat-metadata.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"],
+      "flat_hash_metadata_all": ["unique-compact-sort.wgsl.jinja", "unique-flat-metadata.wgsl.jinja", "unique-hash-build.wgsl.jinja", "unique-hash-init.wgsl.jinja", "unique-hash-mark.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/scan-block-prefix-u32.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,44 @@

+// Single-workgroup exclusive scan over per-block sums for stream compaction.
+// The dispatch is (1, 1, 1). It walks blockSums in workgroup-sized chunks,
+// scans each chunk with Hillis-Steele or subgroup collectives, and links chunks
+// with a running carry. blockPrefix[b] is therefore the number of set flags in
+// all blocks before b. The total count is the last prefix plus the last sum.
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+var<workgroup> wgScan: array<u32, WG>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
+  let tid = lid.x;
+  var carry = 0u;
+  let chunks = (params.numBlocks + WG - 1u) / WG;
+  for (var c = 0u; c < chunks; c = c + 1u) {
+    let j = c * WG + tid;
+    var v = 0u;
+    if (j < params.numBlocks) {
+      v = blockSums[j];
+    }
+    workgroupBarrier();
+    wgScan[tid] = v;
+    for (var step = 1u; step < WG; step = step << 1u) {
+      workgroupBarrier();
+      var prev = 0u;
+      if (tid >= step) {
+        prev = wgScan[tid - step];
+      }
+      workgroupBarrier();
+      wgScan[tid] = wgScan[tid] + prev;
+    }
+    workgroupBarrier();
+    if (j < params.numBlocks) {
+      var value = carry;
+      if (tid > 0u) {
+        value = carry + wgScan[tid - 1u];
+      }
+      blockPrefix[j] = value;
+    }
+    carry = carry + wgScan[WG - 1u];
+  }
+}

build/webgpu/test.json CHANGED Viewed

The diff for this file is too large to render. See raw diff

build/webgpu/unique-axis-compact-sort.wgsl.jinja CHANGED Viewed

@@ -1,16 +1,16 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-{% set emitCount = source.emitCount | default(false) %}
 // Compact every first-occurrence slice index, optionally bitonic-sort them,
 // then scatter the exact result. Small results keep slots in workgroup memory;
 // larger results use storage-buffer scratch with storage barriers.
-const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
-const CAP: u32 = {{ source.capacity }}u;
-{% if sorted or not source.globalScratch %}
-const SORT_N: u32 = {{ source.sortN }}u;
 {% endif %}
 {% if sorted %}
@@ -38,6 +38,7 @@ fn canon_key(v: {{ scalar }}) -> u32 {
 }
 {% endif %}
 fn less_value(a: {{ scalar }}, b: {{ scalar }}) -> bool {
 {% if isFloat %}
   return canon_key(a) < canon_key(b);
@@ -45,9 +46,10 @@ fn less_value(a: {{ scalar }}, b: {{ scalar }}) -> bool {
   return a < b;
 {% endif %}
 }
 {% endif %}
-{% if sorted or not source.compactOnly %}
 fn slice_at(o: u32, k: u32, n: u32) -> {{ scalar }} {
   return x[(o * params.axisDim + k) * params.inner + n];
 }
@@ -56,6 +58,7 @@ fn slice_at(o: u32, k: u32, n: u32) -> {{ scalar }} {
 {% if sorted %}
 // True iff slice a < slice b lexicographically over (outer, inner) element order.
 // The serial insertion sort uses the same comparator.
 fn slice_less(a: u32, b: u32) -> bool {
   for (var o = 0u; o < params.outer; o = o + 1u) {
     for (var n = 0u; n < params.inner; n = n + 1u) {
@@ -68,6 +71,7 @@ fn slice_less(a: u32, b: u32) -> bool {
   return false;
 }
 // Cache the first element's monotonic integer-order key alongside each slot so
 // the bitonic network normally compares only threadgroup u32s; equal keys still
 // use the exact full-slice comparator.
@@ -83,7 +87,7 @@ fn slice_primary_key(k: u32) -> u32 {
 }
 {% endif %}
-{% if not source.compactOnly %}
 // Keep exact-result axis scatter identical across static- and dynamic-shape
 // dispatch strategies. The zero branch is a defensive guard for invalid shapes.
 fn unique_axis_zero_value() -> {{ scalar }} {
@@ -91,10 +95,17 @@ fn unique_axis_zero_value() -> {{ scalar }} {
 }
 fn unique_axis_scatter_element(g: u32, written: u32) {
   let n = g % params.inner;
   let tmp = g / params.inner;
   let p = tmp % params.outputAxisDim;
   let o = tmp / params.outputAxisDim;
   if (p < written) {
     y[g] = slice_at(o, slots[p], n);
   } else {
@@ -105,11 +116,11 @@ fn unique_axis_scatter_element(g: u32, written: u32) {
 {% endif %}
 var<workgroup> wgScan: array<u32, WG>;
-{% if not source.globalScratch %}
 var<workgroup> slots: array<u32, SORT_N>;   // compacted (then sorted) slice indices
 {% endif %}
 var<workgroup> wgCarry: u32;
-{% if sorted and not source.globalScratch %}
 var<workgroup> sortPad: array<u32, SORT_N>;  // 1 = padding slot, sorts after every real
 var<workgroup> sortKey: array<u32, SORT_N>;  // first-element total-order key
@@ -156,7 +167,7 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
   }
   let total = wgCarry;
-{% if sorted or not source.compactOnly %}
   let written = min(total, CAP);
 {% endif %}
 {% if emitCount %}
@@ -164,7 +175,7 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
     count[0] = total;
   }
 {% endif %}
-{% if source.globalScratch %}
   storageBarrier();
 {% endif %}
@@ -181,12 +192,13 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
       sortKey[k] = slice_primary_key(slots[k]);
     }
   }
-{% if source.globalScratch %}
   storageBarrier();
 {% else %}
   workgroupBarrier();
 {% endif %}
   var size = 2u;
   loop {
     if (size > SORT_N) { break; }
@@ -228,7 +240,7 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
           }
         }
       }
-{% if source.globalScratch %}
       storageBarrier();
 {% else %}
       workgroupBarrier();
@@ -238,12 +250,10 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
     size = size * 2u;
   }
 {% endif %}
-{% if not source.compactOnly %}
-{% if source.globalScratch %}
-  storageBarrier();
-{% else %}
-  workgroupBarrier();
 {% endif %}
   // Scatter the distinct slices into the exact Y allocation: one lane per
   // output element, flat index g = (o*outputAxisDim + p)*inner + n.

 {{ env.wgsl.resourceDeclarations }}
+{% set emitCount = emitCount | default(false) %}
 // Compact every first-occurrence slice index, optionally bitonic-sort them,
 // then scatter the exact result. Small results keep slots in workgroup memory;
 // larger results use storage-buffer scratch with storage barriers.
+// Global-scratch variants reserve workgroup memory only for the scan buffer and
+// provide an independent sort width. Workgroup-scratch variants use the shared
+// workgroup-width setting.
+const WG: u32 = {{ sortWg | default(tunables.WORKGROUP_SIZE) }}u;
+const CAP: u32 = {{ capacity }}u;
+{% if sorted or not globalScratch %}
+const SORT_N: u32 = {{ sortN }}u;
 {% endif %}
 {% if sorted %}
 }
 {% endif %}
+{% if not (sortExternally | default(false)) %}
 fn less_value(a: {{ scalar }}, b: {{ scalar }}) -> bool {
 {% if isFloat %}
   return canon_key(a) < canon_key(b);
   return a < b;
 {% endif %}
 }
+{% endif %}
 {% endif %}
+{% if sorted or not compactOnly %}
 fn slice_at(o: u32, k: u32, n: u32) -> {{ scalar }} {
   return x[(o * params.axisDim + k) * params.inner + n];
 }
 {% if sorted %}
 // True iff slice a < slice b lexicographically over (outer, inner) element order.
 // The serial insertion sort uses the same comparator.
+{% if not (sortExternally | default(false)) %}
 fn slice_less(a: u32, b: u32) -> bool {
   for (var o = 0u; o < params.outer; o = o + 1u) {
     for (var n = 0u; n < params.inner; n = n + 1u) {
   return false;
 }
+{% endif %}
 // Cache the first element's monotonic integer-order key alongside each slot so
 // the bitonic network normally compares only threadgroup u32s; equal keys still
 // use the exact full-slice comparator.
 }
 {% endif %}
+{% if not compactOnly %}
 // Keep exact-result axis scatter identical across static- and dynamic-shape
 // dispatch strategies. The zero branch is a defensive guard for invalid shapes.
 fn unique_axis_zero_value() -> {{ scalar }} {
 }
 fn unique_axis_scatter_element(g: u32, written: u32) {
+{% if staticShape is defined and staticShape %}
+  let n = g % INNER;
+  let tmp = g / INNER;
+  let p = tmp % OUTPUT_AXIS_DIM;
+  let o = tmp / OUTPUT_AXIS_DIM;
+{% else %}
   let n = g % params.inner;
   let tmp = g / params.inner;
   let p = tmp % params.outputAxisDim;
   let o = tmp / params.outputAxisDim;
+{% endif %}
   if (p < written) {
     y[g] = slice_at(o, slots[p], n);
   } else {
 {% endif %}
 var<workgroup> wgScan: array<u32, WG>;
+{% if not globalScratch %}
 var<workgroup> slots: array<u32, SORT_N>;   // compacted (then sorted) slice indices
 {% endif %}
 var<workgroup> wgCarry: u32;
+{% if sorted and not globalScratch %}
 var<workgroup> sortPad: array<u32, SORT_N>;  // 1 = padding slot, sorts after every real
 var<workgroup> sortKey: array<u32, SORT_N>;  // first-element total-order key
   }
   let total = wgCarry;
+{% if sorted or not compactOnly %}
   let written = min(total, CAP);
 {% endif %}
 {% if emitCount %}
     count[0] = total;
   }
 {% endif %}
+{% if globalScratch %}
   storageBarrier();
 {% endif %}
       sortKey[k] = slice_primary_key(slots[k]);
     }
   }
+{% if globalScratch %}
   storageBarrier();
 {% else %}
   workgroupBarrier();
 {% endif %}
+{% if not (sortExternally | default(false)) %}
   var size = 2u;
   loop {
     if (size > SORT_N) { break; }
           }
         }
       }
+{% if globalScratch %}
       storageBarrier();
 {% else %}
       workgroupBarrier();
     size = size * 2u;
   }
 {% endif %}
 {% endif %}
+{% if not compactOnly %}
+  // The workgroup-scratch form scatters its compacted slots directly.
+  workgroupBarrier();
   // Scatter the distinct slices into the exact Y allocation: one lane per
   // output element, flat index g = (o*outputAxisDim + p)*inner + n.

build/webgpu/unique-axis-dedup.wgsl.jinja CHANGED Viewed

@@ -7,17 +7,15 @@
 // (o in outer, n in inner). This parallel path compares only integer storage
 // values.
 //
-// The quadratic scan's common case only needs to prove that two slices differ.
-// Loading the first element of every candidate directly from X made that scan a
-// large, cache-hostile gather (the axis stride can be kilobytes). Instead, every
-// workgroup stages a chunk of small sampled slice fingerprints in threadgroup
-// memory. A fingerprint mismatch proves inequality; collisions fall through to
-// the exact element-wise comparator, so this is an optimization only and cannot
-// change Unique semantics.
 fn slice_at(o: u32, k: u32, n: u32) -> {{ scalar }} {
   return x[(o * params.axisDim + k) * params.inner + n];
 }
 fn value_key(value: {{ scalar }}) -> u32 {
   // Integer conversion is injective for the 8/32-bit integer storage types.
   return u32(value);
@@ -29,23 +27,26 @@ fn mix_key(hash: u32, value: u32) -> u32 {
   return h;
 }
 fn slice_key(k: u32) -> u32 {
-  let elements = params.outer * params.inner;
-  if (elements == 0u) { return 0u; }
-  // Four evenly-spaced samples make the prefilter useful for structured rows
-  // whose leading value is shared, while keeping the staging pass tiny.
-  let last = elements - 1u;
   var hash = 0x811c9dc5u;
-  for (var sample = 0u; sample < 4u; sample = sample + 1u) {
-    let logical = (sample * last) / 3u;
-    let o = logical / params.inner;
-    let n = logical % params.inner;
-    hash = mix_key(hash, value_key(slice_at(o, k, n)));
   }
   return hash;
 }
 fn slice_eq(a: u32, b: u32) -> bool {
   for (var o = 0u; o < params.outer; o = o + 1u) {
     for (var n = 0u; n < params.inner; n = n + 1u) {
@@ -66,7 +67,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>,
   let laneValid = k < params.axisDim;
   var kKey = 0u;
   if (laneValid) {
-    kKey = slice_key(k);
   }
   var seen = false;
@@ -78,7 +79,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>,
     let loadK = chunkBase + lid.x;
     var loadedKey = 0u;
     if (loadK < params.axisDim) {
-      loadedKey = slice_key(loadK);
     }
     keyCache[lid.x] = loadedKey;
     workgroupBarrier();
@@ -103,3 +104,4 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>,
     firstFlag[k] = select(1u, 0u, seen);
   }
 }

 // (o in outer, n in inner). This parallel path compares only integer storage
 // values.
 //
+// A separate pass fingerprints each slice into `sliceKeys`, and each workgroup
+// stages a chunk of those fingerprints in shared memory. A mismatch proves
+// inequality; collisions fall through to the exact element-wise comparator, so
+// fingerprints cannot change Unique semantics.
 fn slice_at(o: u32, k: u32, n: u32) -> {{ scalar }} {
   return x[(o * params.axisDim + k) * params.inner + n];
 }
+{% if stage == "keys" %}
 fn value_key(value: {{ scalar }}) -> u32 {
   // Integer conversion is injective for the 8/32-bit integer storage types.
   return u32(value);
   return h;
 }
+// The fingerprint covers every element of the slice and is computed once per
+// slice before the pairwise scan. Equal fingerprints still require the exact
+// comparator because collisions are possible.
 fn slice_key(k: u32) -> u32 {
   var hash = 0x811c9dc5u;
+  for (var o = 0u; o < params.outer; o = o + 1u) {
+    for (var n = 0u; n < params.inner; n = n + 1u) {
+      hash = mix_key(hash, value_key(slice_at(o, k, n)));
+    }
   }
   return hash;
 }
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let k = gid.x;
+  if (k >= params.axisDim) { return; }
+  sliceKeys[k] = slice_key(k);
+}
+{% else %}
 fn slice_eq(a: u32, b: u32) -> bool {
   for (var o = 0u; o < params.outer; o = o + 1u) {
     for (var n = 0u; n < params.inner; n = n + 1u) {
   let laneValid = k < params.axisDim;
   var kKey = 0u;
   if (laneValid) {
+    kKey = sliceKeys[k];
   }
   var seen = false;
     let loadK = chunkBase + lid.x;
     var loadedKey = 0u;
     if (loadK < params.axisDim) {
+      loadedKey = sliceKeys[loadK];
     }
     keyCache[lid.x] = loadedKey;
     workgroupBarrier();
     firstFlag[k] = select(1u, 0u, seen);
   }
 }
+{% endif %}

build/webgpu/unique-axis-hash.wgsl.jinja CHANGED Viewed

@@ -1,3 +1,4 @@
 // Hash-backed axis Unique front end. Initialization computes one full-slice
 // hash per axis index and clears the open-addressed table. Build inserts
 // exact-equality buckets and atomically retains the earliest representative.
@@ -5,19 +6,17 @@
 // Hash collisions always use the exact slice comparator, so hashes never
 // affect semantics. Floats hash and compare canonical FTZ-safe bits (one NaN
 // class, -0 folded to +0), so bucket equality is transitive and exact.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
-const TABLE_SIZE: u32 = {{ source.tableSize }}u;
-{% if source.stage == "build" %}
 const TABLE_MASK: u32 = TABLE_SIZE - 1u;
 {% endif %}
 const EMPTY: u32 = 0xffffffffu;
-{% if source.stage == "init" %}
 {% if isFloat %}
 fn float_bits(v: {{ scalar }}) -> u32 {
 {% if usesF16 %}
@@ -53,30 +52,51 @@ fn slice_hash(k: u32) -> u32 {
   // FNV-1a over the complete slice. The axis index is deliberately excluded:
   // equal slices must hash identically regardless of where they occur.
   var hash = 0x811c9dc5u;
   for (var o = 0u; o < params.outer; o += 1u) {
     for (var n = 0u; n < params.inner; n += 1u) {
       hash = (hash ^ value_key(slice_at(o, k, n))) * 0x01000193u;
     }
   }
   // One final avalanche reduces clustering for short structured slices.
   hash ^= hash >> 16u;
   hash *= 0x7feb352du;
   hash ^= hash >> 15u;
   return hash;
 }
 @compute @workgroup_size(WG)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let i = gid.x;
   if (i < TABLE_SIZE) {
     atomicStore(&hashSlot[i], EMPTY);
   }
   if (i < params.axisDim) {
     firstFlag[i] = 0u;
     sliceHash[i] = slice_hash(i);
   }
 }
-{% elif source.stage == "build" %}
 {% if isFloat %}
 fn float_bits(v: {{ scalar }}) -> u32 {
@@ -118,6 +138,11 @@ fn slice_eq(a: u32, b: u32) -> bool {
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let k = gid.x;
   if (k >= params.axisDim) { return; }
   let key = sliceHash[k];
   var slot = key & TABLE_MASK;
   for (var probe = 0u; probe < TABLE_SIZE; probe += 1u) {

+{% set trackNaN = trackNaN | default(false) %}
 // Hash-backed axis Unique front end. Initialization computes one full-slice
 // hash per axis index and clears the open-addressed table. Build inserts
 // exact-equality buckets and atomically retains the earliest representative.
 // Hash collisions always use the exact slice comparator, so hashes never
 // affect semantics. Floats hash and compare canonical FTZ-safe bits (one NaN
 // class, -0 folded to +0), so bucket equality is transitive and exact.
+{% set usesF16 = usesF16 and stage != "mark" %}
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const TABLE_SIZE: u32 = {{ tableSize }}u;
+{% if stage == "build" %}
 const TABLE_MASK: u32 = TABLE_SIZE - 1u;
 {% endif %}
 const EMPTY: u32 = 0xffffffffu;
+{% if stage == "init" %}
 {% if isFloat %}
 fn float_bits(v: {{ scalar }}) -> u32 {
 {% if usesF16 %}
   // FNV-1a over the complete slice. The axis index is deliberately excluded:
   // equal slices must hash identically regardless of where they occur.
   var hash = 0x811c9dc5u;
+{% if trackNaN %}
+  var containsNaN = false;
+{% endif %}
   for (var o = 0u; o < params.outer; o += 1u) {
     for (var n = 0u; n < params.inner; n += 1u) {
+{% if trackNaN %}
+      let value = slice_at(o, k, n);
+      hash = (hash ^ value_key(value)) * 0x01000193u;
+      let bits = float_bits(value);
+      containsNaN = containsNaN || ((bits & {{ "0x7c00u" if usesF16 else "0x7f800000u" }}) == {{ "0x7c00u" if usesF16 else "0x7f800000u" }} && (bits & {{ "0x03ffu" if usesF16 else "0x007fffffu" }}) != 0u);
+{% else %}
       hash = (hash ^ value_key(slice_at(o, k, n))) * 0x01000193u;
+{% endif %}
     }
   }
   // One final avalanche reduces clustering for short structured slices.
   hash ^= hash >> 16u;
   hash *= 0x7feb352du;
   hash ^= hash >> 15u;
+{% if trackNaN %}
+  // Each invocation owns its slice flag. Build consumes and clears it in the
+  // next dispatch before hash mark uses this same allocation for first flags.
+  firstFlag[k] = select(0u, 1u, containsNaN);
+{% endif %}
   return hash;
 }
 @compute @workgroup_size(WG)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let i = gid.x;
+{% if trackNaN %}
+  // No invocation ORs this flag until the following build dispatch.
+  if (i == 0u) { atomicStore(&nanFlag[0], 0u); }
+{% endif %}
   if (i < TABLE_SIZE) {
     atomicStore(&hashSlot[i], EMPTY);
   }
   if (i < params.axisDim) {
+{% if not trackNaN %}
     firstFlag[i] = 0u;
+{% endif %}
     sliceHash[i] = slice_hash(i);
   }
 }
+{% elif stage == "build" %}
 {% if isFloat %}
 fn float_bits(v: {{ scalar }}) -> u32 {
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let k = gid.x;
   if (k >= params.axisDim) { return; }
+{% if trackNaN %}
+  let containsNaN = firstFlag[k] != 0u;
+  firstFlag[k] = 0u;
+  if (containsNaN) { atomicOr(&nanFlag[0], 1u); }
+{% endif %}
   let key = sliceHash[k];
   var slot = key & TABLE_MASK;
   for (var probe = 0u; probe < TABLE_SIZE; probe += 1u) {

build/webgpu/unique-axis-scalar-inverse.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,109 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Open-addressing table constants and key/hash helpers.
+const EMPTY: u32 = 0xffffffffu;
+const MASK: u32 = {{ tableSize }}u - 1u;
+{% if isFloat %}
+// IEEE total-order comparators for Unique. Floating-point equality and ordering
+// use raw bits because GPUs may flush subnormals in float comparisons, which
+// would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
+// ordering uses a monotonic u32 key over the real line. Integers use == and <.
+{% macro float_bits_def() %}
+fn float_bits(v: {{ scalar }}) -> u32 {
+{% if usesF16 %}
+  // WGSL has no scalar u16 type. Packing v into the low component preserves
+  // its binary16 representation while producing a bitcast-compatible 32 bits.
+  return bitcast<u32>(vec2<f16>(v, 0.0h)) & 0xffffu;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+{%- endmacro -%}
+{%- macro is_nan_bits_def() %}
+fn is_nan_bits(v: {{ scalar }}) -> bool {
+  let b = float_bits(v);
+{% if usesF16 %}
+  return (b & 0x7c00u) == 0x7c00u && (b & 0x03ffu) != 0u;
+{% else %}
+  return (b & 0x7f800000u) == 0x7f800000u && (b & 0x007fffffu) != 0u;
+{% endif %}
+}
+{%- endmacro %}
+{{ float_bits_def() }}
+{{ is_nan_bits_def() }}
+{% endif %}
+fn key_bits(v: {{ scalar }}) -> u32 {
+{% if isFloat %}
+  let bits = float_bits(v);
+  return select(bits, 0u, bits == {{ "0x8000u" if usesF16 else "0x80000000u" }});
+{% elif isUnsigned %}
+  return v;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+fn hash_key(k: u32) -> u32 {
+  var x = k;
+  x = x ^ (x >> 16u);
+  x = x * 0x7feb352du;
+  x = x ^ (x >> 15u);
+  x = x * 0x846ca68bu;
+  x = x ^ (x >> 16u);
+  return x;
+}
+{% if not sorted %}
+fn ordered_key(key: u32) -> u32 {
+{% if usesF16 %}
+  return select(key ^ 0x8000u, (~key) & 0xffffu, (key & 0x8000u) != 0u);
+{% else %}
+  return select(key ^ 0x80000000u, ~key, (key & 0x80000000u) != 0u);
+{% endif %}
+}
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= {{ axisDim }}u) { return; }
+  if (is_nan_bits(x[0])) { inverse_indices[i] = 0u; return; }
+  // Later NaNs joined the smallest representative that already existed.
+  // Scan distinct representatives, excluding ones first seen after this input.
+  // Sorted slots can stop at their first eligible representative; unsorted
+  // slots appear in first-occurrence order, so later slots can be excluded.
+  if (is_nan_bits(x[i])) {
+    let written = min(count[0], {{ axisOutputDim }}u);
+{% if not sorted %}
+    var best = 0u;
+    var bestKey = ordered_key(key_bits(x[slots[0]]));
+{% endif %}
+    for (var p = 0u; p < written; p += 1u) {
+      let first = slots[p];
+{% if sorted %}
+      if (first <= i) { inverse_indices[i] = p; return; }
+{% else %}
+      if (first > i) { break; }
+      let candidate = ordered_key(key_bits(x[first]));
+      if (candidate < bestKey) { best = p; bestKey = candidate; }
+{% endif %}
+    }
+    inverse_indices[i] = {{ "0u" if sorted else "best" }};
+    return;
+  }
+  let key = key_bits(x[i]);
+  var slot = hash_key(key) & MASK;
+  loop {
+    let stored = tableKey[slot];
+    if (stored == key) {
+      inverse_indices[i] = rankOfFirst[tableIdx[slot]];
+      return;
+    }
+    if (stored == EMPTY) { inverse_indices[i] = 0u; return; }
+    slot = (slot + 1u) & MASK;
+  }
+}

build/webgpu/unique-axis-scalar-ranks.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,7 @@

+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let p = gid.x;
+  if (p >= min(count[0], {{ axisOutputDim }}u)) { return; }
+  rankOfFirst[slots[p]] = p;
+}

build/webgpu/unique-axis-scatter.wgsl.jinja CHANGED Viewed

@@ -1,17 +1,14 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // Grid-parallel output stage for axis Unique. Compact/sort must remain a single
-// globally synchronized workgroup, but copying its selected slice indices into a
-// wide Y does not: distribute that bandwidth-heavy tail across the device.
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
-const CAP: u32 = {{ source.capacity }}u;
-const AXIS_DIM: u32 = {{ source.axisDim }}u;
-const INNER: u32 = {{ source.inner }}u;
-const OUTPUT_AXIS_DIM: u32 = {{ source.outputAxisDim }}u;
-const TOTAL_OUT: u32 = {{ source.totalOut }}u;
 fn slice_at(o: u32, k: u32, n: u32) -> {{ scalar }} {
   return x[(o * AXIS_DIM + k) * INNER + n];

 {{ env.wgsl.resourceDeclarations }}
 // Grid-parallel output stage for axis Unique. Compact/sort must remain a single
+// globally synchronized workgroup, but copying selected slice indices into a
+// wide Y can be distributed across independent workgroups.
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const CAP: u32 = {{ capacity }}u;
+const AXIS_DIM: u32 = {{ axisDimSpec }}u;
+const INNER: u32 = {{ innerSize }}u;
+const OUTPUT_AXIS_DIM: u32 = {{ outputAxisDim }}u;
+const TOTAL_OUT: u32 = {{ totalOut }}u;
 fn slice_at(o: u32, k: u32, n: u32) -> {{ scalar }} {
   return x[(o * AXIS_DIM + k) * INNER + n];

build/webgpu/unique-axis.wgsl.jinja CHANGED Viewed

@@ -1,3 +1,4 @@
 // IEEE total-order comparators for Unique. Floating-point equality and ordering
 // use raw bits because GPUs may flush subnormals in float comparisons, which
 // would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
@@ -55,12 +56,9 @@ fn zero_value() -> {{ scalar }} {
 }
 {%- endmacro %}
-{% set emitIndices = source.hasIndices | default(false) %}
-{% set emitInverseIndices = source.hasInverseIndices | default(false) %}
-{% set emitCounts = source.hasCounts | default(false) %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // Axis-mode Unique views the input as outer x axisDim x inner. Two axis slices
@@ -68,14 +66,12 @@ enable f16;
 // each distinct slice is retained, optionally sorted lexicographically, then
 // scattered into the data-dependent output allocation.
 //
-// Float equality and ordering use the same bit-preserving rules as the other
-// Unique kernels, including canonicalized signed zero and FTZ-safe subnormals.
-const CAP: u32 = {{ source.capacity }}u;
-// The walk is one sequential step per slice and every step carries several
-// workgroup barriers, so this width is barrier latency, not throughput: the
-// head test settles in one round on real data and the chunked scan and shift
-// both adapt to any width. Tuned separately from the op-wide WORKGROUP_SIZE.
 const WG: u32 = {{ axisSerialWg }}u;
 {% if isFloat %}
@@ -88,7 +84,7 @@ var<workgroup> hitIndex: u32;
 var<workgroup> diffAt: atomic<u32>;
 var<workgroup> cmpOut: u32;
 {% endif %}
-{% set headCache = isFloat and (source.headCacheSlots | default(0)) > 0 %}
 {% set floatSliceLess = isFloat and (emitInverseIndices or emitCounts) %}
 {% set headOf = "headCache[candidate]" if headCache else "head_bits(order[candidate])" %}
 {% set signBit = "0x8000u" if usesF16 else "0x80000000u" %}
@@ -114,16 +110,10 @@ fn less_head(ba: u32, bb: u32) -> bool {
 {% endif %}
 {% if headCache %}
-// Leading element of each written representative, kept beside `order`. The scan
-// below compares one slice against every representative, and reading those
-// leading elements out of X is a gather whose stride is the slice width, which
-// can be a kilobyte, so every candidate lands on its own cache line. Widening
-// the workgroup makes that worse rather than better, which is the tell: the
-// limit is cache capacity, not latency. Caching the leading elements densely
-// means the common case, where two slices already differ at element zero, never
-// touches X. The integer dedup path solves the same problem the same way, with
-// sampled slice fingerprints.
-var<workgroup> headCache: array<u32, {{ source.headCacheSlots }}>;
 {% endif %}
 {% if isFloat %}
@@ -154,12 +144,9 @@ fn head_qualifies(hc: u32, hk: u32) -> bool {
 // equivalent (identical, or decided at a coordinate holding a NaN, which
 // `slice_less` reports as neither-less), 1 when a < b, 2 when b < a.
 //
-// The serial `slice_less` walks a slice one element at a time on a single lane,
-// and its exit is data dependent, so a pair of EQUAL slices -- what a duplicate
-// input produces, and the reason this comparison exists -- runs the full slice
-// length at one memory latency per element, and that walk dominates the kernel.
-// Here every lane takes one coordinate and the first difference falls out of a
-// min-reduction, so an equal pair costs one round instead of one per element.
 fn slice_cmp_coop(a: u32, b: u32, lane: u32) -> u32 {
   if (lane == 0u) { atomicStore(&diffAt, NO_HIT); }
   workgroupBarrier();
@@ -225,11 +212,8 @@ fn slice_less(a: u32, b: u32) -> bool {
     let baseA = (o * params.axisDim + a) * params.inner;
     let baseB = (o * params.axisDim + b) * params.inner;
     for (var n = 0u; n < params.inner; n = n + 1u) {
-      // Same three tests as `is_nan_bits` / `eq_value` / `less_value`, over one
-      // pair of bit patterns instead of six independent `float_bits` calls --
-      // which on f16 is a vec2 pack and mask each time. Two equal slices, the
-      // case a duplicate input hits, walk to the end of the slice, so this loop
-      // body runs once per element and was the kernel's dominant cost.
       let ba = float_bits(x[baseA + n]);
       let bb = float_bits(x[baseB + n]);
       if (is_nan_head(ba) || is_nan_head(bb)) { return false; }
@@ -328,19 +312,22 @@ fn find_bucket(k: u32, written: u32) -> u32 {
 {% endif %}
 @compute @workgroup_size(WG)
 fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
 {% if isFloat %}
   // Reproduce ordered lower-bound insertion. NaN makes the comparator
   // non-transitive, so pairwise equality and conventional hash tables cannot
   // preserve its stateful bucket behavior.
   //
-  // Every representative written so far is a lower-bound candidate, so the scan
-  // over them is the whole cost, and each step is a strided pair of slice loads
-  // that pays full memory latency. Scanning on one lane left the rest of the
-  // workgroup idle and issued those loads one at a time. Here the workgroup
-  // tests a chunk of candidates at once and keeps the SMALLEST index whose
-  // comparison fails -- the index the serial scan stops at, since it stops at
-  // the first failure too. Testing the remainder of a chunk is redundant, never
-  // observable: `slice_less` reads x and writes nothing.
   //
   // Both inner loops are barrier-free, so their bounds may depend on the live
   // count; every barrier sits at the top level of this loop, whose bound comes
@@ -349,10 +336,8 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
   for (var k = 0u; k < params.axisDim; k = k + 1u) {
     let headK = head_bits(k);
     // Walk the representatives for the first one that is not less than this
-    // slice. The head-only test is exact whenever the leading elements differ,
-    // so on real data this settles in a single round; a leading-element tie is
-    // resolved by one exact comparison, and the scan resumes past it only if
-    // that representative really does sort first.
     var searchFrom = 0u;
     var lowerBound = written;
     var equivalent = false;
@@ -361,8 +346,8 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
       workgroupBarrier();
       for (var base = searchFrom; base < written; base = base + WG) {
         // Chunks ascend, so an index already recorded is smaller than anything
-        // this chunk holds. Reading the atomic without a barrier can only see
-        // the hit late, which costs a redundant chunk and never a wrong bound.
         if (atomicLoad(&lowerBoundHit) != NO_HIT) { break; }
         let candidate = base + lid.x;
         if (candidate >= searchFrom && candidate < written
@@ -447,9 +432,8 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
 {% endif %}
 {% else %}
   if (params.axisDim <= CAP) {
-    // When every input slice can be represented in the order scratch, test
-    // first-occurrence status independently. Duplicate-heavy inputs usually
-    // exit after only a handful of comparisons, and all lanes participate.
     for (var k = lid.x; k < params.axisDim; k = k + WG) {
       var seen = false;
       for (var j = 0u; j < k; j = j + 1u) {
@@ -502,8 +486,8 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
   // `order` is storage-backed so legal data-dependent capacities are not
   // constrained by maxComputeWorkgroupStorageSize.
   storageBarrier();
-  // The public contract requires the caller-provided Y axis dimension to equal
-  // the exact unique count, so it is also the cross-lane written count.
   let writtenCount = params.outputAxisDim;
   // Scattering the selected slices is independent once `order` is ready, so

+{% set nanFallback = nanFallback | default(false) %}
 // IEEE total-order comparators for Unique. Floating-point equality and ordering
 // use raw bits because GPUs may flush subnormals in float comparisons, which
 // would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
 }
 {%- endmacro %}
+{% set emitIndices = hasIndices | default(false) %}
+{% set emitInverseIndices = hasInverseIndices | default(false) %}
+{% set emitCounts = hasCounts | default(false) %}
 {{ env.wgsl.resourceDeclarations }}
 // Axis-mode Unique views the input as outer x axisDim x inner. Two axis slices
 // each distinct slice is retained, optionally sorted lexicographically, then
 // scattered into the data-dependent output allocation.
 //
+// Float comparisons canonicalize signed zero and preserve subnormal magnitudes
+// when testing equality and lexicographic order.
+const CAP: u32 = {{ capacity }}u;
+// The slice walk, chunked scan, and shift all use this independently tunable
+// workgroup width.
 const WG: u32 = {{ axisSerialWg }}u;
 {% if isFloat %}
 var<workgroup> diffAt: atomic<u32>;
 var<workgroup> cmpOut: u32;
 {% endif %}
+{% set headCache = isFloat and (headCacheSlots | default(0)) > 0 %}
 {% set floatSliceLess = isFloat and (emitInverseIndices or emitCounts) %}
 {% set headOf = "headCache[candidate]" if headCache else "head_bits(order[candidate])" %}
 {% set signBit = "0x8000u" if usesF16 else "0x80000000u" %}
 {% endif %}
 {% if headCache %}
+// Canonicalized leading-element bits for each written representative, kept
+// beside `order`. A different leading value proves that two slices differ;
+// equal leading values fall through to the exact full-slice comparison.
+var<workgroup> headCache: array<u32, {{ headCacheSlots }}>;
 {% endif %}
 {% if isFloat %}
 // equivalent (identical, or decided at a coordinate holding a NaN, which
 // `slice_less` reports as neither-less), 1 when a < b, 2 when b < a.
 //
+// Lanes compare separate coordinates and an atomic minimum identifies the first
+// difference, preserving lexicographic order. Equal slices examine every
+// coordinate and return zero.
 fn slice_cmp_coop(a: u32, b: u32, lane: u32) -> u32 {
   if (lane == 0u) { atomicStore(&diffAt, NO_HIT); }
   workgroupBarrier();
     let baseA = (o * params.axisDim + a) * params.inner;
     let baseB = (o * params.axisDim + b) * params.inner;
     for (var n = 0u; n < params.inner; n = n + 1u) {
+      // Load each bit pattern once so the NaN, equality, and ordering tests use
+      // the same representation.
       let ba = float_bits(x[baseA + n]);
       let bb = float_bits(x[baseB + n]);
       if (is_nan_head(ba) || is_nan_head(bb)) { return false; }
 {% endif %}
 @compute @workgroup_size(WG)
 fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
+{% if nanFallback %}
+  // Publish through workgroup memory so the early return is uniform before
+  // entering the cooperative ordered insertion's barrier-containing loops.
+  if (lid.x == 0u) { hitIndex = nanFlag[0]; }
+  let requiresOrdered = workgroupUniformLoad(&hitIndex);
+  if (requiresOrdered == 0u) { return; }
+{% endif %}
 {% if isFloat %}
   // Reproduce ordered lower-bound insertion. NaN makes the comparator
   // non-transitive, so pairwise equality and conventional hash tables cannot
   // preserve its stateful bucket behavior.
   //
+  // Every prior representative is a lower-bound candidate. The workgroup tests
+  // candidates in chunks and retains the smallest index whose comparison fails,
+  // matching the stopping point of a serial lower-bound scan. Comparisons after
+  // that candidate cannot affect output because `slice_less` has no side effects.
   //
   // Both inner loops are barrier-free, so their bounds may depend on the live
   // count; every barrier sits at the top level of this loop, whose bound comes
   for (var k = 0u; k < params.axisDim; k = k + 1u) {
     let headK = head_bits(k);
     // Walk the representatives for the first one that is not less than this
+    // slice. Different leading elements settle the ordering directly; equal
+    // leading elements require an exact comparison.
     var searchFrom = 0u;
     var lowerBound = written;
     var equivalent = false;
       workgroupBarrier();
       for (var base = searchFrom; base < written; base = base + WG) {
         // Chunks ascend, so an index already recorded is smaller than anything
+        // in this chunk. A late atomic observation may repeat work but cannot
+        // change the selected bound.
         if (atomicLoad(&lowerBoundHit) != NO_HIT) { break; }
         let candidate = base + lid.x;
         if (candidate >= searchFrom && candidate < written
 {% endif %}
 {% else %}
   if (params.axisDim <= CAP) {
+    // When every input slice fits in the order scratch, lanes independently
+    // test whether each slice matches an earlier slice.
     for (var k = lid.x; k < params.axisDim; k = k + WG) {
       var seen = false;
       for (var j = 0u; j < k; j = j + 1u) {
   // `order` is storage-backed so legal data-dependent capacities are not
   // constrained by maxComputeWorkgroupStorageSize.
   storageBarrier();
+  // The output axis dimension equals the exact unique count and therefore also
+  // supplies the cross-lane written count.
   let writtenCount = params.outputAxisDim;
   // Scattering the selected slices is independent once `order` is ready, so

build/webgpu/unique-compact-sort.wgsl.jinja CHANGED Viewed

@@ -1,19 +1,26 @@
 {{ env.wgsl.resourceDeclarations }}
 // Pass 2 of parallel Unique: compact flagged first occurrences in appearance
 // order, optionally bitonic-sort them, and write the exact result. Small results
 // keep scratch in workgroup memory, while larger results use storage buffers
 // selected from device limits.
-const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
-const CAP: u32 = {{ source.capacity }}u;
 {% if sorted %}
 // The power-of-two padded tail sorts after every real value.
-const SORT_N: u32 = {{ source.sortN }}u;
 {% endif %}
 var<workgroup> wgScan: array<u32, WG>;
 var<workgroup> wgCarry: u32;
-{% if not source.globalScratch %}
 {% if sorted %}
 var<workgroup> sortKey: array<u32, SORT_N>;
 var<workgroup> sortVal: array<{{ scalar }}, SORT_N>;
@@ -23,20 +30,46 @@ var<workgroup> compacted: array<{{ scalar }}, CAP>;
 {% endif %}
 {% endif %}
 fn zero_value() -> {{ scalar }} {
   return {{ scalar }}(0);
 }
 {% if sorted %}
-// Monotonic key matching unsigned or signed integer order. Global scratch
-// already stores raw bits.
-fn sort_key(v: {% if source.globalScratch %}u32{% else %}{{ scalar }}{% endif %}) -> u32 {
-{% if source.globalScratch %}
   let b = v;
 {% else %}
   let b = bitcast<u32>(v);
 {% endif %}
-{% if isUnsigned %}
   return b;
 {% else %}
   return b ^ 0x80000000u;
@@ -45,20 +78,11 @@ fn sort_key(v: {% if source.globalScratch %}u32{% else %}{{ scalar }}{% endif %}
 {% endif %}
 {% macro scratch_barrier() %}
-{% if source.globalScratch %}storageBarrier();{% else %}workgroupBarrier();{% endif %}
 {% endmacro %}
-@compute @workgroup_size(WG)
-fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
-  let tid = lid.x;
-  if (tid == 0u) {
-    wgCarry = 0u;
-  }
-  workgroupBarrier();
-  // Chunked Hillis-Steele exclusive scan maps flags to appearance-order ranks.
-  let chunks = (params.inputCount + WG - 1u) / WG;
-  for (var c = 0u; c < chunks; c = c + 1u) {
     let i = c * WG + tid;
     var f = 0u;
     if (i < params.inputCount) {
@@ -81,7 +105,7 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
     let excl = wgScan[tid] - f;
     let pos = wgCarry + excl;
     if (i < params.inputCount && f == 1u && pos < CAP) {
-{% if source.globalScratch %}
       sortVal[pos] = bitcast<u32>(x[i]);
 {% elif sorted %}
       sortVal[pos] = x[i];
@@ -94,11 +118,39 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
       wgCarry = wgCarry + wgScan[WG - 1u];
     }
     workgroupBarrier();
   }
   let total = wgCarry;
   let written = min(total, CAP);
-{% if source.globalScratch %}
   storageBarrier();
 {% endif %}
@@ -111,7 +163,7 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
     if (isPad == 0u) {
       sortKey[k] = sort_key(sortVal[k]);
     } else {
-{% if source.globalScratch %}
       sortVal[k] = 0u;
 {% else %}
       sortVal[k] = zero_value();
@@ -121,6 +173,7 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
   }
   {{ scratch_barrier() }}
   // Batcher bitonic network. Each lane owns one side of a compare-exchange.
   var size = 2u;
   loop {
@@ -158,12 +211,14 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
   }
 {% endif %}
-{% if not source.globalScratch %}
   workgroupBarrier();
 {% endif %}
   for (var k = tid; k < CAP; k = k + WG) {
     if (k < written) {
-{% if source.globalScratch %}
       y[k] = bitcast<{{ scalar }}>(sortVal[k]);
 {% elif sorted %}
       y[k] = sortVal[k];
@@ -174,4 +229,5 @@ fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
       y[k] = zero_value();
     }
   }
 }

+{% if usesF16 %}
+enable f16;
+{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // Pass 2 of parallel Unique: compact flagged first occurrences in appearance
 // order, optionally bitonic-sort them, and write the exact result. Small results
 // keep scratch in workgroup memory, while larger results use storage buffers
 // selected from device limits.
+// Global-scratch variants reserve workgroup memory only for the scan buffer and
+// provide an independent sort width. Workgroup-scratch variants use the shared
+// workgroup-width setting.
+{% set compactWg = sortWg | default(tunables.WORKGROUP_SIZE) %}
+const WG: u32 = {{ compactWg }}u;
+const CAP: u32 = {{ capacity }}u;
 {% if sorted %}
 // The power-of-two padded tail sorts after every real value.
+const SORT_N: u32 = {{ sortN }}u;
 {% endif %}
 var<workgroup> wgScan: array<u32, WG>;
 var<workgroup> wgCarry: u32;
+{% if not globalScratch %}
 {% if sorted %}
 var<workgroup> sortKey: array<u32, SORT_N>;
 var<workgroup> sortVal: array<{{ scalar }}, SORT_N>;
 {% endif %}
 {% endif %}
+{% if sortExternally is not defined or not sortExternally %}
 fn zero_value() -> {{ scalar }} {
   return {{ scalar }}(0);
 }
+{% endif %}
 {% if sorted %}
+{% if isFloat and not globalScratch %}
+// IEEE total-order comparators for Unique. Floating-point equality and ordering
+// use raw bits because GPUs may flush subnormals in float comparisons, which
+// would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
+// ordering uses a monotonic u32 key over the real line. Integers use == and <.
+{% macro float_bits_def() %}
+fn float_bits(v: {{ scalar }}) -> u32 {
+{% if usesF16 %}
+  // WGSL has no scalar u16 type. Packing v into the low component preserves
+  // its binary16 representation while producing a bitcast-compatible 32 bits.
+  return bitcast<u32>(vec2<f16>(v, 0.0h)) & 0xffffu;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+{%- endmacro %}
+{{ float_bits_def() }}
+{% endif %}
+// Monotonic raw-bit key preserves finite subnormals on FTZ devices. Global
+// scratch already stores raw bits.
+fn sort_key(v: {% if globalScratch %}u32{% else %}{{ scalar }}{% endif %}) -> u32 {
+{% if globalScratch %}
   let b = v;
+{% elif isFloat %}
+  let b = float_bits(v);
 {% else %}
   let b = bitcast<u32>(v);
 {% endif %}
+{% if isFloat %}
+  return select(b | {{ "0x8000u" if usesF16 else "0x80000000u" }}, (~b) & {{ "0xffffu" if usesF16 else "0xffffffffu" }}, (b & {{ "0x8000u" if usesF16 else "0x80000000u" }}) != 0u);
+{% elif isUnsigned %}
   return b;
 {% else %}
   return b ^ 0x80000000u;
 {% endif %}
 {% macro scratch_barrier() %}
+{% if globalScratch %}storageBarrier();{% else %}workgroupBarrier();{% endif %}
 {% endmacro %}
+// Compact one input chunk while all lanes keep the scan and carry coherent.
+fn compact_chunk(tid: u32, c: u32) {
     let i = c * WG + tid;
     var f = 0u;
     if (i < params.inputCount) {
     let excl = wgScan[tid] - f;
     let pos = wgCarry + excl;
     if (i < params.inputCount && f == 1u && pos < CAP) {
+{% if globalScratch %}
       sortVal[pos] = bitcast<u32>(x[i]);
 {% elif sorted %}
       sortVal[pos] = x[i];
       wgCarry = wgCarry + wgScan[WG - 1u];
     }
     workgroupBarrier();
+}
+@compute @workgroup_size(WG)
+fn main(@builtin(local_invocation_id) lid: vec3<u32>) {
+  let tid = lid.x;
+  if (tid == 0u) {
+    wgCarry = 0u;
+  }
+  workgroupBarrier();
+  let chunks = (params.inputCount + WG - 1u) / WG;
+{% if capacity <= compactWg %}
+  // Test completion once, after the first chunk. If representatives occur
+  // later, the remaining scan retains its original barrier-only loop body.
+  if (chunks > 0u) {
+    compact_chunk(tid, 0u);
+    let first_count = workgroupUniformLoad(&wgCarry);
+    if (first_count < CAP) {
+      for (var c = 1u; c < chunks; c = c + 1u) {
+        compact_chunk(tid, c);
+      }
+    }
   }
+{% else %}
+  // A chunk cannot fill a capacity larger than its invocation count.
+  for (var c = 0u; c < chunks; c = c + 1u) {
+    compact_chunk(tid, c);
+  }
+{% endif %}
   let total = wgCarry;
   let written = min(total, CAP);
+{% if globalScratch %}
   storageBarrier();
 {% endif %}
     if (isPad == 0u) {
       sortKey[k] = sort_key(sortVal[k]);
     } else {
+{% if globalScratch %}
       sortVal[k] = 0u;
 {% else %}
       sortVal[k] = zero_value();
   }
   {{ scratch_barrier() }}
+{% if sortExternally is not defined or not sortExternally %}
   // Batcher bitonic network. Each lane owns one side of a compare-exchange.
   var size = 2u;
   loop {
   }
 {% endif %}
+{% endif %}
+{% if sortExternally is not defined or not sortExternally %}
+{% if not globalScratch %}
   workgroupBarrier();
 {% endif %}
   for (var k = tid; k < CAP; k = k + WG) {
     if (k < written) {
+{% if globalScratch %}
       y[k] = bitcast<{{ scalar }}>(sortVal[k]);
 {% elif sorted %}
       y[k] = sortVal[k];
       y[k] = zero_value();
     }
   }
+{% endif %}
 }

build/webgpu/unique-flag-block-scan.wgsl.jinja ADDED Viewed

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

+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+var<workgroup> scan: array<u32, WG>;
+@compute @workgroup_size(WG)
+fn main(@builtin(local_invocation_id) lid: vec3<u32>, @builtin(global_invocation_id) gid: vec3<u32>, @builtin(workgroup_id) wid: vec3<u32>) {
+  let tid = lid.x;
+  var flag = 0u;
+  if (gid.x < {{ scanN }}u) { flag = flags[gid.x]; }
+  scan[tid] = flag;
+  for (var step = 1u; step < WG; step <<= 1u) {
+    workgroupBarrier();
+    var previous = 0u;
+    if (tid >= step) { previous = scan[tid - step]; }
+    workgroupBarrier();
+    scan[tid] += previous;
+  }
+  workgroupBarrier();
+  if (gid.x < {{ scanN }}u) { flags[gid.x] = 2u * (scan[tid] - flag) + flag; }
+  if (tid == WG - 1u) { blockSums[wid.x] = scan[tid]; }
+}

build/webgpu/unique-flat-metadata.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,129 @@

+// IEEE total-order comparators for Unique. Floating-point equality and ordering
+// use raw bits because GPUs may flush subnormals in float comparisons, which
+// would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
+// ordering uses a monotonic u32 key over the real line. Integers use == and <.
+{% macro eq_value_def() %}
+fn eq_value(a: {{ scalar }}, b: {{ scalar }}) -> bool {
+{% if isFloat %}
+  var ba = float_bits(a);
+  var bb = float_bits(b);
+{% if usesF16 %}
+  if (ba == 0x8000u) { ba = 0u; }  // -0 -> +0
+  if (bb == 0x8000u) { bb = 0u; }
+{% else %}
+  if (ba == 0x80000000u) { ba = 0u; }  // -0 -> +0
+  if (bb == 0x80000000u) { bb = 0u; }
+{% endif %}
+  return ba == bb;
+{% else %}
+  return a == b;
+{% endif %}
+}
+{%- endmacro -%}
+{%- macro less_value_def() %}
+fn less_value(a: {{ scalar }}, b: {{ scalar }}) -> bool {
+{% if isFloat %}
+  let ba = float_bits(a);
+  let bb = float_bits(b);
+{% if usesF16 %}
+  let ka = select(ba | 0x8000u, (~ba) & 0xffffu, (ba & 0x8000u) != 0u);
+  let kb = select(bb | 0x8000u, (~bb) & 0xffffu, (bb & 0x8000u) != 0u);
+{% else %}
+  let ka = select(ba | 0x80000000u, ~ba, (ba & 0x80000000u) != 0u);
+  let kb = select(bb | 0x80000000u, ~bb, (bb & 0x80000000u) != 0u);
+{% endif %}
+  return ka < kb;
+{% else %}
+  return a < b;
+{% endif %}
+}
+{%- endmacro %}
+{% set emitIndices = hasIndices | default(false) %}
+{% set emitInverseIndices = hasInverseIndices | default(false) %}
+{{ env.wgsl.resourceDeclarations }}
+// Metadata for the parallel integer Unique routes. Y and the first-occurrence
+// flags are already final when this runs, so every output is a pure function of
+// (x, y, flags) and one thread can own one element. The single-invocation kernel
+// this replaces resolved the same buckets by scanning the whole input once per
+// (input, output) pair.
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+{% if stage != "scatter" or not sorted %}
+{{ eq_value_def() }}
+{% endif %}
+{% if stage == "scatter" %}
+{% if sorted %}
+{{ less_value_def() }}
+// Y is ascending on this route, so a bucket is a binary search rather than a
+// walk over every unique value.
+fn find_slot(value: {{ scalar }}) -> u32 {
+  var lo = 0u;
+  var hi = params.capacity;
+  loop {
+    if (lo >= hi) { break; }
+    let mid = lo + (hi - lo) / 2u;
+    if (less_value(y[mid], value)) {
+      lo = mid + 1u;
+    } else {
+      hi = mid;
+    }
+  }
+  return lo;
+}
+{% else %}
+// Y is in first-appearance order, so a bucket is a scan of the unique values.
+fn find_slot(value: {{ scalar }}) -> u32 {
+  for (var k = 0u; k < params.capacity; k = k + 1u) {
+    if (eq_value(y[k], value)) { return k; }
+  }
+  return params.capacity;
+}
+{% endif %}
+@compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= params.inputCount) {
+    return;
+  }
+  let slot = find_slot(x[i]);
+  if (slot >= params.capacity) {
+    return;
+  }
+{% if emitInverseIndices %}
+  inverse_indices[i] = slot;
+{% endif %}
+{% if emitIndices %}
+  // Exactly one input element is the first occurrence of its value, so each
+  // slot has a single writer and the store needs no ordering.
+  if (flags[i] == 1u) {
+    indices[slot] = i;
+  }
+{% endif %}
+}
+{% else %}
+// One thread per unique value counts its own occurrences. Accumulating per
+// output rather than per input keeps the tally free of atomics, and so free of
+// the zero-fill pass an atomic accumulation would need first.
+@compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let k = gid.x;
+  if (k >= params.capacity) {
+    return;
+  }
+  let value = y[k];
+  var total = 0u;
+  for (var i = 0u; i < params.inputCount; i = i + 1u) {
+    if (eq_value(x[i], value)) {
+      total = total + 1u;
+    }
+  }
+  counts[k] = total;
+}
+{% endif %}

build/webgpu/unique-global-sort-exchange.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,36 @@

+// Integer Unique's existing (padding, monotonic-key) bitonic comparison.
+// One invocation owns a disjoint pair. A dispatch boundary publishes each
+// completed distance to the next distance across every workgroup.
+{{ env.wgsl.resourceDeclarations }}
+{% macro exchange_integer(keys, values, pads, writeState=true) %}
+  let pk = {{ pads }}[k];
+  let pp = {{ pads }}[partner];
+  let kk = {{ keys }}[k];
+  let kp = {{ keys }}[partner];
+  let kBeforeP = (pk < pp) || (pk == pp && kk <= kp);
+  let needSwap = select(kBeforeP, !kBeforeP, ascending);
+  if (needSwap) {
+    let tVal = {{ values }}[k];
+    {{ values }}[k] = {{ values }}[partner];
+    {{ values }}[partner] = tVal;
+{% if writeState %}
+    {{ keys }}[k] = kp;
+    {{ keys }}[partner] = kk;
+    {{ pads }}[k] = pp;
+    {{ pads }}[partner] = pk;
+{% endif %}
+  }
+{% endmacro %}
+const SIZE: u32 = {{ sortSize }}u;
+const STRIDE: u32 = {{ sortStride }}u;
+const PAIRS: u32 = {{ (sortLength / 2)|int }}u;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let pair = gid.x;
+  if (pair >= PAIRS) { return; }
+  let k = (pair / STRIDE) * (2u * STRIDE) + pair % STRIDE;
+  let partner = k + STRIDE;
+  let ascending = (k & SIZE) == 0u;
+{{ exchange_integer("sortKey", "sortVal", "sortPad", writeSortState | default(true)) }}
+}

build/webgpu/unique-global-sort-output.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,9 @@

+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let k = gid.x;
+  if (k >= {{ capacity }}u) { return; }
+  var value = {{ scalar }}(0);
+  if (k < {{ sortLength }}u && sortPad[k] == 0u) { value = bitcast<{{ scalar }}>(sortVal[k]); }
+  y[k] = value;
+}

build/webgpu/unique-global-sort-shared.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,69 @@

+// Integer-only bitonic stages whose pairs lie in one contiguous shared tile.
+// Global indices, including the tile base, retain the network's direction.
+{{ env.wgsl.resourceDeclarations }}
+{% macro exchange_integer(keys, values, pads, writeState=true) %}
+  let pk = {{ pads }}[k];
+  let pp = {{ pads }}[partner];
+  let kk = {{ keys }}[k];
+  let kp = {{ keys }}[partner];
+  let kBeforeP = (pk < pp) || (pk == pp && kk <= kp);
+  let needSwap = select(kBeforeP, !kBeforeP, ascending);
+  if (needSwap) {
+    let tVal = {{ values }}[k];
+    {{ values }}[k] = {{ values }}[partner];
+    {{ values }}[partner] = tVal;
+{% if writeState %}
+    {{ keys }}[k] = kp;
+    {{ keys }}[partner] = kk;
+    {{ pads }}[k] = pp;
+    {{ pads }}[partner] = pk;
+{% endif %}
+  }
+{% endmacro %}
+const TILE: u32 = {{ tileSize }}u;
+const WG: u32 = {{ workgroupSize }}u;
+var<workgroup> tileKey: array<u32, TILE>;
+var<workgroup> tileVal: array<u32, TILE>;
+var<workgroup> tilePad: array<u32, TILE>;
+fn exchange_pair(pair: u32, stride: u32, size: u32, base: u32) {
+  let k = (pair / stride) * (2u * stride) + pair % stride;
+  let partner = k + stride;
+  let ascending = ((base + k) & size) == 0u;
+{{ exchange_integer("tileKey", "tileVal", "tilePad") }}
+}
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(local_invocation_index) tid: u32) {
+  let base = wg.x * TILE;
+  for (var i = tid; i < TILE; i += WG) {
+    tileKey[i] = sortKey[base + i];
+    tileVal[i] = sortVal[base + i];
+    tilePad[i] = sortPad[base + i];
+  }
+  workgroupBarrier();
+{% if initialSort %}
+  for (var size = 2u; size <= TILE; size *= 2u) {
+    for (var stride = size / 2u; stride > 0u; stride /= 2u) {
+      for (var pair = tid; pair < TILE / 2u; pair += WG) {
+        exchange_pair(pair, stride, size, base);
+      }
+      workgroupBarrier();
+    }
+  }
+{% else %}
+  for (var stride = TILE / 2u; stride > 0u; stride /= 2u) {
+    for (var pair = tid; pair < TILE / 2u; pair += WG) {
+      exchange_pair(pair, stride, {{ sortSize }}u, base);
+    }
+    workgroupBarrier();
+  }
+{% endif %}
+  for (var i = tid; i < TILE; i += WG) {
+{% if writeSortState | default(true) %}
+    sortKey[base + i] = tileKey[i];
+    sortPad[base + i] = tilePad[i];
+{% endif %}
+    sortVal[base + i] = tileVal[i];
+  }
+}

build/webgpu/unique-hash-build.wgsl.jinja CHANGED Viewed

@@ -1,3 +1,6 @@
 {{ env.wgsl.resourceDeclarations }}
 // One thread per input element inserts its value into an open-addressing hash
@@ -15,10 +18,45 @@
 // its minimum index uses `special`.
 // Open-addressing table constants and key/hash helpers.
 const EMPTY: u32 = 0xffffffffu;
-const MASK: u32 = {{ source.tableSize }}u - 1u;
 fn key_bits(v: {{ scalar }}) -> u32 {
-{% if isUnsigned %}
   return v;
 {% else %}
   return bitcast<u32>(v);
@@ -41,6 +79,11 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   if (i >= params.inputCount) {
     return;
   }
   let k = key_bits(x[i]);
   if (k == EMPTY) {
     atomicMin(&special[0], i);
@@ -50,12 +93,14 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   loop {
     let res = atomicCompareExchangeWeak(&tableKey[h], EMPTY, k);
     if (res.exchanged || res.old_value == k) {
-{% if not source.keyOnly %}
       // This slot now holds k (we just placed it, or it already held k).
       atomicMin(&tableIdx[h], i);
 {% endif %}
       break;
     }
-    h = (h + 1u) & MASK;
   }
 }

+{% if usesF16 %}
+enable f16;
+{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // One thread per input element inserts its value into an open-addressing hash
 // its minimum index uses `special`.
 // Open-addressing table constants and key/hash helpers.
 const EMPTY: u32 = 0xffffffffu;
+const MASK: u32 = {{ tableSize }}u - 1u;
+{% if isFloat %}
+// IEEE total-order comparators for Unique. Floating-point equality and ordering
+// use raw bits because GPUs may flush subnormals in float comparisons, which
+// would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
+// ordering uses a monotonic u32 key over the real line. Integers use == and <.
+{% macro float_bits_def() %}
+fn float_bits(v: {{ scalar }}) -> u32 {
+{% if usesF16 %}
+  // WGSL has no scalar u16 type. Packing v into the low component preserves
+  // its binary16 representation while producing a bitcast-compatible 32 bits.
+  return bitcast<u32>(vec2<f16>(v, 0.0h)) & 0xffffu;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+{%- endmacro -%}
+{%- macro is_nan_bits_def() %}
+fn is_nan_bits(v: {{ scalar }}) -> bool {
+  let b = float_bits(v);
+{% if usesF16 %}
+  return (b & 0x7c00u) == 0x7c00u && (b & 0x03ffu) != 0u;
+{% else %}
+  return (b & 0x7f800000u) == 0x7f800000u && (b & 0x007fffffu) != 0u;
+{% endif %}
+}
+{%- endmacro %}
+{{ float_bits_def() }}
+{{ is_nan_bits_def() }}
+{% endif %}
 fn key_bits(v: {{ scalar }}) -> u32 {
+{% if isFloat %}
+  let bits = float_bits(v);
+  return select(bits, 0u, bits == {{ "0x8000u" if usesF16 else "0x80000000u" }});
+{% elif isUnsigned %}
   return v;
 {% else %}
   return bitcast<u32>(v);
   if (i >= params.inputCount) {
     return;
   }
+{% if isFloat %}
+  // A leading NaN is the ordered map's only representative. Otherwise every
+  // NaN is equivalent to an already-present bucket and creates no new value.
+  if (is_nan_bits(x[0]) || is_nan_bits(x[i])) { return; }
+{% endif %}
   let k = key_bits(x[i]);
   if (k == EMPTY) {
     atomicMin(&special[0], i);
   loop {
     let res = atomicCompareExchangeWeak(&tableKey[h], EMPTY, k);
     if (res.exchanged || res.old_value == k) {
+{% if not keyOnly %}
       // This slot now holds k (we just placed it, or it already held k).
       atomicMin(&tableIdx[h], i);
 {% endif %}
       break;
     }
+    // Weak CAS may fail spuriously while the slot is still empty.
+    // Keep probing this slot until it is claimed or contains another key.
+    if (res.old_value != EMPTY) { h = (h + 1u) & MASK; }
   }
 }

build/webgpu/unique-hash-collect.wgsl.jinja CHANGED Viewed

@@ -1,4 +1,4 @@
-{% if source.useSubgroups %}
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
@@ -7,14 +7,14 @@ enable subgroups;
 // occupied hash keys directly and reserve output positions once per workgroup.
 // This avoids serializing every distinct value through one global atomic.
 const EMPTY: u32 = 0xffffffffu;
-const TABLE_SIZE: u32 = {{ source.tableSize }}u;
-const CAP: u32 = {{ source.capacity }}u;
-{% if source.useSubgroups %}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 {% endif %}
 var<workgroup> groupBase: u32;
-{% if source.useSubgroups %}
 // The first subgroup lane writes its total. Lane zero converts these totals to
 // subgroup offsets and performs the workgroup's sole global count reservation.
 var<workgroup> subgroupOffsets: array<u32, WG>;
@@ -26,7 +26,7 @@ var<workgroup> localCount: atomic<u32>;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>,
-        @builtin(local_invocation_id) lid: vec3<u32>{% if source.useSubgroups %},
         @builtin(subgroup_invocation_id) sgLane: u32,
         @builtin(subgroup_size) sgSize: u32{% endif %}) {
   let h = gid.x;
@@ -54,7 +54,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>,
     }
   }
-{% if source.useSubgroups %}
   let subgroupPrefix = subgroupExclusiveAdd(itemCount);
   let subgroupTotal = subgroupAdd(itemCount);
   let subgroupId = lid.x / sgSize;

+{% if useSubgroups %}
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // occupied hash keys directly and reserve output positions once per workgroup.
 // This avoids serializing every distinct value through one global atomic.
 const EMPTY: u32 = 0xffffffffu;
+const TABLE_SIZE: u32 = {{ tableSize }}u;
+const CAP: u32 = {{ capacity }}u;
+{% if useSubgroups %}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 {% endif %}
 var<workgroup> groupBase: u32;
+{% if useSubgroups %}
 // The first subgroup lane writes its total. Lane zero converts these totals to
 // subgroup offsets and performs the workgroup's sole global count reservation.
 var<workgroup> subgroupOffsets: array<u32, WG>;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>{% if useSubgroups %},
         @builtin(subgroup_invocation_id) sgLane: u32,
         @builtin(subgroup_size) sgSize: u32{% endif %}) {
   let h = gid.x;
     }
   }
+{% if useSubgroups %}
   let subgroupPrefix = subgroupExclusiveAdd(itemCount);
   let subgroupTotal = subgroupAdd(itemCount);
   let subgroupId = lid.x / sgSize;

build/webgpu/unique-hash-init.wgsl.jinja CHANGED Viewed

@@ -6,11 +6,11 @@
 // folds each key's minimum input index into tableIdx, and it leaves the output
 // counter alone -- only the sorted-collect path counts from zero here.
 const EMPTY: u32 = 0xffffffffu;
-const TABLE_SIZE: u32 = {{ source.tableSize }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{% if source.keyOnlyVec4 %}
   let h4 = gid.x;
   if (h4 >= TABLE_SIZE / 4u) {
     return;

 // folds each key's minimum input index into tableIdx, and it leaves the output
 // counter alone -- only the sorted-collect path counts from zero here.
 const EMPTY: u32 = 0xffffffffu;
+const TABLE_SIZE: u32 = {{ tableSize }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{% if keyOnlyVec4 %}
   let h4 = gid.x;
   if (h4 >= TABLE_SIZE / 4u) {
     return;

build/webgpu/unique-hash-mark.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,9 @@
 {{ env.wgsl.resourceDeclarations }}
-// Pass 3 of the hash-set parallel Unique (integer dtypes only). One thread per
 // input element re-probes the (now fully built) hash table for its value and
 // reads the stored minimum index. flags[i] = 1 iff i is that minimum, i.e. i is
 // the first occurrence of its value — exactly the predicate the O(n^2) scan
@@ -9,10 +12,45 @@
 // hitting an EMPTY slot.
 // Open-addressing table constants and key/hash helpers.
 const EMPTY: u32 = 0xffffffffu;
-const MASK: u32 = {{ source.tableSize }}u - 1u;
 fn key_bits(v: {{ scalar }}) -> u32 {
-{% if isUnsigned %}
   return v;
 {% else %}
   return bitcast<u32>(v);
@@ -34,6 +72,17 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   if (i >= params.inputCount) {
     return;
   }
   let k = key_bits(x[i]);
   if (k == EMPTY) {
     flags[i] = select(0u, 1u, atomicLoad(&special[0]) == i);

+{% if usesF16 %}
+enable f16;
+{% endif %}
 {{ env.wgsl.resourceDeclarations }}
+// Pass 3 of the hash-set parallel Unique. One thread per
 // input element re-probes the (now fully built) hash table for its value and
 // reads the stored minimum index. flags[i] = 1 iff i is that minimum, i.e. i is
 // the first occurrence of its value — exactly the predicate the O(n^2) scan
 // hitting an EMPTY slot.
 // Open-addressing table constants and key/hash helpers.
 const EMPTY: u32 = 0xffffffffu;
+const MASK: u32 = {{ tableSize }}u - 1u;
+{% if isFloat %}
+// IEEE total-order comparators for Unique. Floating-point equality and ordering
+// use raw bits because GPUs may flush subnormals in float comparisons, which
+// would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
+// ordering uses a monotonic u32 key over the real line. Integers use == and <.
+{% macro float_bits_def() %}
+fn float_bits(v: {{ scalar }}) -> u32 {
+{% if usesF16 %}
+  // WGSL has no scalar u16 type. Packing v into the low component preserves
+  // its binary16 representation while producing a bitcast-compatible 32 bits.
+  return bitcast<u32>(vec2<f16>(v, 0.0h)) & 0xffffu;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+{%- endmacro -%}
+{%- macro is_nan_bits_def() %}
+fn is_nan_bits(v: {{ scalar }}) -> bool {
+  let b = float_bits(v);
+{% if usesF16 %}
+  return (b & 0x7c00u) == 0x7c00u && (b & 0x03ffu) != 0u;
+{% else %}
+  return (b & 0x7f800000u) == 0x7f800000u && (b & 0x007fffffu) != 0u;
+{% endif %}
+}
+{%- endmacro %}
+{{ float_bits_def() }}
+{{ is_nan_bits_def() }}
+{% endif %}
 fn key_bits(v: {{ scalar }}) -> u32 {
+{% if isFloat %}
+  let bits = float_bits(v);
+  return select(bits, 0u, bits == {{ "0x8000u" if usesF16 else "0x80000000u" }});
+{% elif isUnsigned %}
   return v;
 {% else %}
   return bitcast<u32>(v);
   if (i >= params.inputCount) {
     return;
   }
+{% if isFloat %}
+  // Preserve the leading NaN's exact payload through the original X value.
+  if (is_nan_bits(x[0])) {
+    flags[i] = select(0u, 1u, i == 0u);
+    return;
+  }
+  if (is_nan_bits(x[i])) {
+    flags[i] = 0u;
+    return;
+  }
+{% endif %}
   let k = key_bits(x[i]);
   if (k == EMPTY) {
     flags[i] = select(0u, 1u, atomicLoad(&special[0]) == i);

build/webgpu/unique-hash-sort-collected-key-only.wgsl.jinja CHANGED Viewed

@@ -1,8 +1,8 @@
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
-const CAP: u32 = {{ source.capacity }}u;
-const SORT_N: u32 = {{ source.sortN }}u;
 // The monotonic key transform is invertible, so the workgroup sort only needs
 // one u32 per slot. Padding uses the maximum key; if a real maximum key exists,

 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const CAP: u32 = {{ capacity }}u;
+const SORT_N: u32 = {{ sortN }}u;
 // The monotonic key transform is invertible, so the workgroup sort only needs
 // one u32 per slot. Padding uses the maximum key; if a real maximum key exists,

build/webgpu/unique-scalar-compact.wgsl.jinja ADDED Viewed

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

+{% if sorted and usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% if sorted %}
+{% if isFloat %}
+// IEEE total-order comparators for Unique. Floating-point equality and ordering
+// use raw bits because GPUs may flush subnormals in float comparisons, which
+// would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
+// ordering uses a monotonic u32 key over the real line. Integers use == and <.
+{% macro float_bits_def() %}
+fn float_bits(v: {{ scalar }}) -> u32 {
+{% if usesF16 %}
+  // WGSL has no scalar u16 type. Packing v into the low component preserves
+  // its binary16 representation while producing a bitcast-compatible 32 bits.
+  return bitcast<u32>(vec2<f16>(v, 0.0h)) & 0xffffu;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+{%- endmacro %}
+{{ float_bits_def() }}
+{% endif %}
+fn ordered_key(value: {{ scalar }}) -> u32 {
+{% if isFloat %}
+  let bits = float_bits(value);
+  return select(bits | {{ "0x8000u" if usesF16 else "0x80000000u" }}, (~bits) & {{ "0xffffu" if usesF16 else "0xffffffffu" }}, (bits & {{ "0x8000u" if usesF16 else "0x80000000u" }}) != 0u);
+{% elif isUnsigned %}
+  return value;
+{% else %}
+  return bitcast<u32>(value) ^ 0x80000000u;
+{% endif %}
+}
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i < {{ scanN }}u) {
+    let encoded = flags[i];
+    if ((encoded & 1u) != 0u) {
+      let position = blockPrefix[i / {{ tunables.WORKGROUP_SIZE }}u] + (encoded >> 1u);
+      if (position < {{ capacity }}u) {
+        sortVal[position] = i;
+{% if sorted %}
+        sortKey[position] = ordered_key(x[i]);
+        sortPad[position] = 0u;
+{% endif %}
+      }
+    }
+  }
+{% if sorted %}
+  // Exact output cardinality is part of Unique's required request contract.
+  if (i >= {{ capacity }}u && i < {{ sortN }}u) {
+    sortPad[i] = 1u;
+    sortKey[i] = 0u;
+    sortVal[i] = 0u;
+  }
+{% endif %}
+}

build/webgpu/unique-scalar-metadata.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,122 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Open-addressing table constants and key/hash helpers.
+const EMPTY: u32 = 0xffffffffu;
+const MASK: u32 = {{ tableSize }}u - 1u;
+{% if isFloat %}
+// IEEE total-order comparators for Unique. Floating-point equality and ordering
+// use raw bits because GPUs may flush subnormals in float comparisons, which
+// would otherwise collapse distinct values. Equality canonicalizes -0 to +0;
+// ordering uses a monotonic u32 key over the real line. Integers use == and <.
+{% macro float_bits_def() %}
+fn float_bits(v: {{ scalar }}) -> u32 {
+{% if usesF16 %}
+  // WGSL has no scalar u16 type. Packing v into the low component preserves
+  // its binary16 representation while producing a bitcast-compatible 32 bits.
+  return bitcast<u32>(vec2<f16>(v, 0.0h)) & 0xffffu;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+{%- endmacro -%}
+{%- macro is_nan_bits_def() %}
+fn is_nan_bits(v: {{ scalar }}) -> bool {
+  let b = float_bits(v);
+{% if usesF16 %}
+  return (b & 0x7c00u) == 0x7c00u && (b & 0x03ffu) != 0u;
+{% else %}
+  return (b & 0x7f800000u) == 0x7f800000u && (b & 0x007fffffu) != 0u;
+{% endif %}
+}
+{%- endmacro %}
+{{ float_bits_def() }}
+{{ is_nan_bits_def() }}
+{% endif %}
+fn key_bits(v: {{ scalar }}) -> u32 {
+{% if isFloat %}
+  let bits = float_bits(v);
+  return select(bits, 0u, bits == {{ "0x8000u" if usesF16 else "0x80000000u" }});
+{% elif isUnsigned %}
+  return v;
+{% else %}
+  return bitcast<u32>(v);
+{% endif %}
+}
+fn hash_key(k: u32) -> u32 {
+  var x = k;
+  x = x ^ (x >> 16u);
+  x = x * 0x7feb352du;
+  x = x ^ (x >> 15u);
+  x = x * 0x846ca68bu;
+  x = x ^ (x >> 16u);
+  return x;
+}
+{% if isFloat and not sorted %}
+fn ordered_key(key: u32) -> u32 {
+{% if usesF16 %}
+  return select(key ^ 0x8000u, (~key) & 0xffffu, (key & 0x8000u) != 0u);
+{% else %}
+  return select(key ^ 0x80000000u, ~key, (key & 0x80000000u) != 0u);
+{% endif %}
+}
+{% endif %}
+fn find_rank(i: u32) -> u32 {
+{% if isFloat %}
+  if (is_nan_bits(x[0])) { return 0u; }
+  // Later NaNs joined the smallest representative that already existed.
+  // Scan distinct representatives, excluding ones first seen after this input.
+  // Sorted slots can stop at their first eligible representative; unsorted
+  // slots appear in first-occurrence order, so later slots can be excluded.
+  if (is_nan_bits(x[i])) {
+    let written = {{ capacity }}u;
+{% if not sorted %}
+    var best = 0u;
+    var bestKey = ordered_key(key_bits(x[slots[0]]));
+{% endif %}
+    for (var p = 0u; p < written; p += 1u) {
+      let first = slots[p];
+{% if sorted %}
+      if (first <= i) { return p; }
+{% else %}
+      if (first > i) { break; }
+      let candidate = ordered_key(key_bits(x[first]));
+      if (candidate < bestKey) { best = p; bestKey = candidate; }
+{% endif %}
+    }
+    return {{ "0u" if sorted else "best" }};
+  }
+{% endif %}
+  let key = key_bits(x[i]);
+  // The first output class is common in repeated-value streams.
+  if (key == key_bits(x[slots[0]])) { return 0u; }
+  if (key == EMPTY) { return rankOfFirst[special[0]]; }
+  var slot = hash_key(key) & MASK;
+  loop {
+    let stored = tableKey[slot];
+    if (stored == key) {
+      return rankOfFirst[tableIdx[slot]];
+    }
+    if (stored == EMPTY) { return 0u; }
+    slot = (slot + 1u) & MASK;
+  }
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= {{ scanN }}u) { return; }
+  let rank = find_rank(i);
+{% if hasInverseIndices %}
+  inverse_indices[i] = rank;
+{% endif %}
+{% if hasCounts %}
+  atomicAdd(&counts[rank], 1u);
+{% endif %}
+}

build/webgpu/unique-scalar-output.wgsl.jinja ADDED Viewed

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

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let p = gid.x;
+  if (p >= {{ capacity }}u) { return; }
+  let first = slots[p];
+  y[p] = x[first];
+{% if hasIndices %}
+  indices[p] = first;
+{% endif %}
+{% if needsRanks %}
+  rankOfFirst[first] = p;
+{% endif %}
+{% if hasCounts %}
+  counts[p] = 0u;
+{% endif %}
+}

build/webgpu/unique-single-class.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,32 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Exact output cardinality is required by the op contract. One class proves
+// that every input belongs to the first representative, including NaNs and
+// whole axis slices. No hash, scan, sort or atomic count is necessary.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i < {{ outputElements }}u) {
+{% if axisPresent %}
+    let source = (i / {{ scalarInner }}u) * {{ inputExtent * scalarInner }}u + i % {{ scalarInner }}u;
+    y[i] = x[source];
+{% else %}
+    y[i] = x[0];
+{% endif %}
+  }
+{% if hasInverseIndices %}
+  if (i < {{ inputExtent }}u) { inverse_indices[i] = 0u; }
+{% endif %}
+{% if hasIndices or hasCounts %}
+  if (i == 0u) {
+{% if hasIndices %}
+    indices[0] = 0u;
+{% endif %}
+{% if hasCounts %}
+    counts[0] = {{ inputExtent }}u;
+{% endif %}
+  }
+{% endif %}
+}

build/webgpu/unique.wgsl.jinja CHANGED Viewed

@@ -65,12 +65,9 @@ fn is_nan_bits(v: {{ scalar }}) -> bool {
 }
 {%- endmacro %}
-{% set emitIndices = source.hasIndices | default(false) %}
-{% set emitInverseIndices = source.hasInverseIndices | default(false) %}
-{% set emitCounts = source.hasCounts | default(false) %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // Floating-point equality/order use the raw IEEE bit pattern, not the float
@@ -103,6 +100,10 @@ fn ordered_less(a: {{ scalar }}, b: {{ scalar }}) -> bool {
 }
 {% if emitIndices or emitInverseIndices or emitCounts %}
 fn representative_input_index(output_i: u32) -> u32 {
   let representative = y[output_i];
   for (var input_i = 0u; input_i < params.inputCount; input_i = input_i + 1u) {
@@ -120,7 +121,7 @@ fn find_bucket(input_i: u32, written: u32) -> u32 {
   let value = x[input_i];
   var candidate = written;
   for (var output_i = 0u; output_i < written; output_i = output_i + 1u) {
-    if (representative_input_index(output_i) > input_i || ordered_less(y[output_i], value)) { continue; }
     if (candidate == written || ordered_less(y[output_i], y[candidate])) {
       candidate = output_i;
     }
@@ -168,9 +169,14 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   for (var i = written; i < params.capacity; i = i + 1u) {
     y[i] = zero_value();
   }
 {% if emitIndices %}
   for (var unique_i = 0u; unique_i < params.capacity; unique_i = unique_i + 1u) {
-    indices[unique_i] = representative_input_index(unique_i);
   }
 {% endif %}
 {% if emitCounts %}

 }
 {%- endmacro %}
+{% set emitIndices = hasIndices | default(false) %}
+{% set emitInverseIndices = hasInverseIndices | default(false) %}
+{% set emitCounts = hasCounts | default(false) %}
 {{ env.wgsl.resourceDeclarations }}
 // Floating-point equality/order use the raw IEEE bit pattern, not the float
 }
 {% if emitIndices or emitInverseIndices or emitCounts %}
+// The representative of an output slot is a pure function of the finished Y and
+// of X, but the bucket resolution below asks for it once per (input, output)
+// pair, which makes an O(n) scan the innermost of three nested loops. It is
+// computed once per slot into `representatives` and read from there instead.
 fn representative_input_index(output_i: u32) -> u32 {
   let representative = y[output_i];
   for (var input_i = 0u; input_i < params.inputCount; input_i = input_i + 1u) {
   let value = x[input_i];
   var candidate = written;
   for (var output_i = 0u; output_i < written; output_i = output_i + 1u) {
+    if (representatives[output_i] > input_i || ordered_less(y[output_i], value)) { continue; }
     if (candidate == written || ordered_less(y[output_i], y[candidate])) {
       candidate = output_i;
     }
   for (var i = written; i < params.capacity; i = i + 1u) {
     y[i] = zero_value();
   }
+{% if emitIndices or emitInverseIndices or emitCounts %}
+  for (var unique_i = 0u; unique_i < params.capacity; unique_i = unique_i + 1u) {
+    representatives[unique_i] = representative_input_index(unique_i);
+  }
+{% endif %}
 {% if emitIndices %}
   for (var unique_i = 0u; unique_i < params.capacity; unique_i = unique_i + 1u) {
+    indices[unique_i] = representatives[unique_i];
   }
 {% endif %}
 {% if emitCounts %}