sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +78 -0
build/webgpu/bench.json +77 -0
build/webgpu/manifest.json +90 -0
build/webgpu/metadata.json +18 -0
build/webgpu/range.wgsl.jinja +36 -0
build/webgpu/test.json +391 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Range
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 27
+## Description
+Generates a 1-D tensor of numbers starting at `start`, incrementing by `delta`, up to but not including `limit`. The output length is `max(ceil((limit - start) / delta), 0)`, and element `i` equals `start + i * delta`. Float16 inputs use float32 intermediate arithmetic when `stash_type` is `1`, then cast each result back to float16.
+See the [ONNX `Range` spec](https://onnx.ai/onnx/operators/onnx__Range.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `start` | `start` | `T` | `0` | — | Scalar first value in the output sequence. | required |
+| `limit` | `limit` | `T` | `0` | — | Scalar exclusive upper bound of the sequence. | required |
+| `delta` | `delta` | `T` | `0` | — | Scalar step size between consecutive output values. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | `1` | — | 1-D tensor containing the generated range of values, with the same dtype as the inputs. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `stash_type` | `1` | TensorProto element type used for float16 intermediate arithmetic; this implementation supports the standard float32 value (`1`). It has no effect for float32 and int32 inputs. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32` |
+## 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
+- [`range.wgsl.jinja`](build/webgpu/range.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:
+- `output`
+Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Range", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({
+  start: { data: startData, shape: [] },
+  limit: { data: limitData, shape: [] },
+  delta: { data: deltaData, shape: [] },
+}, {
+  outputs: { output: { shape: [1], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,77 @@

+{
+  "op": "ai.onnx.Range",
+  "cases": [
+    {
+      "name": "positions_4096",
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [4096] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4096] } }
+    },
+    {
+      "name": "large_fill_i32_16777216",
+      "preset": "smoke",
+      "vars": { "dtype": "int32" },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777216] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [16777216], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "large_fill_f32_16777216",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [16777216.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [16777216], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "dispatch_cliff_2dfold_i32_16777216",
+      "preset": "smoke",
+      "vars": { "dtype": "int32" },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777216] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [16777216], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "dispatch_single_dim_i32_16776960",
+      "preset": "smoke",
+      "vars": { "dtype": "int32" },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16776960] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [16776960], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "tiny_overhead_bound_i32_4096",
+      "preset": "smoke",
+      "vars": { "dtype": "int32" },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [4096] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4096], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * dtypeBytes(args.dtype)" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,90 @@

+{
+  "domain": "ai.onnx",
+  "name": "Range",
+  "sinceVersion": 27,
+  "description": "Generates a 1-D tensor of numbers starting at `start`, incrementing by `delta`, up to but not including `limit`. The output length is `max(ceil((limit - start) / delta), 0)`, and element `i` equals `start + i * delta`. Float16 inputs use float32 intermediate arithmetic when `stash_type` is `1`, then cast each result back to float16.",
+  "inputs": [
+    { "role": "start", "dtype": "T", "description": "Scalar first value in the output sequence.", "rank": 0 },
+    { "role": "limit", "dtype": "T", "description": "Scalar exclusive upper bound of the sequence.", "rank": 0 },
+    { "role": "delta", "dtype": "T", "description": "Scalar step size between consecutive output values.", "rank": 0 }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "description": "1-D tensor containing the generated range of values, with the same dtype as the inputs.",
+      "rank": 1
+    }
+  ],
+  "attributes": { "stash_type": 1 },
+  "attributeDescriptions": {
+    "stash_type": "TensorProto element type used for float16 intermediate arithmetic; this implementation supports the standard float32 value (`1`). It has no effect for float32 and int32 inputs."
+  },
+  "typeConstraints": { "T": ["float32", "float16", "int32"] },
+  "args": {
+    "start": { "kind": "tensor", "semantic": "start", "role": "start" },
+    "limit": { "kind": "tensor", "semantic": "limit", "role": "limit" },
+    "delta": { "kind": "tensor", "semantic": "delta", "role": "delta" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "variants": [
+    {
+      "id": "scalar_inputs",
+      "when": ["ranks.start == 0", "ranks.limit == 0", "ranks.delta == 0", "ranks.output == 1", "(dtypes.T == \"f32\" or dtypes.T == \"f16\" or dtypes.T == \"i32\")", "f16Ok(dtypes.T)", "logicalDtypes.T != \"float16\" or attrs.stash_type == onnxDtypeCode(\"float32\")"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "isI32": "dtypes.T == \"i32\"",
+        "isF16": "dtypes.T == \"f16\"",
+        "usesF16": "dtypes.T == \"f16\"",
+        "vec4": "numel(shapes.output) > 0 and numel(shapes.output) % 4 == 0",
+        "outElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if (numel(shapes.output) > 0 and numel(shapes.output) % 4 == 0) else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Range",
+          "shader": "range.wgsl.jinja",
+          "bindings": [
+            {
+              "name": "start",
+              "arg": "start",
+              "semantic": "start",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 1
+            },
+            {
+              "name": "delta",
+              "arg": "delta",
+              "semantic": "delta",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 1
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$outElement"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "count", "type": "u32", "value": "numel(shapes.output)" },
+                  { "name": "countV4", "type": "u32", "value": "ceilDiv(numel(shapes.output), 4)" }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "threads": "ceilDiv(numel(shapes.output), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,18 @@

+{
+  "name": "ai.onnx.Range",
+  "id": "_ai_onnx_range_webgpu_ee87b02",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "xMGgyqKg5gqzD5cxTqsYuMjzoLMNnUwayCGjdcxlZQ4=",
+      "manifest.json": "KVdMY/jv0Fvuemw67Mwx/pCr0u50yj4Mv8mW96vthrg=",
+      "range.wgsl.jinja": "y8duBaW6JJX+aryiqb6ngPZwW1p9wylUeTbQhzHk7yY=",
+      "test.json": "Tke9pc1jbyjPaMJF3Wxcu129odh9ujdY/1j4rs4SP6M="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Range" }
+}

build/webgpu/range.wgsl.jinja ADDED Viewed

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

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Each output is evaluated independently as start + i*delta rather than by
+// accumulating deltas, preserving the same result across vector widths.
+{% macro value(i) %}
+{% if isI32 %}start[0] + i32({{ i }}) * delta[0]{% elif isF16 %}f16(f32(start[0]) + f32({{ i }}) * f32(delta[0])){% else %}start[0] + f32({{ i }}) * delta[0]{% endif %}
+{% endmacro %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // Generate four adjacent values per invocation while retaining scalar
+  // storage for arbitrary output lengths.
+  let q = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  let base = q * 4u;
+{% if vec4 %}
+  // A whole number of quads lets each invocation write one vec4 without a tail.
+  if (q >= params.countV4) {
+    return;
+  }
+  output[q] = vec4<{{ scalar }}>(
+    {{ value('base') }},
+    {{ value('(base + 1u)') }},
+    {{ value('(base + 2u)') }},
+    {{ value('(base + 3u)') }});
+{% else %}
+  if (base >= params.count) {
+    return;
+  }
+  let end = min(base + 4u, params.count);
+  for (var i = base; i < end; i = i + 1u) {
+    output[i] = {{ value('i') }};
+  }
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,391 @@

+{
+  "op": "ai.onnx.Range",
+  "cases": [
+    {
+      "name": "dispatch_cliff_i32_16777216",
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777216] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [16777216], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_positive",
+      "provenance": {
+        "notes": "Range-27 defines stash_type only for float16/bfloat16 accumulation, so the non-default FLOAT16 code has no effect on this float32 route."
+      },
+      "attrs": { "stash_type": 10 },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.5] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [6.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3] } }
+    },
+    {
+      "name": "i32_scalar_inputs_positive_delta",
+      "provenance": {
+        "notes": "Covers a positive integer step whose interval length is exactly divisible by the step."
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3] } }
+    },
+    {
+      "name": "f32_true_scalar_inputs",
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.5] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_negative_delta",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Float_NegativeDelta_0"
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-8.1] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_negative_delta_empty",
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [9.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_same_start_limit_empty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Float_SameStartAndLimit"
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_almost_same_high_delta_one_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.AlmostSameStartAndLimitHighDelta"
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.01] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1000000.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_subnormal_positive_delta_length_rounding_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal FTZ: subnormal f32 start/limit/delta (1e-40) flush to zero on GPU, so the generated range collapses to zeros instead of the subnormal sequence (tolerance 2e-45)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Float_Delta_NonDefault",
+        "notes": "Valid float32 denormal-step edge: ORT computes length from rounded f32 inputs, producing five values for [0, 4e-40) at delta 1e-40."
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [4e-40] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1e-40] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [5],
+          "tolerance": 2e-45,
+          "data": {
+            "kind": "values",
+            "values": [0.0, 9.99994610111476e-41, 1.999989220222952e-40, 2.999983830334428e-40, 3.999978440445904e-40]
+          }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_negative_delta_length_rounding_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal FTZ: subnormal f32 start/limit/delta (1e-40) flush to zero on GPU, so the generated range collapses to zeros instead of the subnormal sequence (tolerance 2e-45)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Float_NegativeDelta_0",
+        "notes": "Negative denormal-step companion; the output shape and values follow ORT's rounded float32 Range semantics."
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-4e-40] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1e-40] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [5],
+          "tolerance": 2e-45,
+          "data": {
+            "kind": "values",
+            "values": [0.0, -9.99994610111476e-41, -1.999989220222952e-40, -2.999983830334428e-40, -3.999978440445904e-40]
+          }
+        }
+      }
+    },
+    {
+      "name": "f32_subnormal_positive_delta_shifted_start_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal FTZ: subnormal f32 start/limit/delta (1e-40) flush to zero on GPU, so the generated range collapses to zeros instead of the subnormal sequence (tolerance 2e-45)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Float_Delta_NonDefault",
+        "notes": "Shifted-start companion to the denormal-step Range edge: every output value is a valid positive subnormal, including the first element."
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1e-40] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5e-40] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1e-40] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 2e-45 } }
+    },
+    {
+      "name": "f32_subnormal_negative_delta_shifted_start_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal FTZ: subnormal f32 start/limit/delta (1e-40) flush to zero on GPU, so the generated range collapses to zeros instead of the subnormal sequence (tolerance 2e-45)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Float_NegativeDelta_0",
+        "notes": "Negative shifted-start companion: the generated sequence should preserve finite negative subnormal values instead of flushing the range to zeros."
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1e-40] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5e-40] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1e-40] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 2e-45 } }
+    },
+    {
+      "name": "ort_f32_delta_default_style",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Float_DeltaDefault"
+      },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_i32_delta_default_style",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Int32_DeltaDefault"
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_i32_delta_nondefault_even",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Int32_Delta_NonDefault"
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_i32_delta_nondefault_odd_limit",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Int64_Delta_NonDefault_0",
+        "notes": "Native int32 projection of ORT's int64 odd-limit nondefault-delta case."
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [9] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_i32_delta_nondefault_one_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Int64_Delta_NonDefault_1",
+        "notes": "Native int32 projection of ORT's int64 one-value nondefault-delta case."
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_i32_negative_delta",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Int32_NegativeDelta_0"
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-9] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_i32_negative_delta_empty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/range_test.cc",
+        "test": "RangeTest.Int32_NegativeDelta_1"
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [9] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "i32_exact_above_float24_no_f32_cast",
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777217] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777221] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_range_float_type_positive_delta",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_range_float_type_positive_delta" },
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_range_int32_type_negative_delta",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_range_int32_type_negative_delta" },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [6] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_2dfold_dispatch_16777216",
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [16777216.0] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [16777216], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_fractional_delta_nonmult4_len1027",
+      "inputs": {
+        "start": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "limit": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [256.75] } },
+        "delta": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1027], "tolerance": 0 } }
+    },
+    {
+      "name": "i32_negative_start_positive_delta_len8",
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-3] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "i32_negative_start_negative_delta_descending",
+      "provenance": {
+        "notes": "Negative start with negative delta produces a descending negative sequence [-5,-8,-11,-14,-17]; exercises start + i32(i)*delta with a negative base and negative step, a sign combination not covered by existing negative-delta cases (which start positive)."
+      },
+      "inputs": {
+        "start": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-5] } },
+        "limit": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-20] } },
+        "delta": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx27_f16_stash_float32_vec4",
+      "provenance": {
+        "source": "ONNX Operators.md Range-27",
+        "notes": "Pins the float16 type and explicit float32 stash semantics added by Range-27 on the vec4 output route."
+      },
+      "attrs": { "stash_type": 1 },
+      "inputs": {
+        "start": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2048.0] } },
+        "limit": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2052.0] } },
+        "delta": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [2048.0, 2048.0, 2050.0, 2052.0] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "onnx27_f16_stash_float32_scalar_tail",
+      "provenance": {
+        "source": "ONNX Operators.md Range-27",
+        "notes": "Non-multiple-of-four companion which pins float32 accumulation followed by a float16 cast on the scalar-tail route."
+      },
+      "attrs": { "stash_type": 1 },
+      "inputs": {
+        "start": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [-1.5] } },
+        "limit": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [0.0] } },
+        "delta": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float16",
+          "shape": [3],
+          "data": { "kind": "values", "values": [-1.5, -1.0, -0.5] },
+          "tolerance": 0
+        }
+      }
+    }
+  ]
+}