sync 2e7068faf55e

Browse files

Files changed (10) hide show

README.md +76 -0
build/webgpu/bench.json +221 -0
build/webgpu/datamove-flat-copy.wgsl.jinja +18 -0
build/webgpu/datamove-transpose-2d-tiled-scalar.wgsl.jinja +57 -0
build/webgpu/datamove-transpose-2d-tiled.wgsl.jinja +55 -0
build/webgpu/datamove-transpose-vec4.wgsl.jinja +44 -0
build/webgpu/manifest.json +240 -0
build/webgpu/metadata.json +22 -0
build/webgpu/test.json +1097 -0
build/webgpu/transpose.wgsl.jinja +38 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Transpose
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Transposes the input tensor by permuting its axes according to the `perm` attribute. Axis `i` of the output corresponds to axis `perm[i]` of the input; if `perm` is omitted, the axes are reversed (`n-1, ..., 0`).
+See the [ONNX `Transpose` spec](https://onnx.ai/onnx/operators/onnx__Transpose.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `data` | `x` | `T` | — | — | The input tensor to transpose. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `transposed` | `y` | `T` | same as `data` | — | The transposed output tensor with permuted axes. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `perm` | — | Optional permutation of the input axes. It must contain every axis from 0 through rank - 1 exactly once. When omitted, the axes are reversed. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `int16`, `uint32`, `uint8`, `int8`, `bool` |
+## 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
+- [`datamove-flat-copy.wgsl.jinja`](build/webgpu/datamove-flat-copy.wgsl.jinja)
+- [`datamove-transpose-2d-tiled-scalar.wgsl.jinja`](build/webgpu/datamove-transpose-2d-tiled-scalar.wgsl.jinja)
+- [`datamove-transpose-2d-tiled.wgsl.jinja`](build/webgpu/datamove-transpose-2d-tiled.wgsl.jinja)
+- [`datamove-transpose-vec4.wgsl.jinja`](build/webgpu/datamove-transpose-vec4.wgsl.jinja)
+- [`transpose.wgsl.jinja`](build/webgpu/transpose.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.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Transpose", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { y } = await kernel({ x: { data: xData, shape: [] } }, {
+  outputs: { y: { shape: [], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,221 @@

+{
+  "op": "ai.onnx.Transpose",
+  "tunableSpace": {
+    "WORKGROUP_SIZE": [64, 128, 256],
+    "INNER_VEC4_ITEMS_PER_THREAD": [1, 2, 4, 8],
+    "INNER_VEC4_MIN_THREADS": [16384, 65536, 262144]
+  },
+  "cases": [
+    {
+      "name": "transpose-f32-1024x1024",
+      "preset": "smoke",
+      "vars": { "rows": 1024, "cols": 1024 },
+      "attrs": { "perm": [1, 0] },
+      "inputs": { "x": { "shape": [1024, 1024], "dtype": "float32", "dist": "normal", "seed": 120, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [1024, 1024], "dtype": "float32" } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * 4 * 2" }] }
+    },
+    {
+      "name": "transpose-f32-rank1-1m",
+      "preset": "smoke",
+      "vars": { "count": 1048576, "dtype": "float32" },
+      "inputs": { "x": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 121, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [1048576], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-tiled-f32-2048x2048",
+      "preset": "smoke",
+      "vars": { "rows": 2048, "cols": 2048, "dtype": "float32" },
+      "attrs": { "perm": [1, 0] },
+      "inputs": { "x": { "shape": [2048, 2048], "dtype": "float32", "dist": "normal", "seed": 122, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [2048, 2048], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-tiled-f16-4096x1024",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 1024, "dtype": "float16" },
+      "attrs": { "perm": [1, 0] },
+      "inputs": { "x": { "shape": [4096, 1024], "dtype": "float16", "dist": "normal", "seed": 123, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [1024, 4096], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-batched-f32-8x512x256",
+      "preset": "smoke",
+      "vars": { "batch": 8, "rows": 512, "cols": 256, "dtype": "float32" },
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": { "x": { "shape": [8, 512, 256], "dtype": "float32", "dist": "normal", "seed": 124, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [8, 256, 512], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.rows * args.cols * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "transpose-inner-vec4-f32-8x64x32x128",
+      "preset": "smoke",
+      "vars": { "count": 2097152, "dtype": "float32" },
+      "attrs": { "perm": [0, 2, 1, 3] },
+      "inputs": { "x": { "shape": [8, 64, 32, 128], "dtype": "float32", "dist": "normal", "seed": 125, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [8, 32, 64, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-identity-f32-4m",
+      "preset": "edge",
+      "vars": { "count": 4194304, "dtype": "float32" },
+      "attrs": { "perm": [0, 1] },
+      "inputs": { "x": { "shape": [2048, 2048], "dtype": "float32", "dist": "normal", "seed": 126, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [2048, 2048], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-scalar-f32-1024x1023",
+      "preset": "edge",
+      "vars": { "rows": 1024, "cols": 1023, "dtype": "float32" },
+      "attrs": { "perm": [1, 0] },
+      "inputs": { "x": { "shape": [1024, 1023], "dtype": "float32", "dist": "normal", "seed": 130, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [1023, 1024], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-scalar-f32-768x50257-vocab",
+      "preset": "edge",
+      "vars": { "rows": 768, "cols": 50257, "dtype": "float32" },
+      "attrs": { "perm": [1, 0] },
+      "inputs": { "x": { "shape": [768, 50257], "dtype": "float32", "dist": "normal", "seed": 131, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [50257, 768], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-scalar-f32-rank3-8x512x255",
+      "preset": "edge",
+      "vars": { "batch": 8, "rows": 512, "cols": 255, "dtype": "float32" },
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": { "x": { "shape": [8, 512, 255], "dtype": "float32", "dist": "normal", "seed": 132, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [8, 255, 512], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.rows * args.cols * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "transpose-tiled-i32-1024x1024",
+      "preset": "edge",
+      "vars": { "rows": 1024, "cols": 1024, "dtype": "int32" },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": { "shape": [1024, 1024], "dtype": "int32", "dist": "randint", "seed": 133, "min": -1000, "max": 1000 }
+      },
+      "outputs": { "y": { "shape": [1024, 1024], "dtype": "int32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-scalar-f32-rank4-nhwc-to-nchw-8x64x96x96",
+      "preset": "stress",
+      "vars": { "n": 8, "h": 64, "w": 96, "c": 96, "dtype": "float32" },
+      "attrs": { "perm": [0, 3, 1, 2] },
+      "inputs": { "x": { "shape": [8, 64, 96, 96], "dtype": "float32", "dist": "normal", "seed": 140, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [8, 96, 64, 96], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.n * args.h * args.w * args.c * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "transpose-scalar-f32-rank3-default-reverse-96x200x128",
+      "preset": "stress",
+      "vars": { "d0": 96, "d1": 200, "d2": 128, "dtype": "float32" },
+      "inputs": { "x": { "shape": [96, 200, 128], "dtype": "float32", "dist": "normal", "seed": 141, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [128, 200, 96], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.d0 * args.d1 * args.d2 * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "transpose-scalar-f32-rank5-ncdhw-to-ndhwc-volumetric",
+      "preset": "stress",
+      "vars": { "count": 8388608, "dtype": "float32" },
+      "attrs": { "perm": [0, 2, 3, 4, 1] },
+      "inputs": {
+        "x": { "shape": [1, 64, 32, 64, 64], "dtype": "float32", "dist": "normal", "seed": 142, "scale": 0.2 }
+      },
+      "outputs": { "y": { "shape": [1, 32, 64, 64, 64], "dtype": "float32", "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-scalar-f32-rank5-ndhwc-to-ncdhw-volumetric",
+      "preset": "stress",
+      "vars": { "count": 8388608, "dtype": "float32" },
+      "attrs": { "perm": [0, 4, 1, 2, 3] },
+      "inputs": {
+        "x": { "shape": [1, 32, 64, 64, 64], "dtype": "float32", "dist": "normal", "seed": 143, "scale": 0.2 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 32, 64, 64], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "transpose-f32-rank4-nchw-to-nhwc-8x64x96x96",
+      "preset": "stress",
+      "vars": { "n": 8, "h": 64, "w": 96, "c": 96, "dtype": "float32" },
+      "attrs": { "perm": [0, 2, 3, 1] },
+      "inputs": { "x": { "shape": [8, 64, 96, 96], "dtype": "float32", "dist": "normal", "seed": 140, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [8, 96, 96, 64], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.n * args.h * args.w * args.c * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "transpose-unitdim-memcpy-f32-4096x1024x1-perm201",
+      "preset": "smoke",
+      "vars": { "tokens": 4096, "hidden": 1024, "dtype": "float32" },
+      "attrs": { "perm": [2, 0, 1] },
+      "inputs": { "x": { "shape": [4096, 1024, 1], "dtype": "float32", "dist": "normal", "seed": 141, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [1, 4096, 1024], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.tokens * args.hidden * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "transpose-inner-vec4-f32-dinov2-qkv-1814x3x6x64",
+      "preset": "all",
+      "provenance": {
+        "source": "Depth Anything V2 Small backbone",
+        "notes": "Depth Anything V2 Small's stacked q/k/v permute at 518x686, repeated once per DINOv2 layer. It exercises the inner-vec4 route at the model's 1814-row shape."
+      },
+      "vars": { "count": 2089728, "dtype": "float32" },
+      "attrs": { "perm": [1, 2, 0, 3] },
+      "inputs": { "x": { "shape": [1814, 3, 6, 64], "dtype": "float32", "dist": "normal", "seed": 9411, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [3, 6, 1814, 64], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-f32-dpt-patchgrid-1813x384",
+      "preset": "all",
+      "provenance": {
+        "source": "Depth Anything V2 Small neck reassemble",
+        "notes": "Probe pair (a): the real 37x49 patch grid has 1813 rows and is not vec4-aligned. Together with pair (b), it verifies that extent parity does not determine whether the transpose is staged."
+      },
+      "vars": { "count": 696192, "dtype": "float32" },
+      "attrs": { "perm": [1, 0] },
+      "inputs": { "x": { "shape": [1813, 384], "dtype": "float32", "dist": "normal", "seed": 9441, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [384, 1813], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "transpose-f32-dpt-patchgrid-1812x384-aligned",
+      "preset": "all",
+      "provenance": {
+        "source": "Depth Anything V2 Small neck reassemble",
+        "notes": "Probe pair (b): the same transpose one row shorter, making the row extent vec4-aligned. It is the aligned control for pair (a)."
+      },
+      "vars": { "count": 695808, "dtype": "float32" },
+      "attrs": { "perm": [1, 0] },
+      "inputs": { "x": { "shape": [1812, 384], "dtype": "float32", "dist": "normal", "seed": 9451, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [384, 1812], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/datamove-flat-copy.wgsl.jinja ADDED Viewed

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

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const COUNT: u32 = {{ source.count }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let stride = nwg.x * WG;
+  for (var i = gid.x; i < COUNT; i += stride) {
+    dst[i] = src[i];
+  }
+}

build/webgpu/datamove-transpose-2d-tiled-scalar.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,57 @@

+// Tiled, optionally batched 2-D transpose over scalar elements.
+//
+// A 32x32 tile enters workgroup memory through coalesced reads and leaves
+// through coalesced writes. Guarded edges make this valid for any extent,
+// including shapes that cannot use the aligned vec4 path and would otherwise
+// fall back to a generic copy whose reads and writes cannot both be coalesced.
+//
+// Layout contract:
+// - src is row-major [batch, ROWS, COLS]
+// - dst is row-major [batch, COLS, ROWS]
+// - dispatch is x = ceil(COLS / 32), y = ceil(ROWS / 32), z = batch
+//
+// The extra column avoids bank conflicts on 32-bank GPUs.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const ROWS: u32 = {{ source.rows }}u;
+const COLS: u32 = {{ source.cols }}u;
+const MAT: u32 = {{ source.rows * source.cols }}u;
+const TILE: u32 = 32u;
+const ROW_STEP: u32 = 8u;
+var<workgroup> tile: array<array<{{ scalar }}, 33>, 32>;
+@compute @workgroup_size(32, 8, 1)
+fn main(
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let batch_base = wid.z * MAT;
+  let row_base = wid.y * TILE;
+  let col_base = wid.x * TILE;
+  // Read: lid.x runs along the source row, so a subgroup covers 32 consecutive
+  // source elements.
+  for (var i = 0u; i < TILE; i = i + ROW_STEP) {
+    let r = row_base + lid.y + i;
+    let c = col_base + lid.x;
+    if (r < ROWS && c < COLS) {
+      tile[lid.y + i][lid.x] = x[batch_base + r * COLS + c];
+    }
+  }
+  workgroupBarrier();
+  // Write: lid.x runs along the destination row, which is the source column
+  // direction — the transpose happens in the tile index, not in the memory
+  // access, so this side is coalesced too.
+  for (var i = 0u; i < TILE; i = i + ROW_STEP) {
+    let out_row = col_base + lid.y + i;
+    let out_col = row_base + lid.x;
+    if (out_row < COLS && out_col < ROWS) {
+      y[batch_base + out_row * ROWS + out_col] = tile[lid.x][lid.y + i];
+    }
+  }
+}

build/webgpu/datamove-transpose-2d-tiled.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,55 @@

+// Tiled 2D transpose (optionally batched) over vec4 elements.
+//
+// Layout contract:
+// - src is row-major [batch, ROWS, COLS], dst is row-major [batch, COLS, ROWS]
+// - ROWS % 4 == 0 and COLS % 4 == 0, so both buffers bind as vec4<T>
+// - dispatch: x = ceil(COLS / 32), y = ceil(ROWS / 32), z = batch
+//
+// Each 8x32 workgroup moves one 32x32 scalar tile through workgroup memory, so
+// both global loads and global stores are coalesced vec4 accesses. The +1
+// column pad avoids bank conflicts on 32-bank GPUs.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const ROWS: u32 = {{ source.rows }}u;
+const COLS: u32 = {{ source.cols }}u;
+const ROWS_V: u32 = {{ ((source.rows / 4) | int) }}u;
+const COLS_V: u32 = {{ ((source.cols / 4) | int) }}u;
+const MAT_V: u32 = {{ (((source.rows * source.cols) / 4) | int) }}u;
+const TILE: u32 = 32u;
+var<workgroup> tile: array<array<{{ scalar }}, 33>, 32>;
+@compute @workgroup_size(8, 32, 1)
+fn main(
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let batch_base = wid.z * MAT_V;
+  let row_base = wid.y * TILE;
+  let col_base = wid.x * TILE;
+  let r = row_base + lid.y;
+  let cv = (col_base / 4u) + lid.x;
+  if (r < ROWS && cv < COLS_V) {
+    let v = src[batch_base + r * COLS_V + cv];
+    tile[lid.y][lid.x * 4u + 0u] = v.x;
+    tile[lid.y][lid.x * 4u + 1u] = v.y;
+    tile[lid.y][lid.x * 4u + 2u] = v.z;
+    tile[lid.y][lid.x * 4u + 3u] = v.w;
+  }
+  workgroupBarrier();
+  let out_row = col_base + lid.y;
+  let rv = (row_base / 4u) + lid.x;
+  if (out_row < COLS && rv < ROWS_V) {
+    dst[batch_base + out_row * ROWS_V + rv] = {{ vectorScalar }}(
+      tile[lid.x * 4u + 0u][lid.y],
+      tile[lid.x * 4u + 1u][lid.y],
+      tile[lid.x * 4u + 2u][lid.y],
+      tile[lid.x * 4u + 3u][lid.y]
+    );
+  }
+}

build/webgpu/datamove-transpose-vec4.wgsl.jinja ADDED Viewed

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

+// Permuted copy over vec4 elements for permutations that keep the innermost
+// axis in place. Its dimension is a multiple of four, so both buffers bind as
+// vec4<T>. The index fold uses effective shapes whose last dimension is divided
+// by four, enabling 128-bit loads and stores.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const COUNT: u32 = {{ source.count }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+fn input_offset(out_index: u32) -> u32 {
+  var rem = out_index;
+  var offset = 0u;
+{% for out_axis in range(source.rank) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(out_axis + 1, source.rank) %}
+{% set dj = ((source.outputShape[j] / 4) | int) if j == source.rank - 1 else source.outputShape[j] %}
+{% set out_stride.value = out_stride.value * dj %}
+{% endfor %}
+{% set in_axis = source.perm[out_axis] %}
+{% set in_stride = namespace(value=1) %}
+{% for j in range(in_axis + 1, source.rank) %}
+{% set dj = ((source.inputShape[j] / 4) | int) if j == source.rank - 1 else source.inputShape[j] %}
+{% set in_stride.value = in_stride.value * dj %}
+{% endfor %}
+  let coord{{ out_axis }} = rem / {{ out_stride.value }}u;
+  rem = rem % {{ out_stride.value }}u;
+  offset = offset + coord{{ out_axis }} * {{ in_stride.value }}u;
+{% endfor %}
+  return offset;
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let stride = nwg.x * WG;
+  for (var i = gid.x; i < COUNT; i += stride) {
+    dst[i] = src[input_offset(i)];
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,240 @@

+{
+  "domain": "ai.onnx",
+  "name": "Transpose",
+  "sinceVersion": 13,
+  "description": "Transposes the input tensor by permuting its axes according to the `perm` attribute. Axis `i` of the output corresponds to axis `perm[i]` of the input; if `perm` is omitted, the axes are reversed (`n-1, ..., 0`).",
+  "inputs": [{ "role": "data", "dtype": "T", "description": "The input tensor to transpose." }],
+  "outputs": [
+    {
+      "role": "transposed",
+      "dtype": "T",
+      "description": "The transposed output tensor with permuted axes.",
+      "rank": "ranks.data"
+    }
+  ],
+  "attributes": {},
+  "attributeDescriptions": {
+    "perm": "Optional permutation of the input axes. It must contain every axis from 0 through rank - 1 exactly once. When omitted, the axes are reversed."
+  },
+  "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "uint32", "uint8", "int8", "bool"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "data", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "transposed", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256, "INNER_VEC4_ITEMS_PER_THREAD": 4, "INNER_VEC4_MIN_THREADS": 65536 },
+  "tunableDescriptions": {
+    "INNER_VEC4_ITEMS_PER_THREAD": "Sets the target maximum number of vec4 items assigned to each invocation by the inner-axis vector route.",
+    "INNER_VEC4_MIN_THREADS": "Sets the minimum logical thread count retained when shrinking the inner-axis vector dispatch."
+  },
+  "derive": {
+    "dtypeSupported": "f16Ok(dtypes.T)",
+    "rank0IO": "ranks.data == 0 and ranks.transposed == 0",
+    "rank1IO": "ranks.data == 1 and ranks.transposed == 1",
+    "rank2IO": "ranks.data == 2 and ranks.transposed == 2",
+    "rank3IO": "ranks.data == 3 and ranks.transposed == 3",
+    "rank4IO": "ranks.data == 4 and ranks.transposed == 4",
+    "rank5IO": "ranks.data == 5 and ranks.transposed == 5",
+    "rank6IO": "ranks.data == 6 and ranks.transposed == 6",
+    "rank7IO": "ranks.data == 7 and ranks.transposed == 7",
+    "permPresent": "has(attrs, \"perm\")",
+    "permContract": "not permPresent or ((attrs.perm | length) == ranks.data and isUniqueIntList(attrs.perm, 0, ranks.data))",
+    "rank0Permutation": "rank0IO and permContract",
+    "rank1Permutation": "rank1IO and permContract and dim(shapes.transposed, 0) == dim(shapes.data, 0)",
+    "explicitRank2": "rank2IO and permContract and permPresent and dim(shapes.transposed, 0) == dim(shapes.data, attrs.perm[0]) and dim(shapes.transposed, 1) == dim(shapes.data, attrs.perm[1])",
+    "explicitRank3": "rank3IO and permContract and permPresent and dim(shapes.transposed, 0) == dim(shapes.data, attrs.perm[0]) and dim(shapes.transposed, 1) == dim(shapes.data, attrs.perm[1]) and dim(shapes.transposed, 2) == dim(shapes.data, attrs.perm[2])",
+    "explicitRank4": "rank4IO and permContract and permPresent and dim(shapes.transposed, 0) == dim(shapes.data, attrs.perm[0]) and dim(shapes.transposed, 1) == dim(shapes.data, attrs.perm[1]) and dim(shapes.transposed, 2) == dim(shapes.data, attrs.perm[2]) and dim(shapes.transposed, 3) == dim(shapes.data, attrs.perm[3])",
+    "explicitRank5": "rank5IO and permContract and permPresent and dim(shapes.transposed, 0) == dim(shapes.data, attrs.perm[0]) and dim(shapes.transposed, 1) == dim(shapes.data, attrs.perm[1]) and dim(shapes.transposed, 2) == dim(shapes.data, attrs.perm[2]) and dim(shapes.transposed, 3) == dim(shapes.data, attrs.perm[3]) and dim(shapes.transposed, 4) == dim(shapes.data, attrs.perm[4])",
+    "explicitRank6": "rank6IO and permContract and permPresent and dim(shapes.transposed, 0) == dim(shapes.data, attrs.perm[0]) and dim(shapes.transposed, 1) == dim(shapes.data, attrs.perm[1]) and dim(shapes.transposed, 2) == dim(shapes.data, attrs.perm[2]) and dim(shapes.transposed, 3) == dim(shapes.data, attrs.perm[3]) and dim(shapes.transposed, 4) == dim(shapes.data, attrs.perm[4]) and dim(shapes.transposed, 5) == dim(shapes.data, attrs.perm[5])",
+    "explicitRank7": "rank7IO and permContract and permPresent and dim(shapes.transposed, 0) == dim(shapes.data, attrs.perm[0]) and dim(shapes.transposed, 1) == dim(shapes.data, attrs.perm[1]) and dim(shapes.transposed, 2) == dim(shapes.data, attrs.perm[2]) and dim(shapes.transposed, 3) == dim(shapes.data, attrs.perm[3]) and dim(shapes.transposed, 4) == dim(shapes.data, attrs.perm[4]) and dim(shapes.transposed, 5) == dim(shapes.data, attrs.perm[5]) and dim(shapes.transposed, 6) == dim(shapes.data, attrs.perm[6])",
+    "defaultRank2": "rank2IO and permContract and not permPresent and dim(shapes.transposed, 0) == dim(shapes.data, 1) and dim(shapes.transposed, 1) == dim(shapes.data, 0)",
+    "defaultRank3": "rank3IO and permContract and not permPresent and dim(shapes.transposed, 0) == dim(shapes.data, 2) and dim(shapes.transposed, 1) == dim(shapes.data, 1) and dim(shapes.transposed, 2) == dim(shapes.data, 0)",
+    "defaultRank4": "rank4IO and permContract and not permPresent and dim(shapes.transposed, 0) == dim(shapes.data, 3) and dim(shapes.transposed, 1) == dim(shapes.data, 2) and dim(shapes.transposed, 2) == dim(shapes.data, 1) and dim(shapes.transposed, 3) == dim(shapes.data, 0)",
+    "defaultRank5": "rank5IO and permContract and not permPresent and dim(shapes.transposed, 0) == dim(shapes.data, 4) and dim(shapes.transposed, 1) == dim(shapes.data, 3) and dim(shapes.transposed, 2) == dim(shapes.data, 2) and dim(shapes.transposed, 3) == dim(shapes.data, 1) and dim(shapes.transposed, 4) == dim(shapes.data, 0)",
+    "defaultRank6": "rank6IO and permContract and not permPresent and dim(shapes.transposed, 0) == dim(shapes.data, 5) and dim(shapes.transposed, 1) == dim(shapes.data, 4) and dim(shapes.transposed, 2) == dim(shapes.data, 3) and dim(shapes.transposed, 3) == dim(shapes.data, 2) and dim(shapes.transposed, 4) == dim(shapes.data, 1) and dim(shapes.transposed, 5) == dim(shapes.data, 0)",
+    "defaultRank7": "rank7IO and permContract and not permPresent and dim(shapes.transposed, 0) == dim(shapes.data, 6) and dim(shapes.transposed, 1) == dim(shapes.data, 5) and dim(shapes.transposed, 2) == dim(shapes.data, 4) and dim(shapes.transposed, 3) == dim(shapes.data, 3) and dim(shapes.transposed, 4) == dim(shapes.data, 2) and dim(shapes.transposed, 5) == dim(shapes.data, 1) and dim(shapes.transposed, 6) == dim(shapes.data, 0)",
+    "genericTransposeOk": "rank0Permutation or rank1Permutation or explicitRank2 or defaultRank2 or explicitRank3 or defaultRank3 or explicitRank4 or defaultRank4 or explicitRank5 or defaultRank5 or explicitRank6 or defaultRank6 or explicitRank7 or defaultRank7",
+    "defaultPerm": "pick([[rank0IO, []], [rank1IO, [0]], [rank2IO, [1, 0]], [rank3IO, [2, 1, 0]], [rank4IO, [3, 2, 1, 0]], [rank5IO, [4, 3, 2, 1, 0]], [rank6IO, [5, 4, 3, 2, 1, 0]], [rank7IO, [6, 5, 4, 3, 2, 1, 0]]], [])",
+    "genericPerm": "pick([[permPresent, attrs.perm]], defaultPerm)",
+    "innerVec4Rank3Ok": "explicitRank3 and attrs.perm[2] == 2 and dim(shapes.data, 2) % 4 == 0 and numel(shapes.transposed) > 0 and dtypeSupported",
+    "innerVec4Rank4Ok": "explicitRank4 and attrs.perm[3] == 3 and dim(shapes.data, 3) % 4 == 0 and numel(shapes.transposed) > 0 and dtypeSupported",
+    "innerVec4Ok": "innerVec4Rank3Ok or innerVec4Rank4Ok",
+    "innerVec4Perm": "attrs.perm",
+    "tiled2dShapeOk": "(defaultRank2 or (explicitRank2 and attrs.perm[0] == 1 and attrs.perm[1] == 0)) and numel(shapes.data) > 0 and ceil(dim(shapes.data, 0) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 1) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and dtypeSupported",
+    "tiled2dOk": "tiled2dShapeOk and dim(shapes.data, 0) % 4 == 0 and dim(shapes.data, 1) % 4 == 0",
+    "tiled3dShapeOk": "explicitRank3 and attrs.perm[0] == 0 and attrs.perm[1] == 2 and attrs.perm[2] == 1 and numel(shapes.data) > 0 and dim(shapes.data, 0) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 1) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 2) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and dtypeSupported",
+    "tiled3dOk": "tiled3dShapeOk and dim(shapes.data, 1) % 4 == 0 and dim(shapes.data, 2) % 4 == 0",
+    "tiled4dShapeOk": "explicitRank4 and attrs.perm[0] == 0 and attrs.perm[1] == 1 and attrs.perm[2] == 3 and attrs.perm[3] == 2 and numel(shapes.data) > 0 and dim(shapes.data, 0) * dim(shapes.data, 1) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 2) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 3) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and dtypeSupported",
+    "tiled4dOk": "tiled4dShapeOk and dim(shapes.data, 2) % 4 == 0 and dim(shapes.data, 3) % 4 == 0",
+    "tiled5dNcdhwToNdhwcShapeOk": "explicitRank5 and attrs.perm[0] == 0 and attrs.perm[1] == 2 and attrs.perm[2] == 3 and attrs.perm[3] == 4 and attrs.perm[4] == 1 and numel(shapes.data) > 0 and dim(shapes.data, 0) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 1) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceil((dim(shapes.data, 2) * dim(shapes.data, 3) * dim(shapes.data, 4)) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and dtypeSupported",
+    "tiled5dNcdhwToNdhwcOk": "tiled5dNcdhwToNdhwcShapeOk and dim(shapes.data, 1) % 4 == 0 and (dim(shapes.data, 2) * dim(shapes.data, 3) * dim(shapes.data, 4)) % 4 == 0",
+    "tiled5dNdhwcToNcdhwShapeOk": "explicitRank5 and attrs.perm[0] == 0 and attrs.perm[1] == 4 and attrs.perm[2] == 1 and attrs.perm[3] == 2 and attrs.perm[4] == 3 and numel(shapes.data) > 0 and dim(shapes.data, 0) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 1) * dim(shapes.data, 2) * dim(shapes.data, 3) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 4) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and dtypeSupported",
+    "tiled5dNdhwcToNcdhwOk": "tiled5dNdhwcToNcdhwShapeOk and dim(shapes.data, 1) * dim(shapes.data, 2) * dim(shapes.data, 3) % 4 == 0 and dim(shapes.data, 4) % 4 == 0",
+    "memcpyPerm3": "explicitRank3 and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[1]) == 1 or attrs.perm[0] < attrs.perm[1]) and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[2]) == 1 or attrs.perm[0] < attrs.perm[2]) and (dim(shapes.data, attrs.perm[1]) == 1 or dim(shapes.data, attrs.perm[2]) == 1 or attrs.perm[1] < attrs.perm[2])",
+    "memcpyPerm4": "explicitRank4 and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[1]) == 1 or attrs.perm[0] < attrs.perm[1]) and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[2]) == 1 or attrs.perm[0] < attrs.perm[2]) and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[3]) == 1 or attrs.perm[0] < attrs.perm[3]) and (dim(shapes.data, attrs.perm[1]) == 1 or dim(shapes.data, attrs.perm[2]) == 1 or attrs.perm[1] < attrs.perm[2]) and (dim(shapes.data, attrs.perm[1]) == 1 or dim(shapes.data, attrs.perm[3]) == 1 or attrs.perm[1] < attrs.perm[3]) and (dim(shapes.data, attrs.perm[2]) == 1 or dim(shapes.data, attrs.perm[3]) == 1 or attrs.perm[2] < attrs.perm[3])",
+    "memcpyPerm5": "explicitRank5 and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[1]) == 1 or attrs.perm[0] < attrs.perm[1]) and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[2]) == 1 or attrs.perm[0] < attrs.perm[2]) and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[3]) == 1 or attrs.perm[0] < attrs.perm[3]) and (dim(shapes.data, attrs.perm[0]) == 1 or dim(shapes.data, attrs.perm[4]) == 1 or attrs.perm[0] < attrs.perm[4]) and (dim(shapes.data, attrs.perm[1]) == 1 or dim(shapes.data, attrs.perm[2]) == 1 or attrs.perm[1] < attrs.perm[2]) and (dim(shapes.data, attrs.perm[1]) == 1 or dim(shapes.data, attrs.perm[3]) == 1 or attrs.perm[1] < attrs.perm[3]) and (dim(shapes.data, attrs.perm[1]) == 1 or dim(shapes.data, attrs.perm[4]) == 1 or attrs.perm[1] < attrs.perm[4]) and (dim(shapes.data, attrs.perm[2]) == 1 or dim(shapes.data, attrs.perm[3]) == 1 or attrs.perm[2] < attrs.perm[3]) and (dim(shapes.data, attrs.perm[2]) == 1 or dim(shapes.data, attrs.perm[4]) == 1 or attrs.perm[2] < attrs.perm[4]) and (dim(shapes.data, attrs.perm[3]) == 1 or dim(shapes.data, attrs.perm[4]) == 1 or attrs.perm[3] < attrs.perm[4])",
+    "memcpyPermOk": "memcpyPerm3 or memcpyPerm4 or memcpyPerm5",
+    "identityPermutation": "rank1Permutation or (explicitRank2 and attrs.perm[0] == 0 and attrs.perm[1] == 1) or (explicitRank3 and attrs.perm[0] == 0 and attrs.perm[1] == 1 and attrs.perm[2] == 2) or (explicitRank4 and attrs.perm[0] == 0 and attrs.perm[1] == 1 and attrs.perm[2] == 2 and attrs.perm[3] == 3) or (explicitRank5 and attrs.perm[0] == 0 and attrs.perm[1] == 1 and attrs.perm[2] == 2 and attrs.perm[3] == 3 and attrs.perm[4] == 4) or (explicitRank6 and attrs.perm[0] == 0 and attrs.perm[1] == 1 and attrs.perm[2] == 2 and attrs.perm[3] == 3 and attrs.perm[4] == 4 and attrs.perm[5] == 5)",
+    "identityCopyOk": "numel(shapes.transposed) > 0 and numel(shapes.transposed) % 4 == 0 and (identityPermutation or memcpyPermOk) and dtypeSupported",
+    "tiled4dNchwToNhwcShapeOk": "explicitRank4 and attrs.perm[0] == 0 and attrs.perm[1] == 2 and attrs.perm[2] == 3 and attrs.perm[3] == 1 and numel(shapes.data) > 0 and dim(shapes.data, 0) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 1) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceil((dim(shapes.data, 2) * dim(shapes.data, 3)) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and dtypeSupported",
+    "tiled4dNchwToNhwcOk": "tiled4dNchwToNhwcShapeOk and dim(shapes.data, 1) % 4 == 0 and (dim(shapes.data, 2) * dim(shapes.data, 3)) % 4 == 0",
+    "tiled4dNhwcToNchwShapeOk": "explicitRank4 and attrs.perm[0] == 0 and attrs.perm[1] == 3 and attrs.perm[2] == 1 and attrs.perm[3] == 2 and numel(shapes.data) > 0 and dim(shapes.data, 0) <= device.limits.maxComputeWorkgroupsPerDimension and ceil((dim(shapes.data, 1) * dim(shapes.data, 2)) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceil(dim(shapes.data, 3) / 32) <= device.limits.maxComputeWorkgroupsPerDimension and dtypeSupported",
+    "tiled4dNhwcToNchwOk": "tiled4dNhwcToNchwShapeOk and (dim(shapes.data, 1) * dim(shapes.data, 2)) % 4 == 0 and dim(shapes.data, 3) % 4 == 0",
+    "tiledTransposeOk": "tiled2dOk or tiled3dOk or tiled4dOk or tiled5dNcdhwToNdhwcOk or tiled5dNdhwcToNcdhwOk or tiled4dNchwToNhwcOk or tiled4dNhwcToNchwOk",
+    "tiledScalarTransposeOk": "tiled2dShapeOk or tiled3dShapeOk or tiled4dShapeOk or tiled5dNcdhwToNdhwcShapeOk or tiled5dNdhwcToNcdhwShapeOk or tiled4dNchwToNhwcShapeOk or tiled4dNhwcToNchwShapeOk",
+    "tiledRows": "dim(shapes.data, 0) if tiled2dShapeOk else (dim(shapes.data, 1) if tiled3dShapeOk else (dim(shapes.data, 2) if tiled4dShapeOk else (dim(shapes.data, 1) if tiled5dNcdhwToNdhwcShapeOk else (dim(shapes.data, 1) * dim(shapes.data, 2) * dim(shapes.data, 3) if tiled5dNdhwcToNcdhwShapeOk else (dim(shapes.data, 1) if tiled4dNchwToNhwcShapeOk else (dim(shapes.data, 1) * dim(shapes.data, 2) if tiled4dNhwcToNchwShapeOk else 1))))))",
+    "tiledCols": "dim(shapes.data, 1) if tiled2dShapeOk else (dim(shapes.data, 2) if tiled3dShapeOk else (dim(shapes.data, 3) if tiled4dShapeOk else (dim(shapes.data, 2) * dim(shapes.data, 3) * dim(shapes.data, 4) if tiled5dNcdhwToNdhwcShapeOk else (dim(shapes.data, 4) if tiled5dNdhwcToNcdhwShapeOk else (dim(shapes.data, 2) * dim(shapes.data, 3) if tiled4dNchwToNhwcShapeOk else (dim(shapes.data, 3) if tiled4dNhwcToNchwShapeOk else 1))))))",
+    "tiledBatch": "1 if tiled2dShapeOk else (dim(shapes.data, 0) if tiled3dShapeOk else (dim(shapes.data, 0) * dim(shapes.data, 1) if tiled4dShapeOk else (dim(shapes.data, 0) if tiled5dNcdhwToNdhwcShapeOk or tiled5dNdhwcToNcdhwShapeOk or tiled4dNchwToNhwcShapeOk or tiled4dNhwcToNchwShapeOk else 1)))"
+  },
+  "constants": { "usesF16": "dtypes.T == \"f16\"" },
+  "bindingSets": {
+    "vectorData": [
+      {
+        "name": "src",
+        "arg": "x",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "dst",
+        "arg": "y",
+        "semantic": "transposed",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      }
+    ],
+    "scalarData": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "transposed", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.transposed)" }]
+        }
+      }
+    ],
+    "scalarDataNoParams": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "data",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "transposed", "buffer": { "type": "storage" }, "elementType": "$scalar" }
+    ]
+  },
+  "variants": [
+    {
+      "id": "tiled_vec4",
+      "priority": 22,
+      "when": ["tiledTransposeOk"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Transpose.tiled",
+          "source": {
+            "shader": "datamove-transpose-2d-tiled.wgsl.jinja",
+            "inputs": { "rows": "tiledRows", "cols": "tiledCols" }
+          },
+          "bindings": "vectorData",
+          "dispatch": { "x": "ceil(tiledCols / 32)", "y": "ceil(tiledRows / 32)", "z": "tiledBatch" }
+        }
+      ]
+    },
+    {
+      "id": "tiled_scalar",
+      "description": "Stages ragged two-dimensional transpose planes through a padded workgroup tile so reads and writes remain coalesced when vec4 alignment is unavailable.",
+      "priority": 19,
+      "when": ["tiledScalarTransposeOk", "dtypeSupported"],
+      "constants": { "scalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Transpose.tiledScalar",
+          "source": {
+            "shader": "datamove-transpose-2d-tiled-scalar.wgsl.jinja",
+            "inputs": { "rows": "tiledRows", "cols": "tiledCols" }
+          },
+          "bindings": "scalarDataNoParams",
+          "dispatch": { "x": "ceil(tiledCols / 32)", "y": "ceil(tiledRows / 32)", "z": "tiledBatch" }
+        }
+      ]
+    },
+    {
+      "id": "identity_vec4",
+      "priority": 21,
+      "when": ["identityCopyOk"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Transpose.identityVec4",
+          "source": { "shader": "datamove-flat-copy.wgsl.jinja", "inputs": { "count": "numel(shapes.transposed) / 4" } },
+          "bindings": "vectorData",
+          "dispatch": { "gridStride": "numel(shapes.transposed) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "inner_vec4",
+      "description": "Vectorizes the contiguous innermost dimension for permutations that leave that axis in place. It assigns multiple vectors per invocation only when enough invocations remain to keep the dispatch populated.",
+      "priority": 20,
+      "when": ["innerVec4Ok"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Transpose.innerVec4",
+          "source": {
+            "shader": "datamove-transpose-vec4.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.data",
+              "outputShape": "shapes.transposed",
+              "rank": "ranks.data",
+              "perm": "innerVec4Perm",
+              "count": "numel(shapes.transposed) / 4"
+            }
+          },
+          "bindings": "vectorData",
+          "dispatch": {
+            "gridStride": "max(min(numel(shapes.transposed) / 4, tunables.INNER_VEC4_MIN_THREADS), ceilDiv(numel(shapes.transposed) / 4, tunables.INNER_VEC4_ITEMS_PER_THREAD))",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "generic",
+      "when": ["genericTransposeOk", "dtypeSupported"],
+      "constants": { "scalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Transpose",
+          "source": {
+            "shader": "transpose.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.data",
+              "outputShape": "shapes.transposed",
+              "rank": "ranks.data",
+              "perm": "genericPerm"
+            }
+          },
+          "bindings": "scalarData",
+          "dispatch": { "threads": "numel(shapes.transposed)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,22 @@

+{
+  "name": "ai.onnx.Transpose",
+  "id": "_ai_onnx_transpose_webgpu_dff5d66",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "wT5OzWQiJ0OSKnc+JuSj0vOoFzY1Q/JzgJW1jpIbv5Q=",
+      "datamove-flat-copy.wgsl.jinja": "rEpeYMvIZQuQnB2f3lt51jrrklKICsqUTzCJPRTknNk=",
+      "datamove-transpose-2d-tiled-scalar.wgsl.jinja": "zy7rlN8Se+mn8RB9akbBTHq1RK8ot6NBboxG3vX1Xdg=",
+      "datamove-transpose-2d-tiled.wgsl.jinja": "AtmY7NlBNCV1unIf2epvpBO3hfnLBmgGKGB7WiWMCQk=",
+      "datamove-transpose-vec4.wgsl.jinja": "RgwjTbu7SHDVsUW7W2f5S81gv1UJBpT93udQX/MZdJ0=",
+      "manifest.json": "qACTV3+3jqKTescFk98we7Nm2pwajZoBwGAo1n3dL9M=",
+      "test.json": "acNlC/DcnpoXb3ni/Z4QMbpYD7C1YR16iHXX20mNp0Y=",
+      "transpose.wgsl.jinja": "BbRpj9r38KShHUF3RXf518/bYtZ8BHFbGDZLRbOuP9g="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Transpose" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1097 @@

+{
+  "op": "ai.onnx.Transpose",
+  "fixtureArrays": {
+    "onnx_backend_transpose_input_x": [0.54881352186203, 0.7151893377304077, 0.6027633547782898, 0.5448831915855408, 0.42365479469299316, 0.6458941102027893, 0.4375872015953064, 0.891772985458374, 0.9636627435684204, 0.3834415078163147, 0.7917250394821167, 0.5288949012756348, 0.5680445432662964, 0.9255966544151306, 0.07103605568408966, 0.08712930232286453, 0.020218396559357643, 0.832619845867157, 0.7781567573547363, 0.8700121641159058, 0.978618323802948, 0.7991585731506348, 0.4614793658256531, 0.7805292010307312],
+    "ort_transpose021_f32_input_x": [1, 2, 3, 4, 5, 6, 1.1, 2.1, 3.1, 4.1, 5.1, 6.1, 1.2, 2.2, 3.2, 4.2, 5.2, 6.2, 1.3, 2.3, 3.3, 4.3, 5.3, 6.3],
+    "ort_nhwc_to_nchw_uint32_input_x": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16]
+  },
+  "cases": [
+    {
+      "name": "f32_subnormal_rank2_preserve_data",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim",
+        "notes": "Transpose is data movement with reindexing; selected finite subnormal values must survive the permutation."
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 0.0, 1e-39, -1e-39, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank1_default_perm",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "values", "values": [0.0, 1.0, -2.0, 3.5, 4.0, -5.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank1_explicit_identity_perm",
+      "attrs": { "perm": [0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "values", "values": [0.0, 1.0, -2.0, 3.5, 4.0, -5.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "matrix_2d",
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "matrix_2d_default_perm",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 10.0, 11.0, 12.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank3_0_2_1",
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank3_default_perm",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.27 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank4_attention_layout",
+      "attrs": { "perm": [0, 2, 1, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 3, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank4_default_perm",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.21 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [5, 4, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank5_default_perm",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 3, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 1, 2, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_default_perm",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 3, 1, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 10.0, 11.0, 20.0, 21.0, 100.0, 101.0, 110.0, 111.0, 120.0, 121.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3, 1, 2, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_explicit_interleaved_perm",
+      "attrs": { "perm": [0, 2, 4, 1, 3, 5] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 1, 4, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2, 3, 4, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank0_scalar_default_perm",
+      "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [12.25] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank0_scalar_explicit_empty_perm",
+      "attrs": { "perm": [] },
+      "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-7.5] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "zero_dim_middle_to_last_axis_empty_output",
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_two_dim_default_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDimNoAttr"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_two_dim_default_bool",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDimNoAttr",
+        "notes": "Same two-dimensional data-movement fixture with an ONNX-valid bool payload."
+      },
+      "inputs": {
+        "x": { "dtype": "bool", "shape": [2, 3], "data": { "kind": "values", "values": [1, 0, 1, 0, 1, 0] } }
+      },
+      "outputs": { "y": { "dtype": "bool", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_two_dim_default_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDimNoAttr"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [3, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_two_dim_explicit_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim"
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_two_dim_explicit_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim_mlfloat16",
+        "notes": "Uses normal f16 numeric values instead of ORT's raw MLFloat16 bit-pattern payloads."
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [3, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_two_dim_uint32_adapted",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim_int32",
+        "notes": "ORT covers int32; this framework covers the same 2D explicit transpose with supported uint32 data."
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6] } }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_transpose021_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.Transpose021"
+      },
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 2, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_transpose021_f32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_transpose120_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.Transpose120"
+      },
+      "attrs": { "perm": [1, 2, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 2, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_transpose021_f32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_transpose102_suffix_preserved_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.Transpose102"
+      },
+      "attrs": { "perm": [1, 0, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 2, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_transpose021_f32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_transpose_reshape_rank5_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TransposeReshape"
+      },
+      "attrs": { "perm": [1, 3, 2, 4, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 1, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_transpose021_f32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 1, 2, 3, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_six_dim_swap_outer_axes_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.SixDim"
+      },
+      "attrs": { "perm": [1, 0, 2, 3, 4, 5] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0, 33.0, 34.0, 35.0, 36.0, 37.0, 38.0, 39.0, 40.0, 41.0, 42.0, 43.0, 44.0, 45.0, 46.0, 47.0, 48.0, 49.0, 50.0, 51.0, 52.0, 53.0, 54.0, 55.0, 56.0, 57.0, 58.0, 59.0, 60.0, 61.0, 62.0, 63.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_nchw_to_nhwc_uint32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.NCHW2NHWC"
+      },
+      "attrs": { "perm": [0, 2, 3, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [1, 2, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_nhwc_to_nchw_uint32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.NHWC2NCHW"
+      },
+      "attrs": { "perm": [0, 3, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_nhwc_to_nchw_uint32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_two_dim_int32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim_int32"
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_two_dim_int16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim_int16"
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": { "dtype": "int16", "shape": [2, 3], "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6] } }
+      },
+      "outputs": { "y": { "dtype": "int16", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_single_axis_moving_inwards_block_copy_uint32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.SingleAxisMovingInwardsBlockCopy",
+        "notes": "ORT uses uint64 payload values; this framework stores the representable test values as uint32."
+      },
+      "attrs": { "perm": [1, 2, 0, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_nhwc_to_nchw_uint32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_do_transpose_impl_rank4_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.DoTransposeImpl"
+      },
+      "attrs": { "perm": [2, 1, 0, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [5, 2, 1, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_transpose3dimpl_flatten_dims23_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.Transpose3DImpl",
+        "notes": "Compact projection of ORT's rank-4 [0,2,3,1] case that flattens source dims 2 and 3 into one moved block."
+      },
+      "attrs": { "perm": [0, 2, 3, 1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 3, 4, 5], "data": { "kind": "fillFloat32" } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 5, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_transpose3dimpl_flatten_dims12_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.Transpose3DImpl",
+        "notes": "Compact projection of ORT's rank-4 [0,3,1,2] case that flattens source dims 1 and 2 into one moved block."
+      },
+      "attrs": { "perm": [0, 3, 1, 2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 3, 4, 5], "data": { "kind": "fillFloat32" } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_ndim_rank4_swap_outer_axes_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.NDim"
+      },
+      "attrs": { "perm": [1, 0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_nhwc_to_nchw_uint32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_ndim_rank4_swap_outer_and_inner_axes_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.NDim"
+      },
+      "attrs": { "perm": [1, 0, 3, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_nhwc_to_nchw_uint32_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_transpose_all_permutations_0",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_transpose_all_permutations_0" },
+      "attrs": { "perm": [0, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_transpose_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_transpose_all_permutations_1",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_transpose_all_permutations_1" },
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_transpose_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_transpose_all_permutations_2",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_transpose_all_permutations_2" },
+      "attrs": { "perm": [1, 0, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_transpose_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_transpose_all_permutations_3",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_transpose_all_permutations_3" },
+      "attrs": { "perm": [1, 2, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_transpose_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_transpose_all_permutations_4",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_transpose_all_permutations_4" },
+      "attrs": { "perm": [2, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_transpose_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_transpose_all_permutations_5",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_transpose_all_permutations_5" },
+      "attrs": { "perm": [2, 1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_transpose_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_transpose_default",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_transpose_default" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_transpose_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_two_dim_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim_int8",
+        "notes": "Two-dimensional transpose with logical int8 edge values."
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [2, 3], "data": { "kind": "values", "values": [-128, -1, 0, 1, 2, 127] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_transpose021_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.Transpose021",
+        "notes": "Rank-3 non-default permutation with logical uint8 edge values."
+      },
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 250, 251, 252, 253, 254, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "tiled_2d_f32_64x32",
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [64, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "tiled_2d_f32_36x20_partial_tile",
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [36, 20],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [20, 36], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "tiled_2d_f32_256x256",
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 256], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "tiled_2d_f16_32x64",
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [64, 32], "tolerance": 0.001 } }
+    },
+    {
+      "name": "tiled_2d_u32_8x12",
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [8, 12],
+          "data": { "kind": "cycle", "values": [7, 1, 9, 4, 3, 12, 5, 0, 2, 8, 6, 11, 10] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [12, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "tiled_2d_default_perm_f32_16x44",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [16, 44],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [44, 16], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "tiled_3d_batched_f32_3x36x12",
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 36, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 12, 36], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "tiled_3d_batched_f16_2x32x16",
+      "attrs": { "perm": [0, 2, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 16, 32], "tolerance": 0.001 } }
+    },
+    {
+      "name": "tiled_4d_batched_f32_2x3x16x8",
+      "attrs": { "perm": [0, 1, 3, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 16, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 8, 16], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "identity_vec4_rank1_f32_12",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [12], "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [12], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "identity_vec4_rank2_f32_6x10",
+      "attrs": { "perm": [0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [6, 10], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "identity_vec4_rank3_f16_2x3x4",
+      "attrs": { "perm": [0, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3, 4], "tolerance": 0.001 } }
+    },
+    {
+      "name": "identity_vec4_rank4_f32_1x2x2x4",
+      "attrs": { "perm": [0, 1, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "inner_vec4_rank3_f32_5x3x8",
+      "attrs": { "perm": [1, 0, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [5, 3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 5, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "inner_vec4_rank3_f16_4x6x12",
+      "attrs": { "perm": [1, 0, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 6, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [6, 4, 12], "tolerance": 0.001 } }
+    },
+    {
+      "name": "inner_vec4_rank4_attention_f32_2x3x4x8",
+      "attrs": { "perm": [0, 2, 1, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 3, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "inner_vec4_rank4_f32_2x3x4x8_perm2103",
+      "attrs": { "perm": [2, 1, 0, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 3, 2, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank2_scalar_odd_width_5x7",
+      "provenance": {
+        "notes": "Compact correctness sibling for the odd-width rank2 scalar transpose benchmark; the dimensions avoid vec4/tiled alignment while preserving the same rank2 permutation."
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [5, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [7, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank7_explicit_perm",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest (generic N-dim)",
+        "notes": "Rank-7 explicit-permutation coverage for the generic transpose path. The manifest supplies the rank and permutation as template inputs instead of maintaining a rank-specific variant."
+      },
+      "attrs": { "perm": [1, 0, 3, 2, 5, 4, 6] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 2, 2, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [-10.0, -9.5, -9.0, -8.5, -8.0, -7.5, -7.0, -6.5, -6.0, -5.5, -5.0, -4.5, -4.0, -3.5, -3.0, -2.5, -2.0, -1.5, -1.0, -0.5, 0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0, 5.5, 6.0, 6.5, 7.0, 7.5, 8.0, 8.5, 9.0, 9.5, 10.0, 10.5, 11.0, 11.5, 12.0, 12.5, 13.0, 13.5, 14.0, 14.5, 15.0, 15.5, 16.0, 16.5, 17.0, 17.5, 18.0, 18.5, 19.0, 19.5, 20.0, 20.5, 21.0, 21.5, 22.0, 22.5, 23.0, 23.5, 24.0, 24.5, 25.0, 25.5, 26.0, 26.5, 27.0, 27.5, 28.0, 28.5, 29.0, 29.5, 30.0, 30.5, 31.0, 31.5, 32.0, 32.5, 33.0, 33.5, 34.0, 34.5, 35.0, 35.5, 36.0, 36.5, 37.0, 37.5, 38.0, 38.5, 39.0, 39.5, 40.0, 40.5, 41.0, 41.5, 42.0, 42.5, 43.0, 43.5, 44.0, 44.5, 45.0, 45.5, 46.0, 46.5, 47.0, 47.5, 48.0, 48.5, 49.0, 49.5, 50.0, 50.5, 51.0, 51.5, 52.0, 52.5, 53.0, 53.5, 54.0, 54.5, 55.0, 55.5, 56.0, 56.5, 57.0, 57.5, 58.0, 58.5, 59.0, 59.5, 60.0, 60.5, 61.0, 61.5, 62.0, 62.5, 63.0, 63.5, 64.0, 64.5, 65.0, 65.5, 66.0, 66.5, 67.0, 67.5, 68.0, 68.5, 69.0, 69.5, 70.0, 70.5, 71.0, 71.5, 72.0, 72.5, 73.0, 73.5, 74.0, 74.5, 75.0, 75.5, 76.0, 76.5, 77.0, 77.5, 78.0, 78.5, 79.0, 79.5, 80.0, 80.5, 81.0, 81.5, 82.0, 82.5, 83.0, 83.5, 84.0, 84.5, 85.0, 85.5]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2, 2, 2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank2_bool_explicit_tiled",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TwoDim (bool payload)",
+        "notes": "bool is in typeConstraints T and scalarType maps it to u32, so an aligned [4,4] bool transpose selects tiled_vec4 (gate is only f16Ok, true for u32). The pre-existing ort_two_dim_default_bool_gpu_gap fixture is stale-named; this NON-gpu-gapped case gives bool its first real GPU validation on the vec4 path."
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "bool",
+          "shape": [4, 4],
+          "data": { "kind": "values", "values": [1, 0, 1, 0, 0, 1, 1, 0, 1, 1, 0, 1, 0, 0, 1, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "bool", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank5_scrambled_perm_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/transpose_test.cc",
+        "test": "TransposeOpTest.TransposeReshape (scrambled rank5)",
+        "notes": "Genuinely scrambled rank-5 explicit perm [4,2,0,3,1] exercising the generic scalar-gather path with rank-specific template inputs; prior rank-5 coverage only exercised a default reversal and a near-reshape perm."
+      },
+      "attrs": { "perm": [4, 2, 0, 3, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 3, 2],
+          "data": {
+            "kind": "values",
+            "values": [-5.0, -4.75, -4.5, -4.25, -4.0, -3.75, -3.5, -3.25, -3.0, -2.75, -2.5, -2.25, -2.0, -1.75, -1.5, -1.25, -1.0, -0.75, -0.5, -0.25, 0.0, 0.25, 0.5, 0.75, 1.0, 1.25, 1.5, 1.75, 2.0, 2.25, 2.5, 2.75, 3.0, 3.25, 3.5, 3.75, 4.0, 4.25, 4.5, 4.75, 5.0, 5.25, 5.5, 5.75, 6.0, 6.25, 6.5, 6.75, 7.0, 7.25, 7.5, 7.75, 8.0, 8.25, 8.5, 8.75, 9.0, 9.25, 9.5, 9.75, 10.0, 10.25, 10.5, 10.75, 11.0, 11.25, 11.5, 11.75, 12.0, 12.25, 12.5, 12.75]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank7_default_full_reversal",
+      "provenance": {
+        "notes": "A rank-7 default transpose (no perm) reverses all seven dimensions through the generic transpose path."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 3],
+          "data": { "kind": "linspace", "start": -1.0, "end": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 1, 2, 1, 2, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank1_2d_fold_boundary_last_element_guard",
+      "provenance": {
+        "notes": "numel = 2^24+1 = 16777217 forces ceilDiv(numel,256)=65536 > maxComputeWorkgroupsPerDimension(65535), engaging the scalar path's y-axis 2D fold (i = gid.x + gid.y*nwg.x*WG). rank1 numel%4 != 0 bypasses identity_vec4 and lands on the generic scalar path, isolating the fold last-row guard. Referenceless: TS reference supplies truth so a corrupted/dropped tail element in the second y-row fails the compare."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [16777217],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [16777217], "tolerance": 0 } }
+    },
+    {
+      "name": "tiled_2d_vec4_naninf_signedzero_payload_preserved",
+      "provenance": {
+        "notes": "Aligned [8,8] perm[1,0] f32 selects tiled_vec4 (both dims %4==0, %32 partial). Verifies NaN / +-Inf / signed-zero survive the vec4->scalar->workgroup->vec4 staging bit-exact. Expected values are the [8,8] input transposed (column-major read)."
+      },
+      "attrs": { "perm": [1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8, 8],
+          "data": {
+            "kind": "values",
+            "values": ["NaN", "Infinity", "-Infinity", 0.0, 0.0, 1.0, -1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0, 33.0, 34.0, 35.0, 36.0, 37.0, 38.0, 39.0, 40.0, 41.0, 42.0, 43.0, 44.0, 45.0, 46.0, 47.0, 48.0, 49.0, 50.0, 51.0, 52.0, 53.0, 54.0, 55.0, 56.0, 57.0, 58.0]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [8, 8],
+          "allowNaN": true,
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": ["NaN", 3.0, 11.0, 19.0, 27.0, 35.0, 43.0, 51.0, "Infinity", 4.0, 12.0, 20.0, 28.0, 36.0, 44.0, 52.0, "-Infinity", 5.0, 13.0, 21.0, 29.0, 37.0, 45.0, 53.0, 0.0, 6.0, 14.0, 22.0, 30.0, 38.0, 46.0, 54.0, 0.0, 7.0, 15.0, 23.0, 31.0, 39.0, 47.0, 55.0, 1.0, 8.0, 16.0, 24.0, 32.0, 40.0, 48.0, 56.0, -1.0, 9.0, 17.0, 25.0, 33.0, 41.0, 49.0, 57.0, 2.0, 10.0, 18.0, 26.0, 34.0, 42.0, 50.0, 58.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "scalar_rank4_default_reverse_naninf_preserved",
+      "provenance": {
+        "notes": "No perm => default full reversal [3,2,1,0] on the generic scalar path. Values placed so index [i,j,k,l] holds 8*i+4*j+2*k+l; output[l,k,j,i] = input[i,j,k,l]. NaN/Inf/-0 payloads at select positions must survive the scalar input_offset gather."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": ["NaN", 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, "Infinity", "-Infinity", 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 0.0]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2],
+          "allowNaN": true,
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": ["NaN", "-Infinity", 4.0, 12.0, 2.0, 10.0, 6.0, 14.0, 1.0, 9.0, 5.0, 13.0, 3.0, 11.0, "Infinity", 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "rank5_ncdhw_to_ndhwc_tiled_vec4",
+      "provenance": {
+        "notes": "Compact route lock for the flattened batched-2D NCDHW-to-NDHWC tile. C=4 and D*H*W=32 exercise one complete spatial tile boundary while preserving exact copy semantics."
+      },
+      "attrs": { "perm": [0, 2, 3, 4, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 4, 4],
+          "data": { "kind": "linspace", "start": -64.0, "end": 63.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 4, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank5_ndhwc_to_ncdhw_tiled_candidate",
+      "provenance": {
+        "notes": "Inverse rank-5 layout lock: NDHWC [N,D,H,W,C] becomes NCDHW via perm [0,4,1,2,3]. C=4 and D*H*W=32 match one flattened 2D transpose boundary with bit-exact copy semantics."
+      },
+      "attrs": { "perm": [0, 4, 1, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 4, 4, 4],
+          "data": { "kind": "linspace", "start": -64.0, "end": 63.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 2, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank4_nchw_to_nhwc_perm0231_tiled",
+      "provenance": {
+        "source": "onnxruntime/core/providers/webgpu/tensor/transpose.cc",
+        "test": "channels-last conversions use the tiled shared-memory kernel",
+        "notes": "perm (0,2,3,1)/(0,3,1,2) are batched 2D transposes with merged row/col dims and must reach the tiled kernel, not the scalar generic path."
+      },
+      "attrs": { "perm": [0, 2, 3, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8, 4, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 12, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "rank4_nhwc_to_nchw_perm0312_tiled",
+      "provenance": {
+        "source": "onnxruntime/core/providers/webgpu/tensor/transpose.cc",
+        "test": "channels-last conversions use the tiled shared-memory kernel",
+        "notes": "perm (0,2,3,1)/(0,3,1,2) are batched 2D transposes with merged row/col dims and must reach the tiled kernel, not the scalar generic path."
+      },
+      "attrs": { "perm": [0, 3, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4, 12, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8, 4, 12], "tolerance": 0 } }
+    },
+    {
+      "name": "rank4_unitdim_memcpy_perm0213",
+      "provenance": {
+        "source": "onnxruntime/core/providers/webgpu/tensor/transpose.cc",
+        "test": "unit-dim permutations lower to a flat copy",
+        "notes": "A permutation that only moves size-1 axes preserves the row-major layout, so the flat vec4 copy is exact."
+      },
+      "attrs": { "perm": [0, 2, 1, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 6, 1, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 6, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "rank3_unitdim_memcpy_perm201",
+      "provenance": {
+        "source": "onnxruntime/core/providers/webgpu/tensor/transpose.cc",
+        "test": "unit-dim permutations lower to a flat copy",
+        "notes": "A permutation that only moves size-1 axes preserves the row-major layout, so the flat vec4 copy is exact."
+      },
+      "attrs": { "perm": [2, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 8, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 6, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "rank5_unitdim_memcpy_perm20134",
+      "provenance": {
+        "source": "onnxruntime/core/providers/webgpu/tensor/transpose.cc",
+        "test": "unit-dim permutations lower to a flat copy",
+        "notes": "A permutation that only moves size-1 axes preserves the row-major layout, so the flat vec4 copy is exact."
+      },
+      "attrs": { "perm": [2, 0, 1, 3, 4] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.07 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank4_nonmemcpy_perm0213_canary",
+      "provenance": {
+        "source": "onnxruntime/core/providers/webgpu/tensor/transpose.cc",
+        "test": "unit-dim permutations lower to a flat copy",
+        "notes": "A permutation that only moves size-1 axes preserves the row-major layout, so the flat vec4 copy is exact."
+      },
+      "attrs": { "perm": [0, 2, 1, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.11 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 3, 4], "tolerance": 0 } }
+    }
+  ]
+}

build/webgpu/transpose.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,38 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+fn input_offset({% if source.rank > 0 %}out_index: u32{% endif %}) -> u32 {
+{% if source.rank > 0 %}
+  var rem = out_index;
+{% endif %}
+  var offset = 0u;
+{% for out_axis in range(source.rank) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(out_axis + 1, source.rank) %}
+{% set out_stride.value = out_stride.value * source.outputShape[j] %}
+{% endfor %}
+{% set out_stride_divisor = out_stride.value if out_stride.value != 0 else 1 %}
+{% set in_axis = source.perm[out_axis] %}
+{% set in_stride = namespace(value=1) %}
+{% for j in range(in_axis + 1, source.rank) %}
+{% set in_stride.value = in_stride.value * source.inputShape[j] %}
+{% endfor %}
+  let coord{{ out_axis }} = rem / {{ out_stride_divisor }}u;
+  rem = rem % {{ out_stride_divisor }}u;
+  offset = offset + coord{{ out_axis }} * {{ in_stride.value }}u;
+{% endfor %}
+  return offset;
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i >= params.count) {
+    return;
+  }
+  y[i] = x[input_offset({% if source.rank > 0 %}i{% endif %})];
+}