sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +73 -0
build/webgpu/affine-grid.wgsl.jinja +75 -0
build/webgpu/affine-grid3d-w4.wgsl.jinja +54 -0
build/webgpu/bench.json +36 -0
build/webgpu/manifest.json +164 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +567 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.AffineGrid
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 20
+## Description
+Generates a 2-D or 3-D flow field (sampling grid) from a batch of affine matrices `theta`. Each affine matrix is applied to the target output's normalized coordinate grid in `[-1, 1]`, producing sample coordinates at every output position.
+See the [ONNX `AffineGrid` spec](https://onnx.ai/onnx/operators/onnx__AffineGrid.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `theta` | `theta` | `T` | `3` | — | Batch of affine matrices with shape `(N, 2, 3)` for 2-D or `(N, 3, 4)` for 3-D. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `grid` | `grid` | `T` | derived | — | Sampling grid of 2-D coordinates with shape `(N, H, W, 2)`, or 3-D coordinates with shape `(N, D, H, W, 3)`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `align_corners` | `0` | When 1, the values -1 and 1 refer to the centers of the corner pixels; when 0 (default), they refer to the outer edges of the corner pixels. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`affine-grid.wgsl.jinja`](build/webgpu/affine-grid.wgsl.jinja)
+- [`affine-grid3d-w4.wgsl.jinja`](build/webgpu/affine-grid3d-w4.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:
+- `grid`
+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.AffineGrid", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { grid } = await kernel({ theta: { data: thetaData, shape: [1, 2, 3] } }, {
+  outputs: { grid: { shape: [1, 3, 2, 2], dtype: "float32" } },
+});
+```

build/webgpu/affine-grid.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,75 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+fn lin(coord: u32, size: u32) -> f32 {
+  {% if source.alignCorners %}
+  if (size <= 1u) { return 0.0; }
+  return -1.0 + 2.0 * f32(coord) / f32(size - 1u);
+  {% else %}
+  return (2.0 * f32(coord) + 1.0) / f32(size) - 1.0;
+  {% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("p", guardInline=true) }}
+  let w = p % params.W;
+{% if source.spatialRank == 3 %}
+  var remaining = p / params.W;
+  let h = remaining % params.H;
+  remaining /= params.H;
+  let d = remaining % params.D;
+  let n = remaining / params.D;
+{% else %}
+  let h = (p / params.W) % params.H;
+  let n = p / (params.W * params.H);
+{% endif %}
+  let x0 = lin(w, params.W);
+  let y0 = lin(h, params.H);
+{% if source.spatialRank == 3 %}
+  let z0 = lin(d, params.D);
+  let t = n * 12u;
+  grid[p * 3u] = {{ scalar }}(f32(theta[t]) * x0 + f32(theta[t + 1u]) * y0 + f32(theta[t + 2u]) * z0 + f32(theta[t + 3u]));
+  grid[p * 3u + 1u] = {{ scalar }}(f32(theta[t + 4u]) * x0 + f32(theta[t + 5u]) * y0 + f32(theta[t + 6u]) * z0 + f32(theta[t + 7u]));
+  grid[p * 3u + 2u] = {{ scalar }}(f32(theta[t + 8u]) * x0 + f32(theta[t + 9u]) * y0 + f32(theta[t + 10u]) * z0 + f32(theta[t + 11u]));
+{% else %}
+  let t = n * 6u;
+  grid[p * 2u] = {{ scalar }}(f32(theta[t]) * x0 + f32(theta[t + 1u]) * y0 + f32(theta[t + 2u]));
+  grid[p * 2u + 1u] = {{ scalar }}(f32(theta[t + 3u]) * x0 + f32(theta[t + 4u]) * y0 + f32(theta[t + 5u]));
+{% endif %}
+}

build/webgpu/affine-grid3d-w4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,54 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{% set STORE_OPEN = vectorScalar ~ "(vec4<f32>(" if usesF16 else vectorScalar ~ "(" %}
+{% set STORE_CLOSE = "))" if usesF16 else ")" %}
+{{ env.wgsl.resourceDeclarations }}
+const COUNT4: u32 = {{ source.count4 }}u;
+const D: u32 = {{ source.D }}u;
+const H: u32 = {{ source.H }}u;
+const W: u32 = {{ source.W }}u;
+const W4: u32 = {{ source.W4 }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+fn lin(coord: u32, size: u32) -> f32 {
+{% if source.alignCorners %}
+  if (size <= 1u) { return 0.0; }
+  return -1.0 + 2.0 * f32(coord) / f32(size - 1u);
+{% else %}
+  return (2.0 * f32(coord) + 1.0) / f32(size) - 1.0;
+{% endif %}
+}
+fn transform(t: u32, x0: f32, y0: f32, z0: f32) -> vec3<f32> {
+  return vec3<f32>(
+    f32(theta[t]) * x0 + f32(theta[t + 1u]) * y0 + f32(theta[t + 2u]) * z0 + f32(theta[t + 3u]),
+    f32(theta[t + 4u]) * x0 + f32(theta[t + 5u]) * y0 + f32(theta[t + 6u]) * z0 + f32(theta[t + 7u]),
+    f32(theta[t + 8u]) * x0 + f32(theta[t + 9u]) * y0 + f32(theta[t + 10u]) * z0 + f32(theta[t + 11u])
+  );
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  let q = gid.x + gid.y * nwg.x * WG;
+  if (q >= COUNT4) { return; }
+  let w4 = q % W4;
+  var r = q / W4;
+  let h = r % H;
+  r /= H;
+  let d = r % D;
+  let n = r / D;
+  let y0 = lin(h, H);
+  let z0 = lin(d, D);
+  let w0 = w4 * 4u;
+  let t = n * 12u;
+  let p0 = transform(t, lin(w0, W), y0, z0);
+  let p1 = transform(t, lin(w0 + 1u, W), y0, z0);
+  let p2 = transform(t, lin(w0 + 2u, W), y0, z0);
+  let p3 = transform(t, lin(w0 + 3u, W), y0, z0);
+  let o = q * 3u;
+  grid[o] = {{ STORE_OPEN }}p0.x, p0.y, p0.z, p1.x{{ STORE_CLOSE }};
+  grid[o + 1u] = {{ STORE_OPEN }}p1.y, p1.z, p2.x, p2.y{{ STORE_CLOSE }};
+  grid[o + 2u] = {{ STORE_OPEN }}p2.z, p3.x, p3.y, p3.z{{ STORE_CLOSE }};
+}

build/webgpu/bench.json ADDED Viewed

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

+{
+  "op": "ai.onnx.AffineGrid",
+  "cases": [
+    {
+      "name": "1x256x256_f32",
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 256, 256, 2] } },
+      "preset": "smoke"
+    },
+    {
+      "name": "affine-grid3d-f32-b1-64x128x128-medical-volume",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "points": 1048576 },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 64, 128, 128, 3] } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.points * 3 * dtypeBytes(args.dtype)" }]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,164 @@

+{
+  "domain": "ai.onnx",
+  "name": "AffineGrid",
+  "sinceVersion": 20,
+  "description": "Generates a 2-D or 3-D flow field (sampling grid) from a batch of affine matrices `theta`. Each affine matrix is applied to the target output's normalized coordinate grid in `[-1, 1]`, producing sample coordinates at every output position.",
+  "inputs": [
+    {
+      "role": "theta",
+      "dtype": "T",
+      "description": "Batch of affine matrices with shape `(N, 2, 3)` for 2-D or `(N, 3, 4)` for 3-D.",
+      "rank": 3
+    }
+  ],
+  "outputs": [
+    {
+      "role": "grid",
+      "dtype": "T",
+      "rank": "4 if dim(shapes.theta, 1) == 2 else 5",
+      "description": "Sampling grid of 2-D coordinates with shape `(N, H, W, 2)`, or 3-D coordinates with shape `(N, D, H, W, 3)`."
+    }
+  ],
+  "attributes": { "align_corners": 0 },
+  "attributeDescriptions": {
+    "align_corners": "When 1, the values -1 and 1 refer to the centers of the corner pixels; when 0 (default), they refer to the outer edges of the corner pixels."
+  },
+  "attributeConstraints": { "align_corners": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "theta": { "kind": "tensor", "semantic": "theta", "role": "input" },
+    "grid": { "kind": "tensor", "semantic": "grid", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "dtypeSupported": "f16Ok(dtypes.T)",
+    "grid2dContract": "dtypeSupported and ranks.theta == 3 and ranks.grid == 4 and dim(shapes.theta, 1) == 2 and dim(shapes.theta, 2) == 3 and dim(shapes.grid, 0) == dim(shapes.theta, 0) and dim(shapes.grid, 3) == 2",
+    "grid3dContract": "dtypeSupported and ranks.theta == 3 and ranks.grid == 5 and dim(shapes.theta, 1) == 3 and dim(shapes.theta, 2) == 4 and dim(shapes.grid, 0) == dim(shapes.theta, 0) and dim(shapes.grid, 4) == 3"
+  },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "bindingSets": {
+    "grid3dVector": [
+      {
+        "name": "theta",
+        "arg": "theta",
+        "semantic": "theta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "grid",
+        "arg": "grid",
+        "semantic": "grid",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      }
+    ],
+    "grid2dScalar": [
+      {
+        "name": "theta",
+        "arg": "theta",
+        "semantic": "theta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "grid", "arg": "grid", "semantic": "grid", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.grid) / 2" },
+            { "name": "H", "type": "u32", "value": "dim(shapes.grid, 1)" },
+            { "name": "W", "type": "u32", "value": "dim(shapes.grid, 2)" }
+          ]
+        }
+      }
+    ],
+    "grid3dScalar": [
+      {
+        "name": "theta",
+        "arg": "theta",
+        "semantic": "theta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "grid", "arg": "grid", "semantic": "grid", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.grid) / 3" },
+            { "name": "D", "type": "u32", "value": "dim(shapes.grid, 1)" },
+            { "name": "H", "type": "u32", "value": "dim(shapes.grid, 2)" },
+            { "name": "W", "type": "u32", "value": "dim(shapes.grid, 3)" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "3d_w4",
+      "priority": 20,
+      "when": ["grid3dContract", "dim(shapes.grid, 3) > 0", "dim(shapes.grid, 3) % 4 == 0"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"", "usesF16": "dtypes.T == \"f16\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "AffineGrid3D.W4",
+          "source": {
+            "shader": "affine-grid3d-w4.wgsl.jinja",
+            "inputs": {
+              "alignCorners": "attrs.align_corners != 0",
+              "count4": "numel(shapes.grid) / 3 / 4",
+              "D": "dim(shapes.grid, 1)",
+              "H": "dim(shapes.grid, 2)",
+              "W": "dim(shapes.grid, 3)",
+              "W4": "dim(shapes.grid, 3) / 4"
+            }
+          },
+          "bindings": "grid3dVector",
+          "dispatch": { "threads": "numel(shapes.grid) / 3 / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "2d",
+      "when": ["grid2dContract"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "AffineGrid",
+          "source": {
+            "shader": "affine-grid.wgsl.jinja",
+            "inputs": { "alignCorners": "attrs.align_corners != 0", "spatialRank": 2 }
+          },
+          "bindings": "grid2dScalar",
+          "dispatch": { "threads": "numel(shapes.grid) / 2", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "3d",
+      "priority": 10,
+      "when": ["grid3dContract"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "AffineGrid3D",
+          "source": {
+            "shader": "affine-grid.wgsl.jinja",
+            "inputs": { "alignCorners": "attrs.align_corners != 0", "spatialRank": 3 }
+          },
+          "bindings": "grid3dScalar",
+          "dispatch": { "threads": "numel(shapes.grid) / 3", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,19 @@

+{
+  "name": "ai.onnx.AffineGrid",
+  "id": "_ai_onnx_affinegrid_webgpu_10a1b9d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "affine-grid.wgsl.jinja": "hemLmXc03t/I3JyiZ8OneDr0uaL8q7cwVwUZPZcF7QE=",
+      "affine-grid3d-w4.wgsl.jinja": "pYw6LFXYyqXUgA8id761xxtP2UzQieyOpCKoegZFs5E=",
+      "bench.json": "x1krZetoEyoKMLV4F08UOzIemtdXC4MpoDqERBhDQhM=",
+      "manifest.json": "dYhH9vmmJYbl+nExkZfGUwNlgPtp4F6OJnwXwQkWeIw=",
+      "test.json": "eB8AkQmDzUaHcUhNazVmAXned4+KqnFEuiaOX9kL7gE="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.AffineGrid" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,567 @@

+{
+  "op": "ai.onnx.AffineGrid",
+  "fixtureArrays": {
+    "ort_generated_3d_rotation_shear_batch2_no_align_input_theta": [1.409539, 0, 0.51303, 0.3, 0.118782, 1.969615, -0.326352, -0.5, -0.168412, 0.086824, 0.462708, 1.8, 1.409539, 0, 0.51303, 0.3, 0.118782, 1.969615, -0.326352, -0.5, -0.168412, 0.086824, 0.462708, 1.8],
+    "ort_generated_3d_stronger_transform_batch2_no_align_input_theta": [0.259808, 0, -0.15, -0.5, -1.299038, 1.5, -2.25, -0.5, 1.375, 4.76314, 2.38157, 0.3, 0.259808, 0, -0.15, -0.5, -1.299038, 1.5, -2.25, -0.5, 1.375, 4.76314, 2.38157, 0.3],
+    "onnx_backend_affine_grid_3d_input_theta": [2.6830732822418213, -0.7943316102027893, 0.21829216182231903, 5, 0.622253954410553, 3.2880465984344482, -0.5303300619125366, -3.299999952316284, 0.24721935391426086, 1.7241772413253784, 0.07809311151504517, -1.100000023841858, -0.35552558302879333, 1.004422903060913, 1.3995190858840942, 2.5, 0.17578838765621185, 0.06028856709599495, -0.9240381121635437, 1.100000023841858, -1.100000023841858, -0.44999998807907104, 0.44999998807907104, 2.200000047683716]
+  },
+  "cases": [
+    {
+      "name": "affine_grid3d_w4_coarsened",
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [1.0, 0.1, 0.0, 0.05, 0.0, 1.0, -0.1, -0.05, 0.1, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 2, 3, 4, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_subnormal_translation_2d_singleton_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 translation (1e-40) in theta flushes to zero on GPU, so the singleton grid reads 0 instead of the subnormal translation (tolerance 0)."
+      },
+      "provenance": {
+        "notes": "A singleton 2D output grid has normalized coordinates (0, 0), so the grid should preserve theta's subnormal translation exactly."
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [0.0, 0.0, 1e-40, 0.0, 0.0, -1e-40] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 1, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_subnormal_translation_3d_singleton_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 translation (1e-40) in theta flushes to zero on GPU, so the singleton grid reads 0 instead of the subnormal translation (tolerance 0)."
+      },
+      "provenance": {
+        "notes": "A singleton 3D output grid should preserve subnormal x/y/z translations without arithmetic amplification."
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 1e-40, 0.0, 0.0, 0.0, -1e-40, 0.0, 0.0, 0.0, 1e-40] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 1, 1, 1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "dispatch_cliff_2d_over_16m_elements",
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.5, 0.25, -0.5, 2.0, -1.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 4097, 4097, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "identity_2x2_align_corners",
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 2, 2, 2] } },
+      "tolerance": 0.000001
+    },
+    {
+      "name": "translate_f16",
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float16",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.5, 0.0, 1.0, -0.5] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float16", "shape": [1, 2, 2, 2] } },
+      "tolerance": 0.002
+    },
+    {
+      "name": "identity_3x4_no_align_corners",
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 4, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "batch2_scale_translate_align_corners",
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.25, 0.0, 1.0, -0.25, 0.5, 0.0, -0.5, 0.0, 0.5, 0.5] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_generated_2d_rotation_shear_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test_gen.py",
+        "test": "AffineGridTest.test_2d_0",
+        "notes": "Generated ORT fixture materialized in affine_grid_test.cc."
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.477212, -0.173648, 0.3, 0.173648, 0.492404, -0.5] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_2d_batch2_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_2d_7"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.5, -0.866025, -0.5, 0.866025, 2.75, -0.5, 1.5, -0.866025, -0.5, 0.866025, 2.75, -0.5]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_2d_identity_default_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.2d"
+      },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 2, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_2d_batch2_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_2d_1"
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.477212, -0.173648, 0.3, 0.173648, 0.492404, -0.5, 1.477212, -0.173648, 0.3, 0.173648, 0.492404, -0.5]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_2d_stronger_transform_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_2d_2"
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.5, -0.866025, -0.5, 0.866025, 2.75, -0.5] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_2d_rotation_shear_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_2d_4"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.477212, -0.173648, 0.3, 0.173648, 0.492404, -0.5] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_2d_stronger_transform_batch2_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_2d_3"
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.5, -0.866025, -0.5, 0.866025, 2.75, -0.5, 1.5, -0.866025, -0.5, 0.866025, 2.75, -0.5]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_2d_rotation_shear_batch2_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_2d_5"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.477212, -0.173648, 0.3, 0.173648, 0.492404, -0.5, 1.477212, -0.173648, 0.3, 0.173648, 0.492404, -0.5]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_2d_stronger_transform_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_2d_6"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.5, -0.866025, -0.5, 0.866025, 2.75, -0.5] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_rotation_shear_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_0"
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.409539, 0.0, 0.51303, 0.3, 0.118782, 1.969615, -0.326352, -0.5, -0.168412, 0.086824, 0.462708, 1.8]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_rotation_shear_batch2_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_1"
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_generated_3d_rotation_shear_batch2_no_align_input_theta" }
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 2, 3, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_stronger_transform_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_2"
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.259808, 0.0, -0.15, -0.5, -1.299038, 1.5, -2.25, -0.5, 1.375, 4.76314, 2.38157, 0.3]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_stronger_transform_batch2_no_align",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_3"
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_generated_3d_stronger_transform_batch2_no_align_input_theta" }
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 2, 3, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_rotation_shear_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_4"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.409539, 0.0, 0.51303, 0.3, 0.118782, 1.969615, -0.326352, -0.5, -0.168412, 0.086824, 0.462708, 1.8]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_rotation_shear_batch2_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_5"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_generated_3d_rotation_shear_batch2_no_align_input_theta" }
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 2, 3, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_stronger_transform_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_6"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.259808, 0.0, -0.15, -0.5, -1.299038, 1.5, -2.25, -0.5, 1.375, 4.76314, 2.38157, 0.3]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 3, 2, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_generated_3d_stronger_transform_batch2_align_corners",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/affine_grid_test.cc",
+        "test": "AffineGridTest.test_3d_7"
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_generated_3d_stronger_transform_batch2_no_align_input_theta" }
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 2, 2, 3, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_affine_grid_2d",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_affine_grid_2d",
+        "notes": "Official size metadata is represented by the output grid shape in this framework."
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.088944435119629, -3.2880465984344482, 5.0, 2.022325277328491, 1.096015453338623, -3.299999952316284, 0.835788369178772, -0.5544228553771973, 2.5, 0.7876279354095459, 0.8397114276885986, 1.100000023841858]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 5, 6, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_affine_grid_2d_align_corners",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_affine_grid_2d_align_corners",
+        "notes": "Official size metadata is represented by the output grid shape in this framework."
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.088944435119629, -3.2880465984344482, 5.0, 2.022325277328491, 1.096015453338623, -3.299999952316284, 0.835788369178772, -0.5544228553771973, 2.5, 0.7876279354095459, 0.8397114276885986, 1.100000023841858]
+          }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 5, 6, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_affine_grid_3d",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_affine_grid_3d",
+        "notes": "Official size metadata is represented by the output grid shape in this framework."
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_affine_grid_3d_input_theta" } }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 4, 5, 6, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_affine_grid_3d_align_corners",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_affine_grid_3d_align_corners",
+        "notes": "Official size metadata is represented by the output grid shape in this framework."
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_affine_grid_3d_input_theta" } }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [2, 4, 5, 6, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "align_corners": 0 },
+      "inputs": { "theta": { "dtype": "float32", "shape": [0, 2, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "grid": { "dtype": "float32", "shape": [0, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_dim_f16",
+      "attrs": { "align_corners": 0 },
+      "inputs": { "theta": { "dtype": "float16", "shape": [0, 2, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "grid": { "dtype": "float16", "shape": [0, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "align_corners_1_h1_w3_identity",
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 1, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "3d_f16_identity_no_align",
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float16",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float16", "shape": [1, 2, 2, 2, 3] } },
+      "tolerance": 0.002
+    },
+    {
+      "name": "affine_grid3d_w4_coarsened_f16",
+      "provenance": {
+        "notes": "A float16 width-by-four 3D grid checks half-precision inputs and output on the vectorized 3D path."
+      },
+      "attrs": { "align_corners": 0 },
+      "inputs": {
+        "theta": {
+          "dtype": "float16",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [1.0, 0.1, 0.0, 0.05, 0.0, 1.0, -0.1, -0.05, 0.1, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float16", "shape": [1, 2, 3, 4, 3], "tolerance": 0.002 } }
+    },
+    {
+      "name": "affine_grid3d_w4_align_corners",
+      "provenance": {
+        "notes": "align_corners=1 on the width-by-four 3D grid. The flag changes the normalized-coordinate step the kernel bakes in, and every case that reached this route left it at 0."
+      },
+      "attrs": { "align_corners": 1 },
+      "inputs": {
+        "theta": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": { "kind": "values", "values": [1.0, 0.1, 0.0, 0.05, 0.0, 1.0, -0.1, -0.05, 0.1, 0.0, 1.0, 0.0] }
+        }
+      },
+      "outputs": { "grid": { "dtype": "float32", "shape": [1, 2, 3, 4, 3], "tolerance": 0.000001 } }
+    }
+  ]
+}