sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +70 -0
build/webgpu/bench.json +94 -0
build/webgpu/datamove-tile-vec4.wgsl.jinja +52 -0
build/webgpu/manifest.json +115 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +508 -0
build/webgpu/tile.wgsl.jinja +38 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Tile
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Constructs a tensor by tiling a given tensor: each dimension `i` of the input is repeated `repeats[i]` times, so `output_dim[i] = input_dim[i] * repeats[i]`. Equivalent to NumPy `tile` but without broadcasting.
+See the [ONNX `Tile` spec](https://onnx.ai/onnx/operators/onnx__Tile.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | — | — | Input tensor of any shape. | required |
+| `repeats` | `repeats` | `S` | `uint32` | `1` | — | Logical int64 1-D tensor of length equal to the input rank, specifying non-negative repeat counts stored as uint32 by WebGPU. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | same as `input` | — | Output tensor of the same type as the input, with each dimension scaled by the corresponding repeat count. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool` |
+| `S` | `int64` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`datamove-tile-vec4.wgsl.jinja`](build/webgpu/datamove-tile-vec4.wgsl.jinja)
+- [`tile.wgsl.jinja`](build/webgpu/tile.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
+The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
+- `output`
+Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Tile", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({
+  input: { data: inputData, shape: [2, 1, 3] },
+  repeats: { data: repeatsData, shape: [3] },
+}, {
+  outputs: { output: { shape: [2, 1, 3], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,94 @@

+{
+  "op": "ai.onnx.Tile",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "tile-vec4-rows-f32-8x512x1024-repeat4",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "outCount": 16777216 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [8, 512, 1024], "dist": "normal", "seed": 878, "scale": 1 },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 512, 1024] } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.outCount * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "rank3",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 128, 64] },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 128, 64] } }
+    },
+    {
+      "name": "rank1_large_output_from_scalar",
+      "preset": "edge",
+      "vars": { "dtype": "float32", "count": 1048576 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 1 },
+        "repeats": { "dtype": "uint32", "shape": [1], "dist": "constant", "value": 1048576 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1048576] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "rank5_uniform_repeats",
+      "preset": "edge",
+      "vars": { "dtype": "float32", "count": 65536 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 8, 1, 8, 1], "dist": "normal", "seed": 877, "scale": 1 },
+        "repeats": { "dtype": "uint32", "shape": [5], "dist": "constant", "value": 4 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 32, 4, 32, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "tile-vec4-rows-f16-8x512x1024-repeat4",
+      "preset": "edge",
+      "vars": { "dtype": "float16", "outCount": 16777216 },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [8, 512, 1024], "dist": "normal", "seed": 879, "scale": 1 },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [32, 512, 1024] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.outCount * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "tile-lastdim-not-mul4-scalar-cliff-f32-8x512x1023-repeat4",
+      "preset": "edge",
+      "vars": { "dtype": "float32", "outCount": 16760832 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [8, 512, 1023], "dist": "normal", "seed": 880, "scale": 1 },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 512, 1023] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.outCount * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "tile-lastaxis-repeat-broadcast-fallback-f32-16x2048-repeat4",
+      "preset": "edge",
+      "vars": { "dtype": "float32", "outCount": 131072 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [16, 2048], "dist": "normal", "seed": 881, "scale": 1 },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [16, 8192] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.outCount * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "tile-generic-rank6-deep-fold-lastaxis-repeat-f32-4x4x4x4x4x64",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "outCount": 131072 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4, 4, 4, 4, 4, 64], "dist": "normal", "seed": 5301, "scale": 1 },
+        "repeats": { "dtype": "uint32", "shape": [6], "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4, 4, 4, 4, 128] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.outCount * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

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

	@@ -0,0 +1,52 @@

+// Tile over vec4 elements for repeats that leave the innermost axis unchanged.
+// The equal innermost dimensions are multiples of four, so both buffers bind as
+// vec4<T>. The per-axis modulo fold uses effective shapes whose last dimension
+// is divided by four, enabling 128-bit loads and stores. Axes with repeat == 1
+// skip the modulo entirely.
+{% 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 axis in range(source.rank) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(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 safe_out_stride = out_stride.value %}
+{% set input_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.rank) %}
+{% set dj = ((source.inputShape[j] / 4) | int) if j == source.rank - 1 else source.inputShape[j] %}
+{% set input_stride.value = input_stride.value * dj %}
+{% endfor %}
+{% set in_dim = ((source.inputShape[axis] / 4) | int) if axis == source.rank - 1 else source.inputShape[axis] %}
+{% set out_dim = ((source.outputShape[axis] / 4) | int) if axis == source.rank - 1 else source.outputShape[axis] %}
+{% set safe_in_dim = in_dim %}
+  let out_coord{{ axis }} = rem / {{ safe_out_stride }}u;
+  rem = rem % {{ safe_out_stride }}u;
+{% if in_dim == out_dim %}
+  offset = offset + out_coord{{ axis }} * {{ input_stride.value }}u;
+{% else %}
+  offset = offset + (out_coord{{ axis }} % {{ safe_in_dim }}u) * {{ input_stride.value }}u;
+{% endif %}
+{% 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) {
+    output[i] = input[input_offset(i)];
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,115 @@

+{
+  "domain": "ai.onnx",
+  "name": "Tile",
+  "sinceVersion": 13,
+  "description": "Constructs a tensor by tiling a given tensor: each dimension `i` of the input is repeated `repeats[i]` times, so `output_dim[i] = input_dim[i] * repeats[i]`. Equivalent to NumPy `tile` but without broadcasting.",
+  "inputs": [
+    { "role": "input", "dtype": "T", "description": "Input tensor of any shape." },
+    {
+      "role": "repeats",
+      "dtype": "S",
+      "description": "Logical int64 1-D tensor of length equal to the input rank, specifying non-negative repeat counts stored as uint32 by WebGPU.",
+      "rank": 1
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "description": "Output tensor of the same type as the input, with each dimension scaled by the corresponding repeat count.",
+      "rank": "ranks.input"
+    }
+  ],
+  "typeConstraints": {
+    "T": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"],
+    "S": ["int64"]
+  },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "repeats": { "kind": "tensor", "semantic": "repeats", "role": "repeats", "dtype": "uint32", "narrowing": "checked" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "constants": { "usesF16": "dtypes.T == \"f16\"" },
+  "variants": [
+    {
+      "id": "inner_vec4",
+      "priority": 15,
+      "when": ["ranks.input == ranks.output", "ranks.input >= 1", "ranks.repeats == 1", "dim(shapes.repeats, 0) == ranks.input", "numel(shapes.output) > 0", "dim(shapes.input, ranks.input - 1) == dim(shapes.output, ranks.input - 1)", "dim(shapes.input, ranks.input - 1) % 4 == 0", "f16Ok(dtypes.T)"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Tile.innerVec4",
+          "source": {
+            "shader": "datamove-tile-vec4.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.input",
+              "outputShape": "shapes.output",
+              "rank": "ranks.input",
+              "count": "numel(shapes.output) / 4"
+            }
+          },
+          "bindings": [
+            {
+              "name": "input",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            }
+          ],
+          "dispatch": { "gridStride": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "generic",
+      "when": ["ranks.input == ranks.output", "ranks.repeats == 1", "dim(shapes.repeats, 0) == ranks.input", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Tile",
+          "source": {
+            "shader": "tile.wgsl.jinja",
+            "inputs": { "inputShape": "shapes.input", "outputShape": "shapes.output", "rank": "ranks.input" }
+          },
+          "bindings": [
+            {
+              "name": "input",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Tile",
+  "id": "_ai_onnx_tile_webgpu_d5edcce",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "hvniK7wzBksemRGQsj2QyLX/PI60/fACENqZBziTjuM=",
+      "datamove-tile-vec4.wgsl.jinja": "CcDLTkiLHX6FWfMQCqjsFx8l5I7A/hjpKXFNTD3yMU4=",
+      "manifest.json": "BNBWgAaPI8XFTYVR7q3r0zDz2zipeHZwSoaHysMeEmA=",
+      "test.json": "jVFErSBn7O/ciOZlrxjMCymY4EkBVpwYcElrr+Q/Ew4=",
+      "tile.wgsl.jinja": "JsV5aQYvf0xn8l27Ads1lZJpZSPQ1umQuk9B2QC5JT0="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Tile" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,508 @@

+{
+  "op": "ai.onnx.Tile",
+  "cases": [
+    {
+      "name": "rank2",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 6] } }
+    },
+    {
+      "name": "ort_float_2d_two_axes_full_repeat",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileFloatType",
+        "notes": "RunTest<float>({2, 2}, {2, 2}) from ORT's shared Tile wrapper."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_float_2d_first_axis_repeat",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileFloatType",
+        "notes": "RunTest<float>({2, 3}, {2, 1}) from ORT's shared Tile wrapper."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_f16_2d_two_axes_full_repeat",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileMLFloat16Type",
+        "notes": "RunTest<MLFloat16>({2, 2}, {2, 2}) from ORT's shared Tile wrapper."
+      },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 4], "tolerance": 0.001 } }
+    },
+    {
+      "name": "rank3_f16",
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [1, 2, 2] },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 2, 4] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "rank5_nonuniform_repeats",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11 }
+        },
+        "repeats": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [2, 1, 3, 1, 4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3, 2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "large_output_from_scalar",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [7.0] } },
+        "repeats": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [257] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [257], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_nonuniform_repeats",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19 }
+        },
+        "repeats": { "dtype": "uint32", "shape": [6], "data": { "kind": "values", "values": [2, 1, 3, 1, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank0_scalar_empty_repeats_noop",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [42.5] } },
+        "repeats": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "zero_repeat_axis_empty_output",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [0, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_1d_zero_repeats",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileInt32Type",
+        "notes": "RunTest<int32_t>({3}, {0})."
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 1, 2] } },
+        "repeats": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_2d_two_axes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileInt32Type",
+        "notes": "RunTest<int32_t>({2, 2}, {2, 2})."
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, 2, 3] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int16_2d_two_axes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileInt16Type",
+        "notes": "RunTest<int16_t>({2, 2}, {2, 2})."
+      },
+      "inputs": {
+        "input": { "dtype": "int16", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, 2, 3] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "int16", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_3d_middle_repeat",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileUint32Type",
+        "notes": "RunTest<uint32_t>({2, 1, 3}, {1, 2, 1})."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [2, 1, 3], "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5] } },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_large_values_no_float_cast",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileUint32Type",
+        "notes": "Uses the ORT tile copy pattern with large uint32 values to guard against f32 round trips."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [16777217, 4000000000] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 2] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_singleton_prefix_copy_like",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileInt32Type",
+        "notes": "RunTest<int32_t>({1, 1, 3}, {2, 2, 1})."
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [-7, 0, 9] } },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [2, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_singleton_batch_copy_like",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileInt32Type",
+        "notes": "RunTest<int32_t>({1, 1, 3}, {3, 1, 1})."
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [-7, 0, 9] } },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [3, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_batched_copy_like",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileInt32Type",
+        "notes": "RunTest<int32_t>({2, 1, 3}, {2, 2, 1})."
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [2, 1, 3], "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5] } },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_1d_repeat3_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileFloatType",
+        "notes": "RunTest<float>({3}, {3})."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 1.0, 2.0] } },
+        "repeats": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [9], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_2d_second_axis_zero_repeat",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileFloatType",
+        "notes": "RunTest<float>({2, 2}, {2, 0})."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_rank3_all_repeats_one_noop",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileFloatType",
+        "notes": "RunTest<float>({2, 1, 3}, {1, 1, 1})."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_4d_repeated_middle_spatial",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileFloatType",
+        "notes": "RunTest<float>({1, 2, 3, 4}, {2, 1, 2, 1})."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 4],
+          "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]
+          }
+        },
+        "repeats": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [2, 1, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 6, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_rank5_float_nonuniform_repeats",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileFloatType",
+        "notes": "RunTest<float>({2, 3, 2, 3, 2}, {2, 1, 2, 1, 2})."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 3, 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, 64.0, 65.0, 66.0, 67.0, 68.0, 69.0, 70.0, 71.0]
+          }
+        },
+        "repeats": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [2, 1, 2, 1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3, 4, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_tile_precomputed",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_tile_precomputed",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_tile",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_tile",
+        "notes": "ONNX repeats are int64 metadata; this framework stores non-negative repeats as uint32 metadata slots."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [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, 0.11827442795038223, 0.6399210095405579, 0.14335328340530396, 0.9446688890457153, 0.5218483209609985, 0.4146619439125061, 0.26455560326576233, 0.7742336988449097, 0.4561503231525421, 0.568433940410614, 0.018789799883961678, 0.6176354885101318, 0.6120957136154175, 0.6169340014457703, 0.9437480568885803, 0.681820273399353, 0.35950788855552673, 0.43703195452690125, 0.6976311802864075, 0.0602254718542099, 0.6667667031288147, 0.670637845993042, 0.21038256585597992, 0.12892629206180573, 0.31542834639549255, 0.36371076107025146, 0.5701967477798462, 0.4386015236377716, 0.9883738160133362, 0.10204481333494186, 0.20887675881385803, 0.16130951046943665, 0.6531082987785339, 0.25329160690307617, 0.4663107693195343, 0.24442559480667114, 0.15896958112716675, 0.11037514358758926, 0.6563295722007751, 0.13818295300006866, 0.1965823620557785, 0.3687251806259155, 0.8209932446479797, 0.09710127860307693, 0.8379449248313904, 0.0960984081029892, 0.9764594435691833, 0.4686512053012848, 0.9767611026763916, 0.6048455238342285, 0.7392635941505432, 0.03918779268860817, 0.28280696272850037, 0.12019655853509903, 0.296140193939209, 0.11872772127389908, 0.3179831802845001, 0.414262980222702, 0.06414749473333359, 0.6924721002578735, 0.5666014552116394, 0.26538950204849243, 0.5232480764389038, 0.09394051134586334, 0.5759465098381042, 0.9292961955070496, 0.3185689449310303, 0.6674103736877441, 0.13179786503314972, 0.7163271903991699, 0.28940609097480774, 0.18319135904312134, 0.5865129232406616, 0.02010754682123661, 0.8289400339126587, 0.004695476032793522, 0.6778165102005005, 0.2700079679489136, 0.7351940274238586, 0.9621885418891907, 0.2487531453371048, 0.5761573314666748, 0.5920419096946716, 0.5722519159317017, 0.22308163344860077, 0.9527490139007568, 0.4471253752708435, 0.8464086651802063, 0.6994792819023132, 0.2974369525909424, 0.8137978315353394, 0.396505743265152, 0.8811032176017761, 0.5812729001045227, 0.8817353844642639, 0.6925315856933594]
+          }
+        },
+        "repeats": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [7, 6, 4, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [14, 18, 16, 10], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_2d_two_axes_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileInt8Type",
+        "notes": "RunTest<int8_t>({2, 2}, {2, 2}) with signed-byte edge values."
+      },
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [2, 2], "data": { "kind": "values", "values": [-128, -1, 0, 127] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_3d_middle_repeat_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileUint8Type",
+        "notes": "RunTest<uint8_t>({2, 1, 3}, {1, 2, 1}) with high logical-byte values."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "uint8",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0, 1, 127, 128, 254, 255] }
+        },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [2, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "inner_vec4_rank3_f16",
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [2, 3, 8] },
+        "repeats": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 6, 8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "inner_vec4_u32_large_values",
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [16777217, 4000000000, 1, 2, 3, 4, 5, 4294967295] }
+        },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [6, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_bool_2d_two_axes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileBoolType",
+        "notes": "RunTestForBool Tile2D_2Axes."
+      },
+      "inputs": {
+        "input": { "dtype": "bool", "shape": [2, 2], "data": { "kind": "values", "values": [1, 0, 1, 0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [4, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0] }
+        }
+      }
+    },
+    {
+      "name": "ort_bool_2d_zero_repeat_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/tile_op_test.cc",
+        "test": "TensorOpTest.TileBoolType",
+        "notes": "RunTestForBool Tile2DWithZeroRepeats."
+      },
+      "inputs": {
+        "input": { "dtype": "bool", "shape": [2, 2], "data": { "kind": "values", "values": [1, 0, 1, 0] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 0] } }
+      },
+      "outputs": {
+        "output": { "dtype": "bool", "shape": [4, 0], "data": { "kind": "values", "values": [] }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "rank7_unit_input_alternating_repeats",
+      "provenance": {
+        "source": "ONNX Tile spec (no rank limit) + onnxruntime CPU Tile kernel",
+        "notes": "Spec-valid rank-7 Tile. ORT CPU computes it via its generic strided copy."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 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]
+          }
+        },
+        "repeats": { "dtype": "uint32", "shape": [7], "data": { "kind": "values", "values": [1, 2, 1, 2, 1, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 2, 2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "inner_vec4_rank4_f16_middle_repeats",
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [2, 3, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.13 }
+        },
+        "repeats": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [2, 1, 3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 3, 6, 8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "inner_vec4_rank7_lastdim_mul4_outer_repeats",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.23 }
+        },
+        "repeats": { "dtype": "uint32", "shape": [7], "data": { "kind": "values", "values": [2, 1, 2, 1, 2, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 1, 4, 1, 4, 1, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "generic_rank6_lastaxis_repeat_broadcast_fold",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 3, 1, 2, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        },
+        "repeats": { "dtype": "uint32", "shape": [6], "data": { "kind": "values", "values": [2, 3, 1, 2, 1, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3, 3, 2, 2, 15], "tolerance": 0 } }
+    },
+    {
+      "name": "inner_vec4_int8_lastdim_mul4_outer_repeat",
+      "inputs": {
+        "input": {
+          "dtype": "int8",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": [-128, 127, 0, -1, 1, -2, 64, -64, 100, -100, 7, -7, 3, -3, 15, -15] }
+        },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [6, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "rank8_unit_input_alternating_repeats",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 2, 1],
+          "data": { "kind": "linspace", "start": 0.0, "end": 15.0 }
+        },
+        "repeats": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [1, 2, 1, 2, 1, 2, 1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 2, 2, 2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_input_axis_generic",
+      "provenance": {
+        "notes": "An empty INPUT axis, as opposed to the zero-repeat cases which empty the OUTPUT while every input extent stays positive. The generic route divides the output coordinate by the input extent on each axis, and only this shape drives that extent to zero; the vec4 route cannot, because its guard requires a non-empty output."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [0, 3], "data": { "kind": "values", "values": [] } },
+        "repeats": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [0, 6], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    }
+  ]
+}

build/webgpu/tile.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 axis in range(source.rank) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.rank) %}
+{% set out_stride.value = out_stride.value * source.outputShape[j] %}
+{% endfor %}
+{% set safe_out_stride = 1 if out_stride.value == 0 else out_stride.value %}
+{% set input_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.rank) %}
+{% set input_stride.value = input_stride.value * source.inputShape[j] %}
+{% endfor %}
+{% set safe_input_dim = 1 if source.inputShape[axis] == 0 else source.inputShape[axis] %}
+  let out_coord{{ axis }} = rem / {{ safe_out_stride }}u;
+  rem = rem % {{ safe_out_stride }}u;
+  offset = offset + (out_coord{{ axis }} % {{ safe_input_dim }}u) * {{ input_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;
+  }
+  output[i] = input[input_offset({% if source.rank > 0 %}i{% endif %})];
+}