sync 2e7068faf55e
Browse files- README.md +89 -0
- build/webgpu/bench.json +173 -0
- build/webgpu/manifest.json +483 -0
- build/webgpu/metadata.json +25 -0
- build/webgpu/resize-antialias.wgsl.jinja +159 -0
- build/webgpu/resize-coord-transform-5d.wgsl.jinja +224 -0
- build/webgpu/resize-coord-transform.wgsl.jinja +313 -0
- build/webgpu/resize-cubic.wgsl.jinja +118 -0
- build/webgpu/resize-generic.wgsl.jinja +409 -0
- build/webgpu/resize-linear-2x-stencil-x8.wgsl.jinja +81 -0
- build/webgpu/resize-linear-2x-stencil.wgsl.jinja +124 -0
- build/webgpu/resize-nearest-integer-scale.wgsl.jinja +73 -0
- build/webgpu/test.json +0 -0
README.md
CHANGED
|
@@ -1,3 +1,92 @@
|
|
| 1 |
---
|
|
|
|
| 2 |
license: apache-2.0
|
|
|
|
|
|
|
|
|
|
|
|
|
| 3 |
---
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
---
|
| 2 |
+
library_name: kernels
|
| 3 |
license: apache-2.0
|
| 4 |
+
tags:
|
| 5 |
+
- kernel
|
| 6 |
+
- webgpu
|
| 7 |
+
- wgsl
|
| 8 |
---
|
| 9 |
+
# ai.onnx.Resize
|
| 10 |
+
|
| 11 |
+
`ai.onnx` · standard ONNX operator · ONNX opset ≥ 19
|
| 12 |
+
|
| 13 |
+
## Description
|
| 14 |
+
|
| 15 |
+
Resizes the input tensor by sampling neighboring input values. Output dimensions are determined by per-axis scale factors or explicit target sizes. Supports `nearest`, `linear`, and `cubic` interpolation with configurable coordinate transformation modes.
|
| 16 |
+
|
| 17 |
+
See the [ONNX `Resize` spec](https://onnx.ai/onnx/operators/onnx__Resize.html) for the reference semantics.
|
| 18 |
+
|
| 19 |
+
## Inputs
|
| 20 |
+
|
| 21 |
+
| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
|
| 22 |
+
| --- | --- | --- | --- | --- | --- | --- |
|
| 23 |
+
| `X` | `x` | `T` | — | — | N-D input tensor to be resized. | required |
|
| 24 |
+
|
| 25 |
+
## Outputs
|
| 26 |
+
|
| 27 |
+
| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
|
| 28 |
+
| --- | --- | --- | --- | --- | --- | --- |
|
| 29 |
+
| `Y` | `y` | `T` | same as `X` | — | N-D output tensor after resizing. | required |
|
| 30 |
+
|
| 31 |
+
## Attributes
|
| 32 |
+
|
| 33 |
+
Attributes and default values (overridable per request):
|
| 34 |
+
|
| 35 |
+
| Attribute | Default | Description |
|
| 36 |
+
| --- | --- | --- |
|
| 37 |
+
| `mode` | `"nearest"` | Interpolation mode: `"nearest"` (default), `"linear"` (bilinear/N-linear), or `"cubic"` (bicubic/N-cubic). |
|
| 38 |
+
| `coordinate_transformation_mode` | `"half_pixel"` | How to map a coordinate in the output tensor back to the input tensor; supported modes include `"half_pixel"`, `"pytorch_half_pixel"`, `"align_corners"`, `"asymmetric"`, and `"tf_crop_and_resize"`. |
|
| 39 |
+
| `nearest_mode` | `"round_prefer_floor"` | Rounding strategy used when `mode` is `"nearest"`: `"round_prefer_floor"` (default), `"round_prefer_ceil"`, `"floor"`, or `"ceil"`. |
|
| 40 |
+
| `cubic_coeff_a` | `-0.75` | Coefficient `a` in the cubic interpolation filter; the default is -0.75. Valid only when `mode` is `"cubic"`. |
|
| 41 |
+
| `roi` | `[]` | Values of the optional `roi` tensor, supplied as `[starts..., ends...]` with one pair per input axis or per `axes` entry; used only by `"tf_crop_and_resize"`. |
|
| 42 |
+
| `scales` | `[]` | Values of the optional `scales` tensor. Supply one positive value per input axis, or one per `axes` entry when that attribute is present; omit it when the output shape represents the exact `sizes` input. |
|
| 43 |
+
| `extrapolation_value` | `0` | Fill value used for out-of-bounds samples when `coordinate_transformation_mode` is `"tf_crop_and_resize"`. |
|
| 44 |
+
| `antialias` | `0` | When set to 1, stretches the resampling filter during downscaling so that more input pixels contribute to each output pixel, reducing aliasing. |
|
| 45 |
+
| `exclude_outside` | `0` | When set to 1, zero-weights sampling locations that fall outside the input tensor and renormalizes the remaining weights to sum to 1. |
|
| 46 |
+
| `keep_aspect_ratio_policy` | `"stretch"` | How explicit sizes preserve aspect ratio. Only the default `"stretch"` policy is supported; `"not_larger"` and `"not_smaller"` are not yet implemented. |
|
| 47 |
+
| `axes` | — | Optional axes to resize. When omitted, the scale values or requested output size apply to every input axis. |
|
| 48 |
+
|
| 49 |
+
## Type constraints
|
| 50 |
+
|
| 51 |
+
| Variable | Allowed dtypes |
|
| 52 |
+
| --- | --- |
|
| 53 |
+
| `T` | `float32`, `float16`, `uint8`, `int8` |
|
| 54 |
+
|
| 55 |
+
## Files
|
| 56 |
+
|
| 57 |
+
- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
|
| 58 |
+
- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
|
| 59 |
+
- [`test.json`](build/webgpu/test.json) — correctness cases
|
| 60 |
+
- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
|
| 61 |
+
- [`resize-antialias.wgsl.jinja`](build/webgpu/resize-antialias.wgsl.jinja)
|
| 62 |
+
- [`resize-coord-transform-5d.wgsl.jinja`](build/webgpu/resize-coord-transform-5d.wgsl.jinja)
|
| 63 |
+
- [`resize-coord-transform.wgsl.jinja`](build/webgpu/resize-coord-transform.wgsl.jinja)
|
| 64 |
+
- [`resize-cubic.wgsl.jinja`](build/webgpu/resize-cubic.wgsl.jinja)
|
| 65 |
+
- [`resize-generic.wgsl.jinja`](build/webgpu/resize-generic.wgsl.jinja)
|
| 66 |
+
- [`resize-linear-2x-stencil-x8.wgsl.jinja`](build/webgpu/resize-linear-2x-stencil-x8.wgsl.jinja)
|
| 67 |
+
- [`resize-linear-2x-stencil.wgsl.jinja`](build/webgpu/resize-linear-2x-stencil.wgsl.jinja)
|
| 68 |
+
- [`resize-nearest-integer-scale.wgsl.jinja`](build/webgpu/resize-nearest-integer-scale.wgsl.jinja)
|
| 69 |
+
|
| 70 |
+
## Use with `@huggingface/kernels`
|
| 71 |
+
|
| 72 |
+
The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
|
| 73 |
+
|
| 74 |
+
The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
|
| 75 |
+
|
| 76 |
+
- `y`
|
| 77 |
+
|
| 78 |
+
Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
|
| 79 |
+
|
| 80 |
+
The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
|
| 81 |
+
|
| 82 |
+
Replace each `*Data` placeholder with a typed array containing the corresponding input data.
|
| 83 |
+
|
| 84 |
+
```js
|
| 85 |
+
import { getKernel } from "@huggingface/kernels";
|
| 86 |
+
|
| 87 |
+
const kernel = await getKernel("webgpu-kernels/ai.onnx.Resize", { version: 1 });
|
| 88 |
+
// Explicit destinations request optional results or supply metadata that cannot be inferred.
|
| 89 |
+
const { y } = await kernel({ x: { data: xData, shape: [1, 1, 2, 4] } }, {
|
| 90 |
+
outputs: { y: { shape: [1, 1, 1, 2], dtype: "float32" } },
|
| 91 |
+
});
|
| 92 |
+
```
|
build/webgpu/bench.json
ADDED
|
@@ -0,0 +1,173 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
{
|
| 2 |
+
"op": "ai.onnx.Resize",
|
| 3 |
+
"tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
|
| 4 |
+
"cases": [
|
| 5 |
+
{
|
| 6 |
+
"name": "nchw_1x64_256_to_512_linear_decoder",
|
| 7 |
+
"preset": "smoke",
|
| 8 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "half_pixel" },
|
| 9 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 64, 256, 256], "dist": "normal", "seed": 731, "scale": 1 } },
|
| 10 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 64, 512, 512] } },
|
| 11 |
+
"bench": {
|
| 12 |
+
"primary": true,
|
| 13 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 14 |
+
}
|
| 15 |
+
},
|
| 16 |
+
{
|
| 17 |
+
"name": "nchw_1x64_256_to_512_linear_asymmetric_decoder",
|
| 18 |
+
"preset": "smoke",
|
| 19 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "asymmetric" },
|
| 20 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 64, 256, 256], "dist": "normal", "seed": 732, "scale": 1 } },
|
| 21 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 64, 512, 512] } },
|
| 22 |
+
"bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }] }
|
| 23 |
+
},
|
| 24 |
+
{
|
| 25 |
+
"name": "nchw_1x3_256_to_512_linear",
|
| 26 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "half_pixel" },
|
| 27 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 3, 256, 256] } },
|
| 28 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 3, 512, 512] } },
|
| 29 |
+
"bench": {
|
| 30 |
+
"primary": true,
|
| 31 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 32 |
+
}
|
| 33 |
+
},
|
| 34 |
+
{
|
| 35 |
+
"name": "nchw_1x32_256_to_512_linear_align_corners_large",
|
| 36 |
+
"preset": "smoke",
|
| 37 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "align_corners" },
|
| 38 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 32, 256, 256], "dist": "normal", "seed": 751, "scale": 1 } },
|
| 39 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 32, 512, 512] } },
|
| 40 |
+
"bench": {
|
| 41 |
+
"primary": true,
|
| 42 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 43 |
+
}
|
| 44 |
+
},
|
| 45 |
+
{
|
| 46 |
+
"name": "nchw_1x3_256_to_512_linear_align_corners_fallback",
|
| 47 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "align_corners" },
|
| 48 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 3, 256, 256] } },
|
| 49 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 3, 512, 512] } },
|
| 50 |
+
"bench": {
|
| 51 |
+
"primary": true,
|
| 52 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 53 |
+
}
|
| 54 |
+
},
|
| 55 |
+
{
|
| 56 |
+
"name": "nchw_1x3_256_to_512_nearest_2x",
|
| 57 |
+
"attrs": { "mode": "nearest", "coordinate_transformation_mode": "asymmetric", "nearest_mode": "floor" },
|
| 58 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 3, 256, 256] } },
|
| 59 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 3, 512, 512] } },
|
| 60 |
+
"bench": {
|
| 61 |
+
"primary": true,
|
| 62 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 63 |
+
}
|
| 64 |
+
},
|
| 65 |
+
{
|
| 66 |
+
"name": "nchw_1x128_64_to_128_nearest_2x_f16_vae",
|
| 67 |
+
"preset": "smoke",
|
| 68 |
+
"attrs": {
|
| 69 |
+
"mode": "nearest",
|
| 70 |
+
"coordinate_transformation_mode": "asymmetric",
|
| 71 |
+
"nearest_mode": "floor",
|
| 72 |
+
"scales": [1, 1, 2, 2]
|
| 73 |
+
},
|
| 74 |
+
"inputs": { "x": { "dtype": "float16", "shape": [1, 128, 64, 64], "dist": "normal", "seed": 601, "scale": 0.2 } },
|
| 75 |
+
"outputs": { "y": { "dtype": "float16", "shape": [1, 128, 128, 128] } },
|
| 76 |
+
"bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 2" }] }
|
| 77 |
+
},
|
| 78 |
+
{
|
| 79 |
+
"name": "nchw_1x3_176_to_528_nearest_3x",
|
| 80 |
+
"attrs": { "mode": "nearest", "coordinate_transformation_mode": "asymmetric", "nearest_mode": "floor" },
|
| 81 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 3, 176, 176] } },
|
| 82 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 3, 528, 528] } },
|
| 83 |
+
"bench": {
|
| 84 |
+
"primary": true,
|
| 85 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 86 |
+
}
|
| 87 |
+
},
|
| 88 |
+
{
|
| 89 |
+
"name": "nchw_1x3_128_to_512_nearest_4x",
|
| 90 |
+
"attrs": {
|
| 91 |
+
"mode": "nearest",
|
| 92 |
+
"coordinate_transformation_mode": "half_pixel",
|
| 93 |
+
"nearest_mode": "round_prefer_floor"
|
| 94 |
+
},
|
| 95 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 3, 128, 128] } },
|
| 96 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 3, 512, 512] } },
|
| 97 |
+
"bench": {
|
| 98 |
+
"primary": true,
|
| 99 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 100 |
+
}
|
| 101 |
+
},
|
| 102 |
+
{
|
| 103 |
+
"name": "nchw_1x64_256_to_510_linear_w_odd_scalar_cliff",
|
| 104 |
+
"preset": "smoke",
|
| 105 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "half_pixel" },
|
| 106 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 64, 256, 256], "dist": "normal", "seed": 781, "scale": 1 } },
|
| 107 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 64, 512, 510] } },
|
| 108 |
+
"bench": {
|
| 109 |
+
"primary": true,
|
| 110 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 111 |
+
}
|
| 112 |
+
},
|
| 113 |
+
{
|
| 114 |
+
"name": "nchw_1x64_256_to_512_linear_f16_vec4_decoder",
|
| 115 |
+
"preset": "smoke",
|
| 116 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "half_pixel" },
|
| 117 |
+
"inputs": { "x": { "dtype": "float16", "shape": [1, 64, 256, 256], "dist": "normal", "seed": 782, "scale": 1 } },
|
| 118 |
+
"outputs": { "y": { "dtype": "float16", "shape": [1, 64, 512, 512] } },
|
| 119 |
+
"bench": {
|
| 120 |
+
"primary": true,
|
| 121 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 2" }]
|
| 122 |
+
}
|
| 123 |
+
},
|
| 124 |
+
{
|
| 125 |
+
"name": "nchw_1x64_256_to_510_linear_f16_w_odd_scalar_cliff",
|
| 126 |
+
"preset": "smoke",
|
| 127 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "half_pixel" },
|
| 128 |
+
"inputs": { "x": { "dtype": "float16", "shape": [1, 64, 256, 256], "dist": "normal", "seed": 783, "scale": 1 } },
|
| 129 |
+
"outputs": { "y": { "dtype": "float16", "shape": [1, 64, 512, 510] } },
|
| 130 |
+
"bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 2" }] }
|
| 131 |
+
},
|
| 132 |
+
{
|
| 133 |
+
"name": "nchw_1x64_512_to_256_linear_downsample_vec4",
|
| 134 |
+
"preset": "smoke",
|
| 135 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "half_pixel" },
|
| 136 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 64, 512, 512], "dist": "normal", "seed": 784, "scale": 1 } },
|
| 137 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 64, 256, 256] } },
|
| 138 |
+
"bench": {
|
| 139 |
+
"primary": true,
|
| 140 |
+
"metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }]
|
| 141 |
+
}
|
| 142 |
+
},
|
| 143 |
+
{
|
| 144 |
+
"name": "antialias_linear_8x_downsample_tap_explosion_stress",
|
| 145 |
+
"preset": "stress",
|
| 146 |
+
"attrs": {
|
| 147 |
+
"mode": "linear",
|
| 148 |
+
"antialias": 1,
|
| 149 |
+
"coordinate_transformation_mode": "half_pixel",
|
| 150 |
+
"scales": [1, 1, 0.125, 0.125]
|
| 151 |
+
},
|
| 152 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 32, 256, 256], "dist": "normal", "seed": 901, "scale": 1 } },
|
| 153 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 32, 32, 32], "dist": "empty" } },
|
| 154 |
+
"bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }] }
|
| 155 |
+
},
|
| 156 |
+
{
|
| 157 |
+
"name": "generic_nhwc_linear_2x_upsample_scalar_no_vec4_stress",
|
| 158 |
+
"preset": "stress",
|
| 159 |
+
"attrs": { "mode": "linear", "coordinate_transformation_mode": "half_pixel" },
|
| 160 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 256, 256, 16], "dist": "normal", "seed": 902, "scale": 1 } },
|
| 161 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 512, 512, 16], "dist": "empty" } },
|
| 162 |
+
"bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }] }
|
| 163 |
+
},
|
| 164 |
+
{
|
| 165 |
+
"name": "generic_rank5_cubic_64tap_unroll_stress",
|
| 166 |
+
"preset": "stress",
|
| 167 |
+
"attrs": { "mode": "cubic", "coordinate_transformation_mode": "half_pixel", "cubic_coeff_a": -0.75 },
|
| 168 |
+
"inputs": { "x": { "dtype": "float32", "shape": [1, 1, 32, 32, 32], "dist": "normal", "seed": 903, "scale": 1 } },
|
| 169 |
+
"outputs": { "y": { "dtype": "float32", "shape": [1, 1, 64, 64, 64], "dist": "empty" } },
|
| 170 |
+
"bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.x) + numel(shapes.y)) * 4" }] }
|
| 171 |
+
}
|
| 172 |
+
]
|
| 173 |
+
}
|
build/webgpu/manifest.json
ADDED
|
@@ -0,0 +1,483 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
{
|
| 2 |
+
"domain": "ai.onnx",
|
| 3 |
+
"name": "Resize",
|
| 4 |
+
"sinceVersion": 19,
|
| 5 |
+
"description": "Resizes the input tensor by sampling neighboring input values. Output dimensions are determined by per-axis scale factors or explicit target sizes. Supports `nearest`, `linear`, and `cubic` interpolation with configurable coordinate transformation modes.",
|
| 6 |
+
"inputs": [{ "role": "X", "dtype": "T", "description": "N-D input tensor to be resized." }],
|
| 7 |
+
"outputs": [{ "role": "Y", "dtype": "T", "description": "N-D output tensor after resizing.", "rank": "ranks.X" }],
|
| 8 |
+
"attributes": {
|
| 9 |
+
"mode": "nearest",
|
| 10 |
+
"coordinate_transformation_mode": "half_pixel",
|
| 11 |
+
"nearest_mode": "round_prefer_floor",
|
| 12 |
+
"cubic_coeff_a": -0.75,
|
| 13 |
+
"roi": [],
|
| 14 |
+
"scales": [],
|
| 15 |
+
"extrapolation_value": 0,
|
| 16 |
+
"antialias": 0,
|
| 17 |
+
"exclude_outside": 0,
|
| 18 |
+
"keep_aspect_ratio_policy": "stretch"
|
| 19 |
+
},
|
| 20 |
+
"attributeDescriptions": {
|
| 21 |
+
"mode": "Interpolation mode: `\"nearest\"` (default), `\"linear\"` (bilinear/N-linear), or `\"cubic\"` (bicubic/N-cubic).",
|
| 22 |
+
"coordinate_transformation_mode": "How to map a coordinate in the output tensor back to the input tensor; supported modes include `\"half_pixel\"`, `\"pytorch_half_pixel\"`, `\"align_corners\"`, `\"asymmetric\"`, and `\"tf_crop_and_resize\"`.",
|
| 23 |
+
"nearest_mode": "Rounding strategy used when `mode` is `\"nearest\"`: `\"round_prefer_floor\"` (default), `\"round_prefer_ceil\"`, `\"floor\"`, or `\"ceil\"`.",
|
| 24 |
+
"cubic_coeff_a": "Coefficient `a` in the cubic interpolation filter; the default is -0.75. Valid only when `mode` is `\"cubic\"`.",
|
| 25 |
+
"roi": "Values of the optional `roi` tensor, supplied as `[starts..., ends...]` with one pair per input axis or per `axes` entry; used only by `\"tf_crop_and_resize\"`.",
|
| 26 |
+
"scales": "Values of the optional `scales` tensor. Supply one positive value per input axis, or one per `axes` entry when that attribute is present; omit it when the output shape represents the exact `sizes` input.",
|
| 27 |
+
"extrapolation_value": "Fill value used for out-of-bounds samples when `coordinate_transformation_mode` is `\"tf_crop_and_resize\"`.",
|
| 28 |
+
"antialias": "When set to 1, stretches the resampling filter during downscaling so that more input pixels contribute to each output pixel, reducing aliasing.",
|
| 29 |
+
"exclude_outside": "When set to 1, zero-weights sampling locations that fall outside the input tensor and renormalizes the remaining weights to sum to 1.",
|
| 30 |
+
"axes": "Optional axes to resize. When omitted, the scale values or requested output size apply to every input axis.",
|
| 31 |
+
"keep_aspect_ratio_policy": "How explicit sizes preserve aspect ratio. Only the default `\"stretch\"` policy is supported; `\"not_larger\"` and `\"not_smaller\"` are not yet implemented."
|
| 32 |
+
},
|
| 33 |
+
"attributeConstraints": {
|
| 34 |
+
"mode": { "values": ["nearest", "linear", "cubic"] },
|
| 35 |
+
"coordinate_transformation_mode": {
|
| 36 |
+
"values": ["half_pixel", "half_pixel_symmetric", "pytorch_half_pixel", "align_corners", "asymmetric", "tf_crop_and_resize"]
|
| 37 |
+
},
|
| 38 |
+
"nearest_mode": { "values": ["round_prefer_floor", "round_prefer_ceil", "floor", "ceil"] },
|
| 39 |
+
"antialias": { "values": [0, 1] },
|
| 40 |
+
"exclude_outside": { "values": [0, 1] },
|
| 41 |
+
"keep_aspect_ratio_policy": { "values": ["stretch"] }
|
| 42 |
+
},
|
| 43 |
+
"typeConstraints": { "T": ["float32", "float16", "uint8", "int8"] },
|
| 44 |
+
"args": {
|
| 45 |
+
"x": { "kind": "tensor", "semantic": "X", "role": "input" },
|
| 46 |
+
"y": { "kind": "tensor", "semantic": "Y", "role": "output" }
|
| 47 |
+
},
|
| 48 |
+
"tunables": { "WORKGROUP_SIZE": 256 },
|
| 49 |
+
"derive": {
|
| 50 |
+
"deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
|
| 51 |
+
"storageBufferLimit": "min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
|
| 52 |
+
"nchw4dContract": "ranks.X == 4 and ranks.Y == 4 and dim(shapes.X, 0) == dim(shapes.Y, 0) and dim(shapes.X, 1) == dim(shapes.Y, 1)",
|
| 53 |
+
"axesContract": "not has(attrs, \"axes\") or ((attrs.axes | length) > 0 and (attrs.axes | length) <= ranks.X and isUniqueIntList(attrs.axes, 0 - ranks.X, ranks.X, ranks.X))",
|
| 54 |
+
"scalesContract": "(attrs.scales | length) == 0 or ((not has(attrs, \"axes\") and (attrs.scales | length) == ranks.X) or (has(attrs, \"axes\") and (attrs.scales | length) == (attrs.axes | length)))",
|
| 55 |
+
"roiContract": "(attrs.roi | length) == 0 or (attrs.roi | length) == 2 * ((attrs.axes | length) if has(attrs, \"axes\") else ranks.X)",
|
| 56 |
+
"noExplicitScales": "(attrs.scales | length) == 0",
|
| 57 |
+
"spatialSpecializedScales": "noExplicitScales or (not has(attrs, \"axes\") and (attrs.scales | length) == ranks.X and ranks.X >= 2 and attrs.scales[0] == 1 and attrs.scales[1] == 1)",
|
| 58 |
+
"workgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
|
| 59 |
+
"storageFits": "numel(shapes.X) * dtypeBytes(dtypes.T) <= storageBufferLimit and numel(shapes.Y) * dtypeBytes(dtypes.T) <= storageBufferLimit",
|
| 60 |
+
"dispatchFits": "numel(shapes.Y) <= device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension * workgroupSize",
|
| 61 |
+
"portableResources": "axesContract and scalesContract and roiContract and workgroupSize > 0 and storageFits and dispatchFits",
|
| 62 |
+
"f16X8PortableOk": "device.features.has(\"subgroups\") or not (has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 8 and device.adapterInfo.subgroupMaxSize <= 32)",
|
| 63 |
+
"rank4Vec4Preferred": "not has(device.adapterInfo, \"subgroupMaxSize\") or device.adapterInfo.subgroupMaxSize <= 32"
|
| 64 |
+
},
|
| 65 |
+
"constants": { "scalar": "dtypes.T" },
|
| 66 |
+
"bindingSets": {
|
| 67 |
+
"static": [
|
| 68 |
+
{ "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
|
| 69 |
+
{ "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputElement" }
|
| 70 |
+
],
|
| 71 |
+
"rank4": [
|
| 72 |
+
{ "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
|
| 73 |
+
{ "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputElement" },
|
| 74 |
+
{
|
| 75 |
+
"name": "params",
|
| 76 |
+
"semantic": "kernel.params",
|
| 77 |
+
"buffer": { "type": "uniform" },
|
| 78 |
+
"struct": {
|
| 79 |
+
"name": "Params",
|
| 80 |
+
"fields": [
|
| 81 |
+
{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
|
| 82 |
+
{ "name": "C", "type": "u32", "value": "dim(shapes.X, 1)" },
|
| 83 |
+
{ "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
|
| 84 |
+
{ "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
|
| 85 |
+
{ "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
|
| 86 |
+
{ "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" },
|
| 87 |
+
{ "name": "scaleH", "type": "f32", "value": "attrs.scales[2] if not noExplicitScales else 0" },
|
| 88 |
+
{ "name": "scaleW", "type": "f32", "value": "attrs.scales[3] if not noExplicitScales else 0" }
|
| 89 |
+
]
|
| 90 |
+
}
|
| 91 |
+
}
|
| 92 |
+
],
|
| 93 |
+
"rank4Cubic": [
|
| 94 |
+
{ "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
|
| 95 |
+
{ "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
|
| 96 |
+
{
|
| 97 |
+
"name": "params",
|
| 98 |
+
"semantic": "kernel.params",
|
| 99 |
+
"buffer": { "type": "uniform" },
|
| 100 |
+
"struct": {
|
| 101 |
+
"name": "Params",
|
| 102 |
+
"fields": [
|
| 103 |
+
{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
|
| 104 |
+
{ "name": "C", "type": "u32", "value": "dim(shapes.X, 1)" },
|
| 105 |
+
{ "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
|
| 106 |
+
{ "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
|
| 107 |
+
{ "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
|
| 108 |
+
{ "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" },
|
| 109 |
+
{ "name": "scaleH", "type": "f32", "value": "attrs.scales[2] if not noExplicitScales else 0" },
|
| 110 |
+
{ "name": "scaleW", "type": "f32", "value": "attrs.scales[3] if not noExplicitScales else 0" },
|
| 111 |
+
{ "name": "coeffA", "type": "f32", "value": "attrs.cubic_coeff_a" }
|
| 112 |
+
]
|
| 113 |
+
}
|
| 114 |
+
}
|
| 115 |
+
],
|
| 116 |
+
"rank5Scalar": [
|
| 117 |
+
{ "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
|
| 118 |
+
{ "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
|
| 119 |
+
{
|
| 120 |
+
"name": "params",
|
| 121 |
+
"semantic": "kernel.params",
|
| 122 |
+
"buffer": { "type": "uniform" },
|
| 123 |
+
"struct": {
|
| 124 |
+
"name": "Params",
|
| 125 |
+
"fields": [
|
| 126 |
+
{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
|
| 127 |
+
{ "name": "C", "type": "u32", "value": "dim(shapes.X, 1)" },
|
| 128 |
+
{ "name": "inD", "type": "u32", "value": "dim(shapes.X, 2)" },
|
| 129 |
+
{ "name": "inH", "type": "u32", "value": "dim(shapes.X, 3)" },
|
| 130 |
+
{ "name": "inW", "type": "u32", "value": "dim(shapes.X, 4)" },
|
| 131 |
+
{ "name": "outD", "type": "u32", "value": "dim(shapes.Y, 2)" },
|
| 132 |
+
{ "name": "outH", "type": "u32", "value": "dim(shapes.Y, 3)" },
|
| 133 |
+
{ "name": "outW", "type": "u32", "value": "dim(shapes.Y, 4)" },
|
| 134 |
+
{ "name": "scaleD", "type": "f32", "value": "attrs.scales[2] if not noExplicitScales else 0" },
|
| 135 |
+
{ "name": "scaleH", "type": "f32", "value": "attrs.scales[3] if not noExplicitScales else 0" },
|
| 136 |
+
{ "name": "scaleW", "type": "f32", "value": "attrs.scales[4] if not noExplicitScales else 0" }
|
| 137 |
+
]
|
| 138 |
+
}
|
| 139 |
+
}
|
| 140 |
+
]
|
| 141 |
+
},
|
| 142 |
+
"variants": [
|
| 143 |
+
{
|
| 144 |
+
"id": "linear_2x_stencil_f16_x8",
|
| 145 |
+
"priority": 22,
|
| 146 |
+
"when": ["nchw4dContract", "numel(shapes.Y) > 0", "attrs.mode == \"linear\"", "(attrs.coordinate_transformation_mode == \"half_pixel\" or attrs.coordinate_transformation_mode == \"pytorch_half_pixel\" or attrs.coordinate_transformation_mode == \"asymmetric\")", "dim(shapes.Y, 2) == dim(shapes.X, 2) * 2", "dim(shapes.Y, 3) == dim(shapes.X, 3) * 2", "dim(shapes.Y, 3) % 8 == 0", "dtypes.T == \"f16\"", "f16X8PortableOk", "f16Ok(dtypes.T)", "portableResources", "noExplicitScales"],
|
| 147 |
+
"constants": { "vectorScalar": "\"vec4<f16>\"", "outputElement": "\"vec4<f16>\"" },
|
| 148 |
+
"passes": [
|
| 149 |
+
{
|
| 150 |
+
"id": "main",
|
| 151 |
+
"name": "Resize.linear2xStencilF16X8",
|
| 152 |
+
"source": {
|
| 153 |
+
"shader": "resize-linear-2x-stencil-x8.wgsl.jinja",
|
| 154 |
+
"inputs": {
|
| 155 |
+
"xShape": "shapes.X",
|
| 156 |
+
"yShape": "shapes.Y",
|
| 157 |
+
"coordinateMode": "attrs.coordinate_transformation_mode"
|
| 158 |
+
}
|
| 159 |
+
},
|
| 160 |
+
"bindings": "static",
|
| 161 |
+
"dispatch": { "gridStride": "numel(shapes.Y) / 8", "workgroupSize": "workgroupSize" }
|
| 162 |
+
}
|
| 163 |
+
]
|
| 164 |
+
},
|
| 165 |
+
{
|
| 166 |
+
"id": "nearest_integer_scale_vec4",
|
| 167 |
+
"priority": 21,
|
| 168 |
+
"when": ["nchw4dContract", "numel(shapes.Y) > 0", "attrs.mode == \"nearest\"", "((attrs.coordinate_transformation_mode == \"asymmetric\" and attrs.nearest_mode == \"floor\") or (attrs.coordinate_transformation_mode == \"half_pixel\" and (attrs.nearest_mode == \"round_prefer_floor\" or attrs.nearest_mode == \"round_prefer_ceil\")))", "(dim(shapes.Y, 2) == dim(shapes.X, 2) * 2 or dim(shapes.Y, 2) == dim(shapes.X, 2) * 3 or dim(shapes.Y, 2) == dim(shapes.X, 2) * 4)", "(dim(shapes.Y, 3) == dim(shapes.X, 3) * 2 or dim(shapes.Y, 3) == dim(shapes.X, 3) * 3 or dim(shapes.Y, 3) == dim(shapes.X, 3) * 4)", "dim(shapes.Y, 3) % 4 == 0", "f16Ok(dtypes.T)", "portableResources", "noExplicitScales"],
|
| 169 |
+
"constants": {
|
| 170 |
+
"usesF16": "dtypes.T == \"f16\"",
|
| 171 |
+
"vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
|
| 172 |
+
"outputElement": "\"vec4<\" ~ dtypes.T ~ \">\""
|
| 173 |
+
},
|
| 174 |
+
"passes": [
|
| 175 |
+
{
|
| 176 |
+
"id": "main",
|
| 177 |
+
"name": "Resize.nearestIntegerScaleVec4",
|
| 178 |
+
"source": {
|
| 179 |
+
"shader": "resize-nearest-integer-scale.wgsl.jinja",
|
| 180 |
+
"inputs": { "xShape": "shapes.X", "yShape": "shapes.Y", "vec4": true }
|
| 181 |
+
},
|
| 182 |
+
"bindings": "static",
|
| 183 |
+
"dispatch": { "gridStride": "numel(shapes.Y) / 4", "workgroupSize": "workgroupSize" }
|
| 184 |
+
}
|
| 185 |
+
]
|
| 186 |
+
},
|
| 187 |
+
{
|
| 188 |
+
"id": "nearest_integer_scale",
|
| 189 |
+
"priority": 20,
|
| 190 |
+
"when": ["nchw4dContract", "numel(shapes.Y) > 0", "attrs.mode == \"nearest\"", "((attrs.coordinate_transformation_mode == \"asymmetric\" and attrs.nearest_mode == \"floor\") or (attrs.coordinate_transformation_mode == \"half_pixel\" and (attrs.nearest_mode == \"round_prefer_floor\" or attrs.nearest_mode == \"round_prefer_ceil\")))", "(dim(shapes.Y, 2) == dim(shapes.X, 2) * 2 or dim(shapes.Y, 2) == dim(shapes.X, 2) * 3 or dim(shapes.Y, 2) == dim(shapes.X, 2) * 4)", "(dim(shapes.Y, 3) == dim(shapes.X, 3) * 2 or dim(shapes.Y, 3) == dim(shapes.X, 3) * 3 or dim(shapes.Y, 3) == dim(shapes.X, 3) * 4)", "f16Ok(dtypes.T)", "portableResources", "noExplicitScales"],
|
| 191 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 192 |
+
"passes": [
|
| 193 |
+
{
|
| 194 |
+
"id": "main",
|
| 195 |
+
"name": "Resize.nearestIntegerScale",
|
| 196 |
+
"source": {
|
| 197 |
+
"shader": "resize-nearest-integer-scale.wgsl.jinja",
|
| 198 |
+
"inputs": { "xShape": "shapes.X", "yShape": "shapes.Y", "vec4": false }
|
| 199 |
+
},
|
| 200 |
+
"bindings": "static",
|
| 201 |
+
"dispatch": { "gridStride": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 202 |
+
}
|
| 203 |
+
]
|
| 204 |
+
},
|
| 205 |
+
{
|
| 206 |
+
"id": "linear_2x_stencil_vec4",
|
| 207 |
+
"priority": 20,
|
| 208 |
+
"when": ["nchw4dContract", "numel(shapes.Y) > 0", "attrs.mode == \"linear\"", "(attrs.coordinate_transformation_mode == \"half_pixel\" or attrs.coordinate_transformation_mode == \"pytorch_half_pixel\" or attrs.coordinate_transformation_mode == \"asymmetric\")", "dim(shapes.Y, 2) == dim(shapes.X, 2) * 2", "dim(shapes.Y, 3) == dim(shapes.X, 3) * 2", "dim(shapes.Y, 3) % 4 == 0", "f16Ok(dtypes.T)", "portableResources", "noExplicitScales"],
|
| 209 |
+
"constants": {
|
| 210 |
+
"usesF16": "dtypes.T == \"f16\"",
|
| 211 |
+
"vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
|
| 212 |
+
"outputElement": "\"vec4<\" ~ dtypes.T ~ \">\""
|
| 213 |
+
},
|
| 214 |
+
"passes": [
|
| 215 |
+
{
|
| 216 |
+
"id": "main",
|
| 217 |
+
"name": "Resize.linear2xStencilVec4",
|
| 218 |
+
"source": {
|
| 219 |
+
"shader": "resize-linear-2x-stencil.wgsl.jinja",
|
| 220 |
+
"inputs": {
|
| 221 |
+
"xShape": "shapes.X",
|
| 222 |
+
"yShape": "shapes.Y",
|
| 223 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 224 |
+
"vec4": true
|
| 225 |
+
}
|
| 226 |
+
},
|
| 227 |
+
"bindings": "static",
|
| 228 |
+
"dispatch": { "gridStride": "numel(shapes.Y) / 4", "workgroupSize": "workgroupSize" }
|
| 229 |
+
}
|
| 230 |
+
]
|
| 231 |
+
},
|
| 232 |
+
{
|
| 233 |
+
"id": "linear_2x_stencil",
|
| 234 |
+
"priority": 19,
|
| 235 |
+
"when": ["nchw4dContract", "numel(shapes.Y) > 0", "attrs.mode == \"linear\"", "(attrs.coordinate_transformation_mode == \"half_pixel\" or attrs.coordinate_transformation_mode == \"pytorch_half_pixel\" or attrs.coordinate_transformation_mode == \"asymmetric\")", "dim(shapes.Y, 2) == dim(shapes.X, 2) * 2", "dim(shapes.Y, 3) == dim(shapes.X, 3) * 2", "f16Ok(dtypes.T)", "portableResources", "noExplicitScales"],
|
| 236 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 237 |
+
"passes": [
|
| 238 |
+
{
|
| 239 |
+
"id": "main",
|
| 240 |
+
"name": "Resize.linear2xStencil",
|
| 241 |
+
"source": {
|
| 242 |
+
"shader": "resize-linear-2x-stencil.wgsl.jinja",
|
| 243 |
+
"inputs": {
|
| 244 |
+
"xShape": "shapes.X",
|
| 245 |
+
"yShape": "shapes.Y",
|
| 246 |
+
"coordinateMode": "attrs.coordinate_transformation_mode"
|
| 247 |
+
}
|
| 248 |
+
},
|
| 249 |
+
"bindings": "static",
|
| 250 |
+
"dispatch": { "gridStride": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 251 |
+
}
|
| 252 |
+
]
|
| 253 |
+
},
|
| 254 |
+
{
|
| 255 |
+
"id": "nchw_rank4_vec4",
|
| 256 |
+
"priority": 1,
|
| 257 |
+
"when": ["nchw4dContract", "(attrs.mode == \"nearest\" or attrs.mode == \"linear\")", "attrs.coordinate_transformation_mode != \"tf_crop_and_resize\"", "(attrs.antialias != 1 or attrs.mode == \"nearest\")", "dim(shapes.Y, 3) % 4 == 0", "f16Ok(dtypes.T)", "portableResources", "spatialSpecializedScales"],
|
| 258 |
+
"demoteWhen": ["not rank4Vec4Preferred"],
|
| 259 |
+
"constants": {
|
| 260 |
+
"usesF16": "dtypes.T == \"f16\"",
|
| 261 |
+
"vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
|
| 262 |
+
"outputElement": "\"vec4<\" ~ dtypes.T ~ \">\""
|
| 263 |
+
},
|
| 264 |
+
"passes": [
|
| 265 |
+
{
|
| 266 |
+
"id": "main",
|
| 267 |
+
"name": "Resize.Vec4",
|
| 268 |
+
"source": {
|
| 269 |
+
"shader": "resize-coord-transform.wgsl.jinja",
|
| 270 |
+
"inputs": {
|
| 271 |
+
"mode": "attrs.mode",
|
| 272 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 273 |
+
"nearestMode": "attrs.nearest_mode",
|
| 274 |
+
"vec4": true,
|
| 275 |
+
"vecType": "\"vec4<\" ~ dtypes.T ~ \">\""
|
| 276 |
+
}
|
| 277 |
+
},
|
| 278 |
+
"bindings": "rank4",
|
| 279 |
+
"dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "workgroupSize" }
|
| 280 |
+
}
|
| 281 |
+
]
|
| 282 |
+
},
|
| 283 |
+
{
|
| 284 |
+
"id": "nchw_rank4_scalar_x4",
|
| 285 |
+
"priority": 2,
|
| 286 |
+
"when": ["nchw4dContract", "(attrs.mode == \"nearest\" or attrs.mode == \"linear\")", "attrs.coordinate_transformation_mode != \"tf_crop_and_resize\"", "(attrs.antialias != 1 or attrs.mode == \"nearest\")", "dim(shapes.Y, 3) >= 4", "dim(shapes.Y, 3) % 4 != 0", "f16Ok(dtypes.T)", "portableResources", "spatialSpecializedScales"],
|
| 287 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 288 |
+
"passes": [
|
| 289 |
+
{
|
| 290 |
+
"id": "main",
|
| 291 |
+
"name": "Resize.ScalarX4",
|
| 292 |
+
"source": {
|
| 293 |
+
"shader": "resize-coord-transform.wgsl.jinja",
|
| 294 |
+
"inputs": {
|
| 295 |
+
"mode": "attrs.mode",
|
| 296 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 297 |
+
"nearestMode": "attrs.nearest_mode",
|
| 298 |
+
"scalarX4": true
|
| 299 |
+
}
|
| 300 |
+
},
|
| 301 |
+
"bindings": "rank4",
|
| 302 |
+
"dispatch": {
|
| 303 |
+
"threads": "dim(shapes.Y, 0) * dim(shapes.Y, 1) * dim(shapes.Y, 2) * ceil(dim(shapes.Y, 3) / 4)",
|
| 304 |
+
"workgroupSize": "workgroupSize"
|
| 305 |
+
}
|
| 306 |
+
}
|
| 307 |
+
]
|
| 308 |
+
},
|
| 309 |
+
{
|
| 310 |
+
"id": "nchw_rank4",
|
| 311 |
+
"when": ["nchw4dContract", "(attrs.mode == \"nearest\" or attrs.mode == \"linear\")", "attrs.coordinate_transformation_mode != \"tf_crop_and_resize\"", "(attrs.antialias != 1 or attrs.mode == \"nearest\")", "f16Ok(dtypes.T)", "portableResources", "spatialSpecializedScales"],
|
| 312 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 313 |
+
"passes": [
|
| 314 |
+
{
|
| 315 |
+
"id": "main",
|
| 316 |
+
"name": "Resize",
|
| 317 |
+
"source": {
|
| 318 |
+
"shader": "resize-coord-transform.wgsl.jinja",
|
| 319 |
+
"inputs": {
|
| 320 |
+
"mode": "attrs.mode",
|
| 321 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 322 |
+
"nearestMode": "attrs.nearest_mode"
|
| 323 |
+
}
|
| 324 |
+
},
|
| 325 |
+
"bindings": "rank4",
|
| 326 |
+
"dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 327 |
+
}
|
| 328 |
+
]
|
| 329 |
+
},
|
| 330 |
+
{
|
| 331 |
+
"id": "ncdhw_rank5",
|
| 332 |
+
"when": ["ranks.X == 5", "ranks.Y == 5", "dim(shapes.X, 0) == dim(shapes.Y, 0)", "dim(shapes.X, 1) == dim(shapes.Y, 1)", "(attrs.mode == \"nearest\" or attrs.mode == \"linear\")", "attrs.coordinate_transformation_mode != \"tf_crop_and_resize\"", "(attrs.antialias != 1 or attrs.mode == \"nearest\")", "f16Ok(dtypes.T)", "portableResources", "spatialSpecializedScales"],
|
| 333 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"" },
|
| 334 |
+
"passes": [
|
| 335 |
+
{
|
| 336 |
+
"id": "main",
|
| 337 |
+
"name": "Resize.Ncdhw",
|
| 338 |
+
"source": {
|
| 339 |
+
"shader": "resize-coord-transform-5d.wgsl.jinja",
|
| 340 |
+
"inputs": {
|
| 341 |
+
"mode": "attrs.mode",
|
| 342 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 343 |
+
"nearestMode": "attrs.nearest_mode"
|
| 344 |
+
}
|
| 345 |
+
},
|
| 346 |
+
"bindings": "rank5Scalar",
|
| 347 |
+
"dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 348 |
+
}
|
| 349 |
+
]
|
| 350 |
+
},
|
| 351 |
+
{
|
| 352 |
+
"id": "nchw_rank4_cubic",
|
| 353 |
+
"when": ["nchw4dContract", "(attrs.antialias != 1)", "attrs.mode == \"cubic\"", "(attrs.coordinate_transformation_mode == \"half_pixel\" or attrs.coordinate_transformation_mode == \"pytorch_half_pixel\" or attrs.coordinate_transformation_mode == \"asymmetric\" or attrs.coordinate_transformation_mode == \"align_corners\")", "f16Ok(dtypes.T)", "portableResources", "spatialSpecializedScales"],
|
| 354 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"" },
|
| 355 |
+
"passes": [
|
| 356 |
+
{
|
| 357 |
+
"id": "main",
|
| 358 |
+
"name": "Resize.Cubic",
|
| 359 |
+
"source": {
|
| 360 |
+
"shader": "resize-cubic.wgsl.jinja",
|
| 361 |
+
"inputs": {
|
| 362 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 363 |
+
"excludeOutside": "attrs.exclude_outside if attrs.exclude_outside else 0"
|
| 364 |
+
}
|
| 365 |
+
},
|
| 366 |
+
"bindings": "rank4Cubic",
|
| 367 |
+
"dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 368 |
+
}
|
| 369 |
+
]
|
| 370 |
+
},
|
| 371 |
+
{
|
| 372 |
+
"id": "generic_rank_n",
|
| 373 |
+
"priority": 0,
|
| 374 |
+
"when": ["ranks.X == ranks.Y", "ranks.X >= 1", "numel(shapes.Y) > 0", "(attrs.mode == \"nearest\" or attrs.mode == \"linear\")", "(attrs.antialias != 1 or attrs.mode == \"nearest\")", "(attrs.coordinate_transformation_mode == \"tf_crop_and_resize\" or ranks.X < 4 or ranks.X > 5 or dim(shapes.X, 0) != dim(shapes.Y, 0) or dim(shapes.X, 1) != dim(shapes.Y, 1) or not spatialSpecializedScales)", "f16Ok(dtypes.T)", "portableResources"],
|
| 375 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 376 |
+
"passes": [
|
| 377 |
+
{
|
| 378 |
+
"id": "main",
|
| 379 |
+
"name": "Resize.Generic",
|
| 380 |
+
"source": {
|
| 381 |
+
"shader": "resize-generic.wgsl.jinja",
|
| 382 |
+
"inputs": {
|
| 383 |
+
"xShape": "shapes.X",
|
| 384 |
+
"yShape": "shapes.Y",
|
| 385 |
+
"mode": "attrs.mode",
|
| 386 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 387 |
+
"nearestMode": "attrs.nearest_mode",
|
| 388 |
+
"roi": "attrs.roi",
|
| 389 |
+
"scales": "attrs.scales",
|
| 390 |
+
"extrapolationValue": "attrs.extrapolation_value",
|
| 391 |
+
"excludeOutside": "attrs.exclude_outside if attrs.exclude_outside else 0",
|
| 392 |
+
"axes": "attrs.axes if has(attrs, \"axes\") else []"
|
| 393 |
+
}
|
| 394 |
+
},
|
| 395 |
+
"bindings": "static",
|
| 396 |
+
"dispatch": { "gridStride": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 397 |
+
}
|
| 398 |
+
]
|
| 399 |
+
},
|
| 400 |
+
{
|
| 401 |
+
"id": "antialias_linear",
|
| 402 |
+
"priority": 5,
|
| 403 |
+
"when": ["ranks.X == ranks.Y", "ranks.X >= 2", "numel(shapes.Y) > 0", "attrs.mode == \"linear\"", "attrs.antialias == 1", "(attrs.coordinate_transformation_mode == \"half_pixel\" or attrs.coordinate_transformation_mode == \"pytorch_half_pixel\" or attrs.coordinate_transformation_mode == \"asymmetric\" or attrs.coordinate_transformation_mode == \"align_corners\")", "f16Ok(dtypes.T)", "portableResources"],
|
| 404 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 405 |
+
"passes": [
|
| 406 |
+
{
|
| 407 |
+
"id": "main",
|
| 408 |
+
"name": "Resize.AntialiasLinear",
|
| 409 |
+
"source": {
|
| 410 |
+
"shader": "resize-antialias.wgsl.jinja",
|
| 411 |
+
"inputs": {
|
| 412 |
+
"xShape": "shapes.X",
|
| 413 |
+
"yShape": "shapes.Y",
|
| 414 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 415 |
+
"excludeOutside": "attrs.exclude_outside == 1",
|
| 416 |
+
"scales": "attrs.scales",
|
| 417 |
+
"axes": "attrs.axes if has(attrs, \"axes\") else []"
|
| 418 |
+
}
|
| 419 |
+
},
|
| 420 |
+
"bindings": "static",
|
| 421 |
+
"dispatch": { "gridStride": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 422 |
+
}
|
| 423 |
+
]
|
| 424 |
+
},
|
| 425 |
+
{
|
| 426 |
+
"id": "antialias_cubic",
|
| 427 |
+
"priority": 5,
|
| 428 |
+
"when": ["ranks.X == ranks.Y", "ranks.X >= 2", "numel(shapes.Y) > 0", "attrs.mode == \"cubic\"", "attrs.antialias == 1", "(attrs.coordinate_transformation_mode == \"half_pixel\" or attrs.coordinate_transformation_mode == \"pytorch_half_pixel\" or attrs.coordinate_transformation_mode == \"asymmetric\" or attrs.coordinate_transformation_mode == \"align_corners\")", "f16Ok(dtypes.T)", "portableResources"],
|
| 429 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 430 |
+
"passes": [
|
| 431 |
+
{
|
| 432 |
+
"id": "main",
|
| 433 |
+
"name": "Resize.AntialiasCubic",
|
| 434 |
+
"source": {
|
| 435 |
+
"shader": "resize-antialias.wgsl.jinja",
|
| 436 |
+
"inputs": {
|
| 437 |
+
"xShape": "shapes.X",
|
| 438 |
+
"yShape": "shapes.Y",
|
| 439 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 440 |
+
"excludeOutside": "attrs.exclude_outside == 1",
|
| 441 |
+
"scales": "attrs.scales",
|
| 442 |
+
"filter": "\"cubic\"",
|
| 443 |
+
"coeffA": "attrs.cubic_coeff_a",
|
| 444 |
+
"axes": "attrs.axes if has(attrs, \"axes\") else []"
|
| 445 |
+
}
|
| 446 |
+
},
|
| 447 |
+
"bindings": "static",
|
| 448 |
+
"dispatch": { "gridStride": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 449 |
+
}
|
| 450 |
+
]
|
| 451 |
+
},
|
| 452 |
+
{
|
| 453 |
+
"id": "generic_rank_n_cubic",
|
| 454 |
+
"priority": 0,
|
| 455 |
+
"when": ["ranks.X == ranks.Y", "ranks.X >= 1", "numel(shapes.Y) > 0", "(attrs.antialias != 1)", "attrs.mode == \"cubic\"", "(attrs.coordinate_transformation_mode == \"half_pixel\" or attrs.coordinate_transformation_mode == \"pytorch_half_pixel\" or attrs.coordinate_transformation_mode == \"asymmetric\" or attrs.coordinate_transformation_mode == \"align_corners\")", "(ranks.X != 4 or dim(shapes.X, 0) != dim(shapes.Y, 0) or dim(shapes.X, 1) != dim(shapes.Y, 1) or not spatialSpecializedScales)", "f16Ok(dtypes.T)", "portableResources"],
|
| 456 |
+
"constants": { "usesF16": "dtypes.T == \"f16\"", "outputElement": "dtypes.T" },
|
| 457 |
+
"passes": [
|
| 458 |
+
{
|
| 459 |
+
"id": "main",
|
| 460 |
+
"name": "Resize.GenericCubic",
|
| 461 |
+
"source": {
|
| 462 |
+
"shader": "resize-generic.wgsl.jinja",
|
| 463 |
+
"inputs": {
|
| 464 |
+
"xShape": "shapes.X",
|
| 465 |
+
"yShape": "shapes.Y",
|
| 466 |
+
"mode": "attrs.mode",
|
| 467 |
+
"coordinateMode": "attrs.coordinate_transformation_mode",
|
| 468 |
+
"nearestMode": "attrs.nearest_mode",
|
| 469 |
+
"coeffA": "attrs.cubic_coeff_a",
|
| 470 |
+
"roi": "attrs.roi",
|
| 471 |
+
"scales": "attrs.scales",
|
| 472 |
+
"extrapolationValue": "attrs.extrapolation_value",
|
| 473 |
+
"excludeOutside": "attrs.exclude_outside if attrs.exclude_outside else 0",
|
| 474 |
+
"axes": "attrs.axes if has(attrs, \"axes\") else []"
|
| 475 |
+
}
|
| 476 |
+
},
|
| 477 |
+
"bindings": "static",
|
| 478 |
+
"dispatch": { "gridStride": "numel(shapes.Y)", "workgroupSize": "workgroupSize" }
|
| 479 |
+
}
|
| 480 |
+
]
|
| 481 |
+
}
|
| 482 |
+
]
|
| 483 |
+
}
|
build/webgpu/metadata.json
ADDED
|
@@ -0,0 +1,25 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
{
|
| 2 |
+
"name": "ai.onnx.Resize",
|
| 3 |
+
"id": "_ai_onnx_resize_webgpu_0155293",
|
| 4 |
+
"version": 1,
|
| 5 |
+
"license": "Apache-2.0",
|
| 6 |
+
"backend": { "type": "webgpu" },
|
| 7 |
+
"digest": {
|
| 8 |
+
"algorithm": "sha256",
|
| 9 |
+
"files": {
|
| 10 |
+
"bench.json": "gBoDCJxNV/CQ0sFt2VIweuZzVQ5ZE8KiGZxw1hRRmWI=",
|
| 11 |
+
"manifest.json": "kCg0ZXiAdK80uaoitypMEHWn5/vUFkKlUSBqmaKSVGE=",
|
| 12 |
+
"resize-antialias.wgsl.jinja": "Ku4/azULjkzIO0lW2a5sh2LEAt/XGD23dEc/bKcnYKU=",
|
| 13 |
+
"resize-coord-transform-5d.wgsl.jinja": "gq3U+clM7SPPu+VzMf8uIlHm2trJED/SiOABXtgRuKI=",
|
| 14 |
+
"resize-coord-transform.wgsl.jinja": "ekD/ojBZXcHosWDo639tp/Vsi5SemF7OE8FWYNdIvMk=",
|
| 15 |
+
"resize-cubic.wgsl.jinja": "0Z/lgy1+yxfW5GYhV4HuC2biY4nqsTgToHjz/aHUBPQ=",
|
| 16 |
+
"resize-generic.wgsl.jinja": "uSTH4SpnJnVYFc0nRbNiB5fCyUvYi7u10bs+SZw5brw=",
|
| 17 |
+
"resize-linear-2x-stencil-x8.wgsl.jinja": "JkfXP+3P19pHLuYmn/vIF1ccfjjvoRdP/AoMDZWqZWM=",
|
| 18 |
+
"resize-linear-2x-stencil.wgsl.jinja": "/gGRgsSwzZ7qGar2xwj8fPpDyG07eTwLdHJGPIOpu14=",
|
| 19 |
+
"resize-nearest-integer-scale.wgsl.jinja": "fgry1tCB+VAmoiWgqHWsy3EL2G47mtvtZhVPP27Z0Fg=",
|
| 20 |
+
"test.json": "YXNNuR7ZMl9fAc3N2D1p2IeZpN3I4hpffoHROgKbOEQ="
|
| 21 |
+
}
|
| 22 |
+
},
|
| 23 |
+
"provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
|
| 24 |
+
"webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Resize" }
|
| 25 |
+
}
|
build/webgpu/resize-antialias.wgsl.jinja
ADDED
|
@@ -0,0 +1,159 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// Downsample-aware antialiased Resize. A separable filter runs over each resized
|
| 2 |
+
// axis. The filter scale is max(1, downsample ratio), so downsampling widens the
|
| 3 |
+
// support and averages a source window. Linear mode uses a triangle/area filter;
|
| 4 |
+
// cubic mode uses a Keys kernel at distance/filter-scale. Excluding outside
|
| 5 |
+
// samples zeroes out-of-bounds taps before per-window normalization.
|
| 6 |
+
//
|
| 7 |
+
// An explicit per-axis scale drives both coordinate transformation and filter
|
| 8 |
+
// scale. Because out = floor(in*scale), it may differ from the output/input
|
| 9 |
+
// shape ratio. The window start is dynamic, while its maximum size is baked.
|
| 10 |
+
{% if usesF16 %}
|
| 11 |
+
enable f16;
|
| 12 |
+
{% endif %}
|
| 13 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 14 |
+
{% set wg = workgroupSize %}
|
| 15 |
+
{%- set xs = source.xShape %}
|
| 16 |
+
{% set ys = source.yShape %}
|
| 17 |
+
{% set r = xs | length %}
|
| 18 |
+
{% set FILTER = source.filter if source.filter is defined else "linear" %}
|
| 19 |
+
{% set supmul = 2 if FILTER == "cubic" else 1 %}
|
| 20 |
+
{% set resolvedScales = namespace(values = []) %}
|
| 21 |
+
{% for a in range(r) %}
|
| 22 |
+
{% set axisScale = namespace(value = 0) %}
|
| 23 |
+
{% if source.scales is defined and source.scales %}
|
| 24 |
+
{% if source.axes is defined and source.axes %}
|
| 25 |
+
{% for j in range(source.axes | length) %}
|
| 26 |
+
{% set scaleAxis = source.axes[j] if source.axes[j] >= 0 else source.axes[j] + r %}
|
| 27 |
+
{% if scaleAxis == a %}{% set axisScale.value = source.scales[j] %}{% endif %}
|
| 28 |
+
{% endfor %}
|
| 29 |
+
{% else %}
|
| 30 |
+
{% set axisScale.value = source.scales[a] %}
|
| 31 |
+
{% endif %}
|
| 32 |
+
{% endif %}
|
| 33 |
+
{% set resolvedScales.values = resolvedScales.values + [axisScale.value] %}
|
| 34 |
+
{% endfor %}
|
| 35 |
+
{% set sc = resolvedScales.values %}
|
| 36 |
+
|
| 37 |
+
{% set ocount = namespace(v = 1) %}
|
| 38 |
+
{% for a in range(r) %}{% set ocount.v = ocount.v * ys[a] %}{% endfor %}
|
| 39 |
+
|
| 40 |
+
{% set resized = namespace(list = []) %}
|
| 41 |
+
{% for a in range(r) %}
|
| 42 |
+
{% if xs[a] != ys[a] or (sc[a] > 0 and sc[a] != 1) %}{% set resized.list = resized.list + [a] %}{% endif %}
|
| 43 |
+
{% endfor %}
|
| 44 |
+
{% set k = resized.list | length %}
|
| 45 |
+
|
| 46 |
+
const WG: u32 = {{ wg }}u;
|
| 47 |
+
{% if FILTER == "cubic" %}
|
| 48 |
+
|
| 49 |
+
// Keys cubic convolution kernel with coefficient coeffA.
|
| 50 |
+
fn cubic_weight(s_in: f32) -> f32 {
|
| 51 |
+
let a = f32({{ source.coeffA }});
|
| 52 |
+
let s = abs(s_in);
|
| 53 |
+
let s2 = s * s;
|
| 54 |
+
let s3 = s2 * s;
|
| 55 |
+
if (s <= 1.0) { return (a + 2.0) * s3 - (a + 3.0) * s2 + 1.0; }
|
| 56 |
+
if (s < 2.0) { return a * s3 - 5.0 * a * s2 + 8.0 * a * s - 4.0 * a; }
|
| 57 |
+
return 0.0;
|
| 58 |
+
}
|
| 59 |
+
|
| 60 |
+
{% endif %}
|
| 61 |
+
{% macro istride(a) %}
|
| 62 |
+
{% set s = namespace(v = 1) %}
|
| 63 |
+
{% for j in range(a + 1, r) %}{% set s.v = s.v * xs[j] %}{% endfor %}
|
| 64 |
+
{{ s.v }}
|
| 65 |
+
{%- endmacro %}{% macro center(a, oc) %}
|
| 66 |
+
{% set inS = xs[a] %}
|
| 67 |
+
{% set outS = ys[a] %}
|
| 68 |
+
{% set axsc = sc[a] if ((sc | length) > a and sc[a] and sc[a] > 0) else 0 %}
|
| 69 |
+
{% set inv = (1.0 / axsc) if axsc > 0 else (inS / outS) %}
|
| 70 |
+
{% if source.coordinateMode == "align_corners" %}
|
| 71 |
+
{% if outS == 1 %}0.0{% else %}(f32({{ oc }}) * {{ inS - 1 }}.0 / {{ outS - 1 }}.0){% endif %}
|
| 72 |
+
{% elif source.coordinateMode == "asymmetric" -%}
|
| 73 |
+
(f32({{ oc }}) * {{ inv }})
|
| 74 |
+
{%- elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 75 |
+
{% if outS > 1 %}((f32({{ oc }}) + 0.5) * {{ inv }} - 0.5){% else %}0.0{% endif %}
|
| 76 |
+
{% else -%}
|
| 77 |
+
((f32({{ oc }}) + 0.5) * {{ inv }} - 0.5)
|
| 78 |
+
{%- endif %}
|
| 79 |
+
{% endmacro %}
|
| 80 |
+
|
| 81 |
+
@compute @workgroup_size({{ wg }})
|
| 82 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
|
| 83 |
+
let stride = nwg.x * WG;
|
| 84 |
+
for (var i = gid.x; i < {{ ocount.v }}u; i += stride) {
|
| 85 |
+
var rem = i;
|
| 86 |
+
{% for a in range(r) %}
|
| 87 |
+
{% set os = namespace(v = 1) %}
|
| 88 |
+
{% for j in range(a + 1, r) %}{% set os.v = os.v * ys[j] %}{% endfor %}
|
| 89 |
+
{% set constantZeroCenter = a in resized.list and ys[a] == 1 and source.coordinateMode in ["align_corners", "pytorch_half_pixel"] %}
|
| 90 |
+
{% if not constantZeroCenter %}
|
| 91 |
+
let oc{{ a }} = rem / {{ os.v }}u;
|
| 92 |
+
rem = rem % {{ os.v }}u;
|
| 93 |
+
{% endif %}
|
| 94 |
+
{% endfor %}
|
| 95 |
+
{% for ridx in range(k) %}
|
| 96 |
+
{% set a = resized.list[ridx] %}
|
| 97 |
+
{% set inS = xs[a] %}
|
| 98 |
+
{% set outS = ys[a] %}
|
| 99 |
+
{% set axsc = sc[a] if ((sc | length) > a and sc[a] and sc[a] > 0) else 0 %}
|
| 100 |
+
{% set dsr = (1.0 / axsc) if axsc > 0 else (inS / outS) %}
|
| 101 |
+
{% set fscale = dsr if dsr > 1 else 1.0 %}
|
| 102 |
+
{% set support = fscale * supmul %}
|
| 103 |
+
{% set win = (2 * ceil(support)) + 1 %}
|
| 104 |
+
// axis {{ a }}: filterscale={{ fscale }}, support={{ support }}, window<={{ win }}
|
| 105 |
+
let center{{ a }} = {{ center(a, "oc" ~ a) }};
|
| 106 |
+
let start{{ a }} = i32(floor(center{{ a }} - {{ support }} + 0.5));
|
| 107 |
+
var wsum{{ a }} = 0.0;
|
| 108 |
+
var w{{ a }}: array<f32, {{ win }}>;
|
| 109 |
+
{% for t in range(win) %}
|
| 110 |
+
{
|
| 111 |
+
let xx = start{{ a }} + {{ t }};
|
| 112 |
+
{% if FILTER == "cubic" %}
|
| 113 |
+
var ww = cubic_weight((f32(xx) - center{{ a }}) / {{ fscale }});
|
| 114 |
+
{% else %}
|
| 115 |
+
var ww = 1.0 - abs((f32(xx) - center{{ a }}) / {{ fscale }});
|
| 116 |
+
ww = max(0.0, ww);
|
| 117 |
+
{% endif %}
|
| 118 |
+
{% if source.excludeOutside %}
|
| 119 |
+
if (xx < 0 || xx > {{ inS - 1 }}) { ww = 0.0; }
|
| 120 |
+
{% endif %}
|
| 121 |
+
w{{ a }}[{{ t }}] = ww;
|
| 122 |
+
wsum{{ a }} = wsum{{ a }} + ww;
|
| 123 |
+
}
|
| 124 |
+
{% endfor %}
|
| 125 |
+
{% endfor %}
|
| 126 |
+
var acc = 0.0;{% set total = namespace(v = 1) %}
|
| 127 |
+
{% for ridx in range(k) %}
|
| 128 |
+
{% set a = resized.list[ridx] %}
|
| 129 |
+
{% set inS = xs[a] %}{% set outS = ys[a] %}
|
| 130 |
+
{% set axsc = sc[a] if ((sc | length) > a and sc[a] and sc[a] > 0) else 0 %}
|
| 131 |
+
{% set dsr = (1.0 / axsc) if axsc > 0 else (inS / outS) %}
|
| 132 |
+
{% set fscale = dsr if dsr > 1 else 1.0 %}
|
| 133 |
+
{% set win = (2 * ceil(fscale * supmul)) + 1 %}
|
| 134 |
+
{% set total.v = total.v * win %}
|
| 135 |
+
{% endfor %}
|
| 136 |
+
{% for combo in range(total.v) %}
|
| 137 |
+
{
|
| 138 |
+
var cw = 1.0;
|
| 139 |
+
{% set divns = namespace(d = 1) %}
|
| 140 |
+
{% for ridx in range(k) %}
|
| 141 |
+
{% set a = resized.list[ridx] %}
|
| 142 |
+
{% set inS = xs[a] %}{% set outS = ys[a] %}
|
| 143 |
+
{% set axsc = sc[a] if ((sc | length) > a and sc[a] and sc[a] > 0) else 0 %}
|
| 144 |
+
{% set dsr = (1.0 / axsc) if axsc > 0 else (inS / outS) %}
|
| 145 |
+
{% set fscale = dsr if dsr > 1 else 1.0 %}
|
| 146 |
+
{% set win = (2 * ceil(fscale * supmul)) + 1 %}
|
| 147 |
+
{% set t = ((combo / divns.d) | int) % win %}
|
| 148 |
+
{% set divns.d = divns.d * win %}
|
| 149 |
+
let ci{{ a }} = u32(clamp(start{{ a }} + {{ t }}, 0, {{ inS - 1 }}));
|
| 150 |
+
cw = cw * w{{ a }}[{{ t }}];
|
| 151 |
+
{% endfor %}
|
| 152 |
+
let off = 0u{% for a in range(r) %}{% if a in resized.list %} + ci{{ a }} * {{ istride(a) }}u{% else %} + oc{{ a }} * {{ istride(a) }}u{% endif %}{% endfor %};
|
| 153 |
+
acc = acc + cw * f32(x[off]);
|
| 154 |
+
}
|
| 155 |
+
{% endfor %}
|
| 156 |
+
let norm = 1.0{% for ridx in range(k) %}{% set a = resized.list[ridx] %} * wsum{{ a }}{% endfor %};
|
| 157 |
+
y[i] = {{ scalar }}(acc / norm);
|
| 158 |
+
}
|
| 159 |
+
}
|
build/webgpu/resize-coord-transform-5d.wgsl.jinja
ADDED
|
@@ -0,0 +1,224 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// Generic NCDHW (rank-5) Resize with full coordinate-transformation math.
|
| 2 |
+
// The batch N and channel C axes are unchanged; the three trailing axes D/H/W
|
| 3 |
+
// are resized. This mirrors the rank-specialized coordinate transform and
|
| 4 |
+
// clamping, so half_pixel, pytorch_half_pixel, asymmetric, and align_corners
|
| 5 |
+
// stay bit-for-bit consistent with the rank-4 path.
|
| 6 |
+
{% set alignCorners = source.coordinateMode == "align_corners" %}
|
| 7 |
+
{% if usesF16 %}
|
| 8 |
+
enable f16;
|
| 9 |
+
{% endif %}
|
| 10 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 11 |
+
{% set wg = workgroupSize %}
|
| 12 |
+
|
| 13 |
+
{% if source.mode != "nearest" or not alignCorners %}
|
| 14 |
+
fn source_coord(out_coord: u32, in_size: u32, out_size: u32{% if not alignCorners %}, explicit_scale: f32{% endif %}) -> f32 {
|
| 15 |
+
{% if not alignCorners %}
|
| 16 |
+
let inferred_scale = f32(out_size) / f32(in_size);
|
| 17 |
+
let scale = select(inferred_scale, explicit_scale, explicit_scale > 0.0);
|
| 18 |
+
{% endif %}
|
| 19 |
+
{% if alignCorners %}
|
| 20 |
+
if (out_size == 1u) { return 0.0; }
|
| 21 |
+
return f32(out_coord) * f32(in_size - 1u) / f32(out_size - 1u);
|
| 22 |
+
{% elif source.coordinateMode == "asymmetric" %}
|
| 23 |
+
return f32(out_coord) / scale;
|
| 24 |
+
{% elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 25 |
+
return select(0.0, (f32(out_coord) + 0.5) / scale - 0.5, out_size > 1u);
|
| 26 |
+
{% elif source.coordinateMode == "half_pixel_symmetric" %}
|
| 27 |
+
let output_width = f32(in_size) * scale;
|
| 28 |
+
let adjustment = f32(out_size) / output_width;
|
| 29 |
+
let center = f32(in_size) * 0.5;
|
| 30 |
+
let offset = center * (1.0 - adjustment);
|
| 31 |
+
return offset + (f32(out_coord) + 0.5) / scale - 0.5;
|
| 32 |
+
{% else %}
|
| 33 |
+
return (f32(out_coord) + 0.5) / scale - 0.5;
|
| 34 |
+
{% endif %}
|
| 35 |
+
}
|
| 36 |
+
{% endif %}
|
| 37 |
+
|
| 38 |
+
/* Backend-stable nearest-neighbor coordinate selection. Inferred align-corners,
|
| 39 |
+
* asymmetric, and half-pixel transforms use exact integer ratios. This matters
|
| 40 |
+
* at nearest-mode ties: reassociating an f32 division into a reciprocal
|
| 41 |
+
* multiply can move an exact 0.5 to either side of the boundary. */{% set ratioMode = source.coordinateMode in ["align_corners", "asymmetric"] %}
|
| 42 |
+
{% set inferredHalfPixelMode = source.coordinateMode in ["pytorch_half_pixel", "half_pixel", "half_pixel_symmetric"] %}
|
| 43 |
+
{% set noRoi = source.roi is not defined or (source.roi | length) == 0 %}
|
| 44 |
+
{% set usesExplicitScale = source.coordinateMode not in ["align_corners", "tf_crop_and_resize"] %}
|
| 45 |
+
{% if source.mode == "nearest" %}
|
| 46 |
+
fn nearest_index(v: f32) -> i32 {
|
| 47 |
+
{% if source.nearestMode == "floor" %}
|
| 48 |
+
return i32(floor(v));
|
| 49 |
+
{% elif source.nearestMode == "ceil" %}
|
| 50 |
+
return i32(ceil(v));
|
| 51 |
+
{% elif source.nearestMode == "round_prefer_ceil" %}
|
| 52 |
+
// round half toward +inf. WGSL round() is round-half-to-even and therefore
|
| 53 |
+
// cannot express ONNX's preferred direction. The lower-side epsilon absorbs
|
| 54 |
+
// backend division error that lands a mathematical half-tie just below .5
|
| 55 |
+
// (mirror of the round_prefer_floor epsilon below).
|
| 56 |
+
let tie_epsilon = 0.0000003;
|
| 57 |
+
let f = floor(v);
|
| 58 |
+
let frac = v - f;
|
| 59 |
+
return i32(select(f, f + 1.0, frac >= 0.5 - tie_epsilon));
|
| 60 |
+
{% else %}
|
| 61 |
+
// Explicit scales still require f32 coordinate math. Keep a narrow epsilon
|
| 62 |
+
// for their backend rounding while inferred scales take the exact path below.
|
| 63 |
+
let tie_epsilon = 0.0000003;
|
| 64 |
+
let f = floor(v);
|
| 65 |
+
let frac = v - f;
|
| 66 |
+
return i32(select(f, f + 1.0, frac > 0.5 + tie_epsilon));
|
| 67 |
+
{% endif %}
|
| 68 |
+
}
|
| 69 |
+
|
| 70 |
+
{% if noRoi and (ratioMode or inferredHalfPixelMode) %}
|
| 71 |
+
fn rational_value_nearest_index(numerator: u32, denominator: u32) -> i32 {
|
| 72 |
+
if (numerator == 0u || denominator == 0u) {
|
| 73 |
+
return 0;
|
| 74 |
+
}
|
| 75 |
+
let lower = numerator / denominator;
|
| 76 |
+
{% if source.nearestMode != "floor" %}
|
| 77 |
+
let remainder = numerator - lower * denominator;
|
| 78 |
+
{% endif %}
|
| 79 |
+
{% if source.nearestMode == "floor" %}
|
| 80 |
+
return i32(lower);
|
| 81 |
+
{% elif source.nearestMode == "ceil" %}
|
| 82 |
+
return i32(lower + select(0u, 1u, remainder != 0u));
|
| 83 |
+
{% elif source.nearestMode == "round_prefer_ceil" %}
|
| 84 |
+
return i32(lower + select(0u, 1u, remainder >= denominator - remainder));
|
| 85 |
+
{% else %}
|
| 86 |
+
return i32(lower + select(0u, 1u, remainder > denominator - remainder));
|
| 87 |
+
{% endif %}
|
| 88 |
+
}
|
| 89 |
+
|
| 90 |
+
{% endif %}
|
| 91 |
+
{% if noRoi and ratioMode %}
|
| 92 |
+
fn rational_nearest_index(coord: u32, numerator: u32, denominator: u32) -> i32 {
|
| 93 |
+
if (coord == 0u || numerator == 0u || denominator == 0u) {
|
| 94 |
+
return 0;
|
| 95 |
+
}
|
| 96 |
+
// Use 0xfffffffe rather than UINT_MAX for the overflow guard: some shader
|
| 97 |
+
// compilers reject division with a UINT_MAX constant dividend. The smaller
|
| 98 |
+
// bound is conservative and sends at most one extra coordinate to the f32 path.
|
| 99 |
+
if (coord > 0xfffffffeu / numerator) {
|
| 100 |
+
return nearest_index(f32(coord) * f32(numerator) / f32(denominator));
|
| 101 |
+
}
|
| 102 |
+
return rational_value_nearest_index(coord * numerator, denominator);
|
| 103 |
+
}
|
| 104 |
+
|
| 105 |
+
{% endif %}
|
| 106 |
+
{% if noRoi and inferredHalfPixelMode %}
|
| 107 |
+
fn inferred_half_pixel_nearest_index(
|
| 108 |
+
coord: u32,
|
| 109 |
+
in_size: u32,
|
| 110 |
+
out_size: u32,
|
| 111 |
+
subtract_half: bool
|
| 112 |
+
) -> i32 {
|
| 113 |
+
// (coord + 1/2) * in/out [- 1/2]
|
| 114 |
+
// = ((2*coord + 1)*in [- out]) / (2*out)
|
| 115 |
+
// Use the f32 definition only when the u32 intermediates would overflow.
|
| 116 |
+
if (in_size == 0u || out_size == 0u || out_size > 0x7fffffffu ||
|
| 117 |
+
coord > (0xffffffffu - 1u) / 2u) {
|
| 118 |
+
return nearest_index(source_coord(coord, in_size, out_size, 0.0));
|
| 119 |
+
}
|
| 120 |
+
let odd_coord = coord * 2u + 1u;
|
| 121 |
+
// 0xfffffffe, not 0xffffffff — see rational_nearest_index (NVVM crash).
|
| 122 |
+
if (odd_coord > 0xfffffffeu / in_size) {
|
| 123 |
+
return nearest_index(source_coord(coord, in_size, out_size, 0.0));
|
| 124 |
+
}
|
| 125 |
+
let scaled = odd_coord * in_size;
|
| 126 |
+
let shift = select(0u, out_size, subtract_half);
|
| 127 |
+
// A negative half-pixel coordinate is in [-0.5, 0). Every nearest mode may
|
| 128 |
+
// choose -1 or 0 there, but Resize clamps both to source index zero.
|
| 129 |
+
if (scaled <= shift) {
|
| 130 |
+
return 0;
|
| 131 |
+
}
|
| 132 |
+
return rational_value_nearest_index(scaled - shift, out_size * 2u);
|
| 133 |
+
}
|
| 134 |
+
|
| 135 |
+
{% endif %}
|
| 136 |
+
fn nearest_coord_index(
|
| 137 |
+
out_coord: u32,
|
| 138 |
+
in_size: u32,
|
| 139 |
+
out_size: u32{% if usesExplicitScale %},
|
| 140 |
+
explicit_scale: f32
|
| 141 |
+
{% endif %}
|
| 142 |
+
) -> i32 {
|
| 143 |
+
{% if source.coordinateMode == "align_corners" %}
|
| 144 |
+
if (out_size <= 1u || in_size <= 1u) {
|
| 145 |
+
return 0;
|
| 146 |
+
}
|
| 147 |
+
return rational_nearest_index(out_coord, in_size - 1u, out_size - 1u);
|
| 148 |
+
{% else %}
|
| 149 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 150 |
+
if (explicit_scale <= 0.0) {
|
| 151 |
+
return rational_nearest_index(out_coord, in_size, out_size);
|
| 152 |
+
}
|
| 153 |
+
{% elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 154 |
+
if (out_size <= 1u) {
|
| 155 |
+
return 0;
|
| 156 |
+
}
|
| 157 |
+
if (explicit_scale <= 0.0) {
|
| 158 |
+
return inferred_half_pixel_nearest_index(out_coord, in_size, out_size, true);
|
| 159 |
+
}
|
| 160 |
+
{% elif source.coordinateMode == "tf_crop_and_resize" %}
|
| 161 |
+
{% else %}
|
| 162 |
+
if (explicit_scale <= 0.0) {
|
| 163 |
+
return inferred_half_pixel_nearest_index(out_coord, in_size, out_size, true);
|
| 164 |
+
}
|
| 165 |
+
{% endif %}
|
| 166 |
+
return nearest_index(source_coord(out_coord, in_size, out_size{% if usesExplicitScale %}, explicit_scale{% endif %}));
|
| 167 |
+
{% endif %}
|
| 168 |
+
}
|
| 169 |
+
|
| 170 |
+
|
| 171 |
+
{% else %}
|
| 172 |
+
|
| 173 |
+
{% endif -%}
|
| 174 |
+
fn input_at(base: u32, d: i32, h: i32, w: i32) -> f32 {
|
| 175 |
+
let dd = u32(clamp(d, 0, i32(params.inD) - 1));
|
| 176 |
+
let hh = u32(clamp(h, 0, i32(params.inH) - 1));
|
| 177 |
+
let ww = u32(clamp(w, 0, i32(params.inW) - 1));
|
| 178 |
+
return f32(x[base + (dd * params.inH + hh) * params.inW + ww]);
|
| 179 |
+
}
|
| 180 |
+
|
| 181 |
+
@compute @workgroup_size({{ wg }})
|
| 182 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
|
| 183 |
+
let i = gid.x + gid.y * nwg.x * {{ wg }}u;
|
| 184 |
+
if (i >= params.count) { return; }
|
| 185 |
+
let ow = i % params.outW;
|
| 186 |
+
let oh = (i / params.outW) % params.outH;
|
| 187 |
+
let od = (i / (params.outW * params.outH)) % params.outD;
|
| 188 |
+
let c = (i / (params.outW * params.outH * params.outD)) % params.C;
|
| 189 |
+
let n = i / (params.outW * params.outH * params.outD * params.C);
|
| 190 |
+
let base = (n * params.C + c) * params.inD * params.inH * params.inW;
|
| 191 |
+
{% if source.mode != "nearest" %}
|
| 192 |
+
let sd = source_coord(od, params.inD, params.outD{% if not alignCorners %}, params.scaleD{% endif %});
|
| 193 |
+
let sy = source_coord(oh, params.inH, params.outH{% if not alignCorners %}, params.scaleH{% endif %});
|
| 194 |
+
let sx = source_coord(ow, params.inW, params.outW{% if not alignCorners %}, params.scaleW{% endif %});
|
| 195 |
+
{% endif %}
|
| 196 |
+
{% if source.mode == "nearest" %}
|
| 197 |
+
let id = nearest_coord_index(od, params.inD, params.outD{% if not alignCorners %}, params.scaleD{% endif %});
|
| 198 |
+
let ih = nearest_coord_index(oh, params.inH, params.outH{% if not alignCorners %}, params.scaleH{% endif %});
|
| 199 |
+
let iw = nearest_coord_index(ow, params.inW, params.outW{% if not alignCorners %}, params.scaleW{% endif %});
|
| 200 |
+
let result = input_at(base, id, ih, iw);
|
| 201 |
+
{% else %}
|
| 202 |
+
let z0 = i32(floor(sd));
|
| 203 |
+
let y0 = i32(floor(sy));
|
| 204 |
+
let x0 = i32(floor(sx));
|
| 205 |
+
let z1 = z0 + 1;
|
| 206 |
+
let y1 = y0 + 1;
|
| 207 |
+
let x1 = x0 + 1;
|
| 208 |
+
let wz = sd - f32(z0);
|
| 209 |
+
let wy = sy - f32(y0);
|
| 210 |
+
let wx = sx - f32(x0);
|
| 211 |
+
let c000 = input_at(base, z0, y0, x0);
|
| 212 |
+
let c001 = input_at(base, z0, y0, x1);
|
| 213 |
+
let c010 = input_at(base, z0, y1, x0);
|
| 214 |
+
let c011 = input_at(base, z0, y1, x1);
|
| 215 |
+
let c100 = input_at(base, z1, y0, x0);
|
| 216 |
+
let c101 = input_at(base, z1, y0, x1);
|
| 217 |
+
let c110 = input_at(base, z1, y1, x0);
|
| 218 |
+
let c111 = input_at(base, z1, y1, x1);
|
| 219 |
+
let front = (1.0 - wy) * ((1.0 - wx) * c000 + wx * c001) + wy * ((1.0 - wx) * c010 + wx * c011);
|
| 220 |
+
let back = (1.0 - wy) * ((1.0 - wx) * c100 + wx * c101) + wy * ((1.0 - wx) * c110 + wx * c111);
|
| 221 |
+
let result = (1.0 - wz) * front + wz * back;
|
| 222 |
+
{% endif %}
|
| 223 |
+
y[i] = {{ scalar }}(result);
|
| 224 |
+
}
|
build/webgpu/resize-coord-transform.wgsl.jinja
ADDED
|
@@ -0,0 +1,313 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// Generic NCHW Resize/Upsample with full coordinate-transformation math.
|
| 2 |
+
// Explicit scales of zero are inferred from the output/input shape ratio.
|
| 3 |
+
// Nearest-index rounding is load-bearing for ONNX compatibility:
|
| 4 |
+
// round_prefer_ceil ties toward positive infinity, while round_prefer_floor
|
| 5 |
+
// uses its specified epsilon.
|
| 6 |
+
{% set alignCorners = source.coordinateMode == "align_corners" %}
|
| 7 |
+
{% if usesF16 %}
|
| 8 |
+
enable f16;
|
| 9 |
+
{% endif %}
|
| 10 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 11 |
+
{% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
|
| 12 |
+
|
| 13 |
+
{% if source.mode != "nearest" or not alignCorners %}fn source_coord(out_coord: u32, in_size: u32, out_size: u32{% if not alignCorners %}, explicit_scale: f32{% endif %}) -> f32 {
|
| 14 |
+
{% if not alignCorners %}
|
| 15 |
+
{% if source.coordinateMode != "asymmetric" %}
|
| 16 |
+
let inferred_scale = f32(out_size) / f32(in_size);
|
| 17 |
+
let scale = select(inferred_scale, explicit_scale, explicit_scale > 0.0);
|
| 18 |
+
{% endif %}
|
| 19 |
+
{% endif %}
|
| 20 |
+
{% if alignCorners %}
|
| 21 |
+
if (out_size == 1u) { return 0.0; }
|
| 22 |
+
return f32(out_coord) * f32(in_size - 1u) / f32(out_size - 1u);
|
| 23 |
+
{% elif source.coordinateMode == "asymmetric" %}
|
| 24 |
+
// With an inferred scale, divide by sizes (out_coord*in/out) rather than by
|
| 25 |
+
// the precomputed reciprocal out/in: the single division keeps exact ties
|
| 26 |
+
// (e.g. a 7->2 downsample at out col 1 == 3.5) exact instead of landing at
|
| 27 |
+
// 3.4999998 after double rounding. An explicit scale uses coord/scale.
|
| 28 |
+
return select(f32(out_coord) * f32(in_size) / f32(out_size), f32(out_coord) / explicit_scale, explicit_scale > 0.0);
|
| 29 |
+
{% elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 30 |
+
return select(0.0, (f32(out_coord) + 0.5) / scale - 0.5, out_size > 1u);
|
| 31 |
+
{% elif source.coordinateMode == "half_pixel_symmetric" %}
|
| 32 |
+
let output_width = f32(in_size) * scale;
|
| 33 |
+
let adjustment = f32(out_size) / output_width;
|
| 34 |
+
let center = f32(in_size) * 0.5;
|
| 35 |
+
let offset = center * (1.0 - adjustment);
|
| 36 |
+
return offset + (f32(out_coord) + 0.5) / scale - 0.5;
|
| 37 |
+
{% else %}
|
| 38 |
+
return (f32(out_coord) + 0.5) / scale - 0.5;
|
| 39 |
+
{% endif %}
|
| 40 |
+
}{% endif %}
|
| 41 |
+
|
| 42 |
+
/* Backend-stable nearest-neighbor coordinate selection. Inferred align-corners,
|
| 43 |
+
* asymmetric, and half-pixel transforms use exact integer ratios. This matters
|
| 44 |
+
* at nearest-mode ties: reassociating an f32 division into a reciprocal
|
| 45 |
+
* multiply can move an exact 0.5 to either side of the boundary. */{% set ratioMode = source.coordinateMode in ["align_corners", "asymmetric"] %}
|
| 46 |
+
{% set inferredHalfPixelMode = source.coordinateMode in ["pytorch_half_pixel", "half_pixel", "half_pixel_symmetric"] %}
|
| 47 |
+
{% set noRoi = source.roi is not defined or (source.roi | length) == 0 %}
|
| 48 |
+
{% set usesExplicitScale = source.coordinateMode not in ["align_corners", "tf_crop_and_resize"] %}
|
| 49 |
+
{% if source.mode == "nearest" %}
|
| 50 |
+
fn nearest_index(v: f32) -> i32 {
|
| 51 |
+
{% if source.nearestMode == "floor" %}
|
| 52 |
+
return i32(floor(v));
|
| 53 |
+
{% elif source.nearestMode == "ceil" %}
|
| 54 |
+
return i32(ceil(v));
|
| 55 |
+
{% elif source.nearestMode == "round_prefer_ceil" %}
|
| 56 |
+
// round half toward +inf. WGSL round() is round-half-to-even and therefore
|
| 57 |
+
// cannot express ONNX's preferred direction. The lower-side epsilon absorbs
|
| 58 |
+
// backend division error that lands a mathematical half-tie just below .5
|
| 59 |
+
// (mirror of the round_prefer_floor epsilon below).
|
| 60 |
+
let tie_epsilon = 0.0000003;
|
| 61 |
+
let f = floor(v);
|
| 62 |
+
let frac = v - f;
|
| 63 |
+
return i32(select(f, f + 1.0, frac >= 0.5 - tie_epsilon));
|
| 64 |
+
{% else %}
|
| 65 |
+
// Explicit scales still require f32 coordinate math. Keep a narrow epsilon
|
| 66 |
+
// for their backend rounding while inferred scales take the exact path below.
|
| 67 |
+
let tie_epsilon = 0.0000003;
|
| 68 |
+
let f = floor(v);
|
| 69 |
+
let frac = v - f;
|
| 70 |
+
return i32(select(f, f + 1.0, frac > 0.5 + tie_epsilon));
|
| 71 |
+
{% endif %}
|
| 72 |
+
}
|
| 73 |
+
|
| 74 |
+
{% if noRoi and (ratioMode or inferredHalfPixelMode) %}
|
| 75 |
+
fn rational_value_nearest_index(numerator: u32, denominator: u32) -> i32 {
|
| 76 |
+
if (numerator == 0u || denominator == 0u) {
|
| 77 |
+
return 0;
|
| 78 |
+
}
|
| 79 |
+
let lower = numerator / denominator;
|
| 80 |
+
{% if source.nearestMode != "floor" %}
|
| 81 |
+
let remainder = numerator - lower * denominator;
|
| 82 |
+
{% endif %}
|
| 83 |
+
{% if source.nearestMode == "floor" %}
|
| 84 |
+
return i32(lower);
|
| 85 |
+
{% elif source.nearestMode == "ceil" %}
|
| 86 |
+
return i32(lower + select(0u, 1u, remainder != 0u));
|
| 87 |
+
{% elif source.nearestMode == "round_prefer_ceil" %}
|
| 88 |
+
return i32(lower + select(0u, 1u, remainder >= denominator - remainder));
|
| 89 |
+
{% else %}
|
| 90 |
+
return i32(lower + select(0u, 1u, remainder > denominator - remainder));
|
| 91 |
+
{% endif %}
|
| 92 |
+
}
|
| 93 |
+
|
| 94 |
+
{% endif %}
|
| 95 |
+
{% if noRoi and ratioMode %}
|
| 96 |
+
fn rational_nearest_index(coord: u32, numerator: u32, denominator: u32) -> i32 {
|
| 97 |
+
if (coord == 0u || numerator == 0u || denominator == 0u) {
|
| 98 |
+
return 0;
|
| 99 |
+
}
|
| 100 |
+
// Use 0xfffffffe rather than UINT_MAX for the overflow guard: some shader
|
| 101 |
+
// compilers reject division with a UINT_MAX constant dividend. The smaller
|
| 102 |
+
// bound is conservative and sends at most one extra coordinate to the f32 path.
|
| 103 |
+
if (coord > 0xfffffffeu / numerator) {
|
| 104 |
+
return nearest_index(f32(coord) * f32(numerator) / f32(denominator));
|
| 105 |
+
}
|
| 106 |
+
return rational_value_nearest_index(coord * numerator, denominator);
|
| 107 |
+
}
|
| 108 |
+
|
| 109 |
+
{% endif %}
|
| 110 |
+
{% if noRoi and inferredHalfPixelMode %}
|
| 111 |
+
fn inferred_half_pixel_nearest_index(
|
| 112 |
+
coord: u32,
|
| 113 |
+
in_size: u32,
|
| 114 |
+
out_size: u32,
|
| 115 |
+
subtract_half: bool
|
| 116 |
+
) -> i32 {
|
| 117 |
+
// (coord + 1/2) * in/out [- 1/2]
|
| 118 |
+
// = ((2*coord + 1)*in [- out]) / (2*out)
|
| 119 |
+
// Use the f32 definition only when the u32 intermediates would overflow.
|
| 120 |
+
if (in_size == 0u || out_size == 0u || out_size > 0x7fffffffu ||
|
| 121 |
+
coord > (0xffffffffu - 1u) / 2u) {
|
| 122 |
+
return nearest_index(source_coord(coord, in_size, out_size, 0.0));
|
| 123 |
+
}
|
| 124 |
+
let odd_coord = coord * 2u + 1u;
|
| 125 |
+
// 0xfffffffe, not 0xffffffff — see rational_nearest_index (NVVM crash).
|
| 126 |
+
if (odd_coord > 0xfffffffeu / in_size) {
|
| 127 |
+
return nearest_index(source_coord(coord, in_size, out_size, 0.0));
|
| 128 |
+
}
|
| 129 |
+
let scaled = odd_coord * in_size;
|
| 130 |
+
let shift = select(0u, out_size, subtract_half);
|
| 131 |
+
// A negative half-pixel coordinate is in [-0.5, 0). Every nearest mode may
|
| 132 |
+
// choose -1 or 0 there, but Resize clamps both to source index zero.
|
| 133 |
+
if (scaled <= shift) {
|
| 134 |
+
return 0;
|
| 135 |
+
}
|
| 136 |
+
return rational_value_nearest_index(scaled - shift, out_size * 2u);
|
| 137 |
+
}
|
| 138 |
+
|
| 139 |
+
{% endif %}
|
| 140 |
+
fn nearest_coord_index(
|
| 141 |
+
out_coord: u32,
|
| 142 |
+
in_size: u32,
|
| 143 |
+
out_size: u32{% if usesExplicitScale %},
|
| 144 |
+
explicit_scale: f32
|
| 145 |
+
{% endif %}
|
| 146 |
+
) -> i32 {
|
| 147 |
+
{% if source.coordinateMode == "align_corners" %}
|
| 148 |
+
if (out_size <= 1u || in_size <= 1u) {
|
| 149 |
+
return 0;
|
| 150 |
+
}
|
| 151 |
+
return rational_nearest_index(out_coord, in_size - 1u, out_size - 1u);
|
| 152 |
+
{% else %}
|
| 153 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 154 |
+
if (explicit_scale <= 0.0) {
|
| 155 |
+
return rational_nearest_index(out_coord, in_size, out_size);
|
| 156 |
+
}
|
| 157 |
+
{% elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 158 |
+
if (out_size <= 1u) {
|
| 159 |
+
return 0;
|
| 160 |
+
}
|
| 161 |
+
if (explicit_scale <= 0.0) {
|
| 162 |
+
return inferred_half_pixel_nearest_index(out_coord, in_size, out_size, true);
|
| 163 |
+
}
|
| 164 |
+
{% elif source.coordinateMode == "tf_crop_and_resize" %}
|
| 165 |
+
{% else %}
|
| 166 |
+
if (explicit_scale <= 0.0) {
|
| 167 |
+
return inferred_half_pixel_nearest_index(out_coord, in_size, out_size, true);
|
| 168 |
+
}
|
| 169 |
+
{% endif %}
|
| 170 |
+
return nearest_index(source_coord(out_coord, in_size, out_size{% if usesExplicitScale %}, explicit_scale{% endif %}));
|
| 171 |
+
{% endif %}
|
| 172 |
+
}
|
| 173 |
+
|
| 174 |
+
|
| 175 |
+
{% else %}
|
| 176 |
+
|
| 177 |
+
{% endif -%}
|
| 178 |
+
fn input_at(base: u32, h: i32, w: i32) -> f32 {
|
| 179 |
+
let hh = u32(clamp(h, 0, i32(params.inH) - 1));
|
| 180 |
+
let ww = u32(clamp(w, 0, i32(params.inW) - 1));
|
| 181 |
+
return f32(x[base + hh * params.inW + ww]);
|
| 182 |
+
}
|
| 183 |
+
|
| 184 |
+
{% if source.mode == "linear" %}
|
| 185 |
+
// Clamp before choosing neighbors. At an edge both neighbors coincide and
|
| 186 |
+
// receive half weight; the four-term order is observable for integer outputs.
|
| 187 |
+
fn bilinear_sample(base: u32, source_y: f32, source_x: f32) -> f32 {
|
| 188 |
+
let sy = clamp(source_y, 0.0, f32(params.inH - 1u));
|
| 189 |
+
let sx = clamp(source_x, 0.0, f32(params.inW - 1u));
|
| 190 |
+
let y0 = i32(floor(sy));
|
| 191 |
+
let x0 = i32(floor(sx));
|
| 192 |
+
let y1 = min(y0 + 1, i32(params.inH) - 1);
|
| 193 |
+
let x1 = min(x0 + 1, i32(params.inW) - 1);
|
| 194 |
+
let same_y = y0 == y1;
|
| 195 |
+
let same_x = x0 == x1;
|
| 196 |
+
let wy0 = select(abs(sy - f32(y1)), 0.5, same_y);
|
| 197 |
+
let wy1 = select(abs(sy - f32(y0)), 0.5, same_y);
|
| 198 |
+
let wx0 = select(abs(sx - f32(x1)), 0.5, same_x);
|
| 199 |
+
let wx1 = select(abs(sx - f32(x0)), 0.5, same_x);
|
| 200 |
+
let v00 = input_at(base, y0, x0);
|
| 201 |
+
let v01 = input_at(base, y0, x1);
|
| 202 |
+
let v10 = input_at(base, y1, x0);
|
| 203 |
+
let v11 = input_at(base, y1, x1);
|
| 204 |
+
return wx0 * wy0 * v00 + wx1 * wy0 * v01 + wx0 * wy1 * v10 + wx1 * wy1 * v11;
|
| 205 |
+
}
|
| 206 |
+
{% endif %}
|
| 207 |
+
|
| 208 |
+
@compute @workgroup_size({{ wg }})
|
| 209 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
|
| 210 |
+
// 2D-folded flat index: gid.y carries the high bits past the
|
| 211 |
+
// maxComputeWorkgroupsPerDimension limit.
|
| 212 |
+
let flat = gid.x + gid.y * nwg.x * {{ wg }}u;
|
| 213 |
+
{% if not alignCorners %}
|
| 214 |
+
let scaleH_val = params.scaleH;
|
| 215 |
+
let scaleW_val = params.scaleW;
|
| 216 |
+
{% endif %}
|
| 217 |
+
{% if source.vec4 %}
|
| 218 |
+
// One thread owns 4 contiguous output columns (outW % 4 == 0, so a 4-group
|
| 219 |
+
// never crosses a row): all 4 lanes share oh/c/n and the source row coord sy;
|
| 220 |
+
// only sx differs per lane. One 128-bit vec4 store. Bit-identical to scalar.
|
| 221 |
+
let count4 = params.count / 4u;
|
| 222 |
+
if (flat >= count4) { return; }
|
| 223 |
+
let i = flat * 4u;
|
| 224 |
+
let ow0 = i % params.outW;
|
| 225 |
+
let row = i / params.outW;
|
| 226 |
+
let oh = row % params.outH;
|
| 227 |
+
let c = (row / params.outH) % params.C;
|
| 228 |
+
let n = row / (params.outH * params.C);
|
| 229 |
+
let base = (n * params.C + c) * params.inH * params.inW;
|
| 230 |
+
{% if source.mode != "nearest" %}
|
| 231 |
+
let sy = source_coord(oh, params.inH, params.outH{% if not alignCorners %}, scaleH_val{% endif %});
|
| 232 |
+
let sx0 = source_coord(ow0, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 233 |
+
let sx1 = source_coord(ow0 + 1u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 234 |
+
let sx2 = source_coord(ow0 + 2u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 235 |
+
let sx3 = source_coord(ow0 + 3u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 236 |
+
{% endif %}
|
| 237 |
+
{% if source.mode == "nearest" %}
|
| 238 |
+
let ih = nearest_coord_index(oh, params.inH, params.outH{% if not alignCorners %}, scaleH_val{% endif %});
|
| 239 |
+
let r = vec4<f32>(
|
| 240 |
+
input_at(base, ih, nearest_coord_index(ow0, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})),
|
| 241 |
+
input_at(base, ih, nearest_coord_index(ow0 + 1u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})),
|
| 242 |
+
input_at(base, ih, nearest_coord_index(ow0 + 2u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})),
|
| 243 |
+
input_at(base, ih, nearest_coord_index(ow0 + 3u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})));
|
| 244 |
+
{% else %}
|
| 245 |
+
let r = vec4<f32>(
|
| 246 |
+
bilinear_sample(base, sy, sx0),
|
| 247 |
+
bilinear_sample(base, sy, sx1),
|
| 248 |
+
bilinear_sample(base, sy, sx2),
|
| 249 |
+
bilinear_sample(base, sy, sx3));
|
| 250 |
+
{% endif %}
|
| 251 |
+
y[flat] = {{ source.vecType }}(r);
|
| 252 |
+
{% elif source.scalarX4 %}
|
| 253 |
+
// Four columns per thread without a vec4 storage binding. Unlike the packed
|
| 254 |
+
// route this remains valid for odd outW: group indexing is row-local and the
|
| 255 |
+
// final 1--3 lanes are guarded scalar stores.
|
| 256 |
+
let groups_per_row = (params.outW + 3u) / 4u;
|
| 257 |
+
let group_count = (params.count / params.outW) * groups_per_row;
|
| 258 |
+
if (flat >= group_count) { return; }
|
| 259 |
+
let group_in_row = flat % groups_per_row;
|
| 260 |
+
let row = flat / groups_per_row;
|
| 261 |
+
let ow0 = group_in_row * 4u;
|
| 262 |
+
let oh = row % params.outH;
|
| 263 |
+
let c = (row / params.outH) % params.C;
|
| 264 |
+
let n = row / (params.outH * params.C);
|
| 265 |
+
let base = (n * params.C + c) * params.inH * params.inW;
|
| 266 |
+
{% if source.mode != "nearest" %}
|
| 267 |
+
let sy = source_coord(oh, params.inH, params.outH{% if not alignCorners %}, scaleH_val{% endif %});
|
| 268 |
+
let sx0 = source_coord(ow0, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 269 |
+
let sx1 = source_coord(ow0 + 1u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 270 |
+
let sx2 = source_coord(ow0 + 2u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 271 |
+
let sx3 = source_coord(ow0 + 3u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 272 |
+
{% endif %}
|
| 273 |
+
{% if source.mode == "nearest" %}
|
| 274 |
+
let ih = nearest_coord_index(oh, params.inH, params.outH{% if not alignCorners %}, scaleH_val{% endif %});
|
| 275 |
+
let r = vec4<f32>(
|
| 276 |
+
input_at(base, ih, nearest_coord_index(ow0, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})),
|
| 277 |
+
input_at(base, ih, nearest_coord_index(ow0 + 1u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})),
|
| 278 |
+
input_at(base, ih, nearest_coord_index(ow0 + 2u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})),
|
| 279 |
+
input_at(base, ih, nearest_coord_index(ow0 + 3u, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %})));
|
| 280 |
+
{% else %}
|
| 281 |
+
let r = vec4<f32>(
|
| 282 |
+
bilinear_sample(base, sy, sx0),
|
| 283 |
+
bilinear_sample(base, sy, sx1),
|
| 284 |
+
bilinear_sample(base, sy, sx2),
|
| 285 |
+
bilinear_sample(base, sy, sx3));
|
| 286 |
+
{% endif %}
|
| 287 |
+
let i = row * params.outW + ow0;
|
| 288 |
+
y[i] = {{ scalar }}(r.x);
|
| 289 |
+
if (ow0 + 1u < params.outW) { y[i + 1u] = {{ scalar }}(r.y); }
|
| 290 |
+
if (ow0 + 2u < params.outW) { y[i + 2u] = {{ scalar }}(r.z); }
|
| 291 |
+
if (ow0 + 3u < params.outW) { y[i + 3u] = {{ scalar }}(r.w); }
|
| 292 |
+
{% else %}
|
| 293 |
+
let i = flat;
|
| 294 |
+
if (i >= params.count) { return; }
|
| 295 |
+
let ow = i % params.outW;
|
| 296 |
+
let oh = (i / params.outW) % params.outH;
|
| 297 |
+
let c = (i / (params.outW * params.outH)) % params.C;
|
| 298 |
+
let n = i / (params.outW * params.outH * params.C);
|
| 299 |
+
let base = (n * params.C + c) * params.inH * params.inW;
|
| 300 |
+
{% if source.mode != "nearest" %}
|
| 301 |
+
let sy = source_coord(oh, params.inH, params.outH{% if not alignCorners %}, scaleH_val{% endif %});
|
| 302 |
+
let sx = source_coord(ow, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 303 |
+
{% endif %}
|
| 304 |
+
{% if source.mode == "nearest" %}
|
| 305 |
+
let ih = nearest_coord_index(oh, params.inH, params.outH{% if not alignCorners %}, scaleH_val{% endif %});
|
| 306 |
+
let iw = nearest_coord_index(ow, params.inW, params.outW{% if not alignCorners %}, scaleW_val{% endif %});
|
| 307 |
+
let result = input_at(base, ih, iw);
|
| 308 |
+
{% else %}
|
| 309 |
+
let result = bilinear_sample(base, sy, sx);
|
| 310 |
+
{% endif %}
|
| 311 |
+
y[i] = {{ scalar }}(result);
|
| 312 |
+
{% endif %}
|
| 313 |
+
}
|
build/webgpu/resize-cubic.wgsl.jinja
ADDED
|
@@ -0,0 +1,118 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// Generic NCHW Resize with separable Keys cubic interpolation. Each output
|
| 2 |
+
// coordinate maps to input space and samples the four taps around floor(src)
|
| 3 |
+
// at offsets {-1, 0, 1, 2}. The rows-then-columns pass uses 16 taps in 2D.
|
| 4 |
+
// Sample coordinates clamp to [0, dim - 1] when exclude_outside is disabled.
|
| 5 |
+
// The coefficient `a` is cubic_coeff_a (default -0.75), and half_pixel,
|
| 6 |
+
// pytorch_half_pixel, asymmetric, and align_corners follow the same coordinate
|
| 7 |
+
// rules as the other Resize kernels.
|
| 8 |
+
{% set EXCL = source.excludeOutside is defined and source.excludeOutside == 1 %}
|
| 9 |
+
{% if usesF16 %}
|
| 10 |
+
enable f16;
|
| 11 |
+
{% endif %}
|
| 12 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 13 |
+
{% set wg = workgroupSize %}
|
| 14 |
+
|
| 15 |
+
fn source_coord(out_coord: u32, in_size: u32, out_size: u32{% if source.coordinateMode != "align_corners" %}, explicit_scale: f32{% endif %}) -> f32 {
|
| 16 |
+
{% if source.coordinateMode != "align_corners" %}
|
| 17 |
+
let inferred_scale = f32(out_size) / f32(in_size);
|
| 18 |
+
let scale = select(inferred_scale, explicit_scale, explicit_scale > 0.0);
|
| 19 |
+
{% endif %}
|
| 20 |
+
{% if source.coordinateMode == "align_corners" %}
|
| 21 |
+
if (out_size == 1u) { return 0.0; }
|
| 22 |
+
return f32(out_coord) * f32(in_size - 1u) / f32(out_size - 1u);
|
| 23 |
+
{% elif source.coordinateMode == "asymmetric" %}
|
| 24 |
+
return f32(out_coord) / scale;
|
| 25 |
+
{% elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 26 |
+
return select(0.0, (f32(out_coord) + 0.5) / scale - 0.5, out_size > 1u);
|
| 27 |
+
{% else %}
|
| 28 |
+
return (f32(out_coord) + 0.5) / scale - 0.5;
|
| 29 |
+
{% endif %}
|
| 30 |
+
}
|
| 31 |
+
|
| 32 |
+
// Keys cubic convolution kernel with coefficient params.coeffA.
|
| 33 |
+
fn cubic_weight(s_in: f32) -> f32 {
|
| 34 |
+
let a = params.coeffA;
|
| 35 |
+
let s = abs(s_in);
|
| 36 |
+
let s2 = s * s;
|
| 37 |
+
let s3 = s2 * s;
|
| 38 |
+
if (s <= 1.0) {
|
| 39 |
+
return (a + 2.0) * s3 - (a + 3.0) * s2 + 1.0;
|
| 40 |
+
}
|
| 41 |
+
if (s < 2.0) {
|
| 42 |
+
return a * s3 - 5.0 * a * s2 + 8.0 * a * s - 4.0 * a;
|
| 43 |
+
}
|
| 44 |
+
return 0.0;
|
| 45 |
+
}
|
| 46 |
+
|
| 47 |
+
fn input_at(base: u32, h: i32, w: i32) -> f32 {
|
| 48 |
+
let hh = u32(clamp(h, 0, i32(params.inH) - 1));
|
| 49 |
+
let ww = u32(clamp(w, 0, i32(params.inW) - 1));
|
| 50 |
+
return f32(x[base + hh * params.inW + ww]);
|
| 51 |
+
}
|
| 52 |
+
|
| 53 |
+
// Cubic interpolation along the W axis for a fixed input row `h`, at source
|
| 54 |
+
// column `sx`. Four taps at floor(sx)-1 .. floor(sx)+2.
|
| 55 |
+
fn cubic_x(base: u32, h: i32, sx: f32) -> f32 {
|
| 56 |
+
let x0 = i32(floor(sx));
|
| 57 |
+
let t = sx - f32(x0);
|
| 58 |
+
var w0 = cubic_weight(1.0 + t);
|
| 59 |
+
var w1 = cubic_weight(t);
|
| 60 |
+
var w2 = cubic_weight(1.0 - t);
|
| 61 |
+
var w3 = cubic_weight(2.0 - t);
|
| 62 |
+
{% if EXCL %}
|
| 63 |
+
// exclude_outside: zero taps outside the input range, renormalize the rest.
|
| 64 |
+
let wlim = i32(params.inW);
|
| 65 |
+
if (x0 - 1 < 0 || x0 - 1 >= wlim) { w0 = 0.0; }
|
| 66 |
+
if (x0 < 0 || x0 >= wlim) { w1 = 0.0; }
|
| 67 |
+
if (x0 + 1 < 0 || x0 + 1 >= wlim) { w2 = 0.0; }
|
| 68 |
+
if (x0 + 2 < 0 || x0 + 2 >= wlim) { w3 = 0.0; }
|
| 69 |
+
let wsum = w0 + w1 + w2 + w3;
|
| 70 |
+
w0 = w0 / wsum;
|
| 71 |
+
w1 = w1 / wsum;
|
| 72 |
+
w2 = w2 / wsum;
|
| 73 |
+
w3 = w3 / wsum;
|
| 74 |
+
{% endif %}
|
| 75 |
+
return w0 * input_at(base, h, x0 - 1)
|
| 76 |
+
+ w1 * input_at(base, h, x0)
|
| 77 |
+
+ w2 * input_at(base, h, x0 + 1)
|
| 78 |
+
+ w3 * input_at(base, h, x0 + 2);
|
| 79 |
+
}
|
| 80 |
+
|
| 81 |
+
@compute @workgroup_size({{ wg }})
|
| 82 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
|
| 83 |
+
// 2D-folded flat index: gid.y carries the high bits past the
|
| 84 |
+
// maxComputeWorkgroupsPerDimension limit.
|
| 85 |
+
let i = gid.x + gid.y * nwg.x * {{ wg }}u;
|
| 86 |
+
if (i >= params.count) { return; }
|
| 87 |
+
let ow = i % params.outW;
|
| 88 |
+
let oh = (i / params.outW) % params.outH;
|
| 89 |
+
let c = (i / (params.outW * params.outH)) % params.C;
|
| 90 |
+
let n = i / (params.outW * params.outH * params.C);
|
| 91 |
+
let base = (n * params.C + c) * params.inH * params.inW;
|
| 92 |
+
let sy = source_coord(oh, params.inH, params.outH{% if source.coordinateMode != "align_corners" %}, params.scaleH{% endif %});
|
| 93 |
+
let sx = source_coord(ow, params.inW, params.outW{% if source.coordinateMode != "align_corners" %}, params.scaleW{% endif %});
|
| 94 |
+
// Separable: cubic along W for each of the 4 source rows, then cubic along H.
|
| 95 |
+
let y0 = i32(floor(sy));
|
| 96 |
+
let t = sy - f32(y0);
|
| 97 |
+
var wy0 = cubic_weight(1.0 + t);
|
| 98 |
+
var wy1 = cubic_weight(t);
|
| 99 |
+
var wy2 = cubic_weight(1.0 - t);
|
| 100 |
+
var wy3 = cubic_weight(2.0 - t);
|
| 101 |
+
{% if EXCL %}
|
| 102 |
+
let hlim = i32(params.inH);
|
| 103 |
+
if (y0 - 1 < 0 || y0 - 1 >= hlim) { wy0 = 0.0; }
|
| 104 |
+
if (y0 < 0 || y0 >= hlim) { wy1 = 0.0; }
|
| 105 |
+
if (y0 + 1 < 0 || y0 + 1 >= hlim) { wy2 = 0.0; }
|
| 106 |
+
if (y0 + 2 < 0 || y0 + 2 >= hlim) { wy3 = 0.0; }
|
| 107 |
+
let wysum = wy0 + wy1 + wy2 + wy3;
|
| 108 |
+
wy0 = wy0 / wysum;
|
| 109 |
+
wy1 = wy1 / wysum;
|
| 110 |
+
wy2 = wy2 / wysum;
|
| 111 |
+
wy3 = wy3 / wysum;
|
| 112 |
+
{% endif %}
|
| 113 |
+
let result = wy0 * cubic_x(base, y0 - 1, sx)
|
| 114 |
+
+ wy1 * cubic_x(base, y0, sx)
|
| 115 |
+
+ wy2 * cubic_x(base, y0 + 1, sx)
|
| 116 |
+
+ wy3 * cubic_x(base, y0 + 2, sx);
|
| 117 |
+
y[i] = {{ scalar }}(result);
|
| 118 |
+
}
|
build/webgpu/resize-generic.wgsl.jinja
ADDED
|
@@ -0,0 +1,409 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// Rank-generic Resize for arbitrary ranks and resized axes, including
|
| 2 |
+
// tf_crop_and_resize with an ROI and extrapolation value. The per-axis
|
| 3 |
+
// coordinate transform follows the rank-specialized paths so half_pixel,
|
| 4 |
+
// pytorch_half_pixel, asymmetric, and align_corners stay
|
| 5 |
+
// consistent. Linear interpolation runs only over resized axes; passthrough
|
| 6 |
+
// axes contribute one tap.
|
| 7 |
+
{% if usesF16 %}
|
| 8 |
+
enable f16;
|
| 9 |
+
{% endif %}
|
| 10 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 11 |
+
{% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
|
| 12 |
+
{%- set xs = source.xShape %}
|
| 13 |
+
{% set ys = source.yShape %}
|
| 14 |
+
{% set r = xs | length -%}
|
| 15 |
+
{% set usesExplicitScale = source.coordinateMode not in ["align_corners", "tf_crop_and_resize"] %}
|
| 16 |
+
|
| 17 |
+
{% set resolvedScales = namespace(values = []) %}
|
| 18 |
+
{% for a in range(r) %}
|
| 19 |
+
{% set axisScale = namespace(value = 0) %}
|
| 20 |
+
{% if source.scales is defined and source.scales %}
|
| 21 |
+
{% if source.axes is defined and source.axes %}
|
| 22 |
+
{% for j in range(source.axes | length) %}
|
| 23 |
+
{% set scaleAxis = source.axes[j] if source.axes[j] >= 0 else source.axes[j] + r %}
|
| 24 |
+
{% if scaleAxis == a %}{% set axisScale.value = source.scales[j] %}{% endif %}
|
| 25 |
+
{% endfor %}
|
| 26 |
+
{% else %}
|
| 27 |
+
{% set axisScale.value = source.scales[a] %}
|
| 28 |
+
{% endif %}
|
| 29 |
+
{% endif %}
|
| 30 |
+
{% set resolvedScales.values = resolvedScales.values + [axisScale.value] %}
|
| 31 |
+
{% endfor %}
|
| 32 |
+
{% set staticScales = resolvedScales.values %}
|
| 33 |
+
|
| 34 |
+
{% set roiStarts = namespace(values = []) %}
|
| 35 |
+
{% set roiEnds = namespace(values = []) %}
|
| 36 |
+
{% for a in range(r) %}
|
| 37 |
+
{% set bounds = namespace(start = 0, end = 1) %}
|
| 38 |
+
{% if source.roi is defined and source.roi %}
|
| 39 |
+
{% if source.axes is defined and source.axes %}
|
| 40 |
+
{% for j in range(source.axes | length) %}
|
| 41 |
+
{% set roiAxis = source.axes[j] if source.axes[j] >= 0 else source.axes[j] + r %}
|
| 42 |
+
{% if roiAxis == a %}
|
| 43 |
+
{% set bounds.start = source.roi[j] %}
|
| 44 |
+
{% set bounds.end = source.roi[j + (source.axes | length)] %}
|
| 45 |
+
{% endif %}
|
| 46 |
+
{% endfor %}
|
| 47 |
+
{% else %}
|
| 48 |
+
{% set bounds.start = source.roi[a] %}
|
| 49 |
+
{% set bounds.end = source.roi[a + r] %}
|
| 50 |
+
{% endif %}
|
| 51 |
+
{% endif %}
|
| 52 |
+
{% set roiStarts.values = roiStarts.values + [bounds.start] %}
|
| 53 |
+
{% set roiEnds.values = roiEnds.values + [bounds.end] %}
|
| 54 |
+
{% endfor %}
|
| 55 |
+
{% set staticRoi = roiStarts.values + roiEnds.values %}{% set ocount = namespace(v = 1) %}
|
| 56 |
+
{% for a in range(r) %}{% set ocount.v = ocount.v * ys[a] %}{% endfor %}{% set resized = namespace(list = []) %}
|
| 57 |
+
{% for a in range(r) %}
|
| 58 |
+
{% if xs[a] != ys[a] or (source.scalesBuffer | default(false)) or (staticScales[a] > 0 and staticScales[a] != 1) or ((source.roi is defined and source.roi) and (staticRoi[a] != 0 or staticRoi[a + r] != 1)) %}
|
| 59 |
+
{% set resized.list = resized.list + [a] %}
|
| 60 |
+
{% endif %}
|
| 61 |
+
{% endfor %}
|
| 62 |
+
{% set k = resized.list | length %}
|
| 63 |
+
|
| 64 |
+
const WG: u32 = {{ wg }}u;
|
| 65 |
+
|
| 66 |
+
// Runtime form used by the shared exact-nearest helper. The generic kernel's
|
| 67 |
+
// interpolation paths keep their baked expressions below; nearest uses this
|
| 68 |
+
// function only when an exact inferred-ratio path is unavailable.
|
| 69 |
+
{% if source.mode == "nearest" and k > 0 and (source.roi is not defined or (source.roi | length) == 0) and source.coordinateMode != "align_corners" %}
|
| 70 |
+
|
| 71 |
+
fn source_coord(out_coord: u32, in_size: u32, out_size: u32{% if usesExplicitScale %}, explicit_scale: f32{% endif %}) -> f32 {
|
| 72 |
+
{% if source.coordinateMode not in ["asymmetric", "tf_crop_and_resize"] %}
|
| 73 |
+
let inferred_scale = f32(out_size) / f32(in_size);
|
| 74 |
+
let scale = select(inferred_scale, explicit_scale, explicit_scale > 0.0);
|
| 75 |
+
{% endif %}
|
| 76 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 77 |
+
return select(f32(out_coord) * f32(in_size) / f32(out_size), f32(out_coord) / explicit_scale, explicit_scale > 0.0);
|
| 78 |
+
{% elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 79 |
+
return select(0.0, (f32(out_coord) + 0.5) / scale - 0.5, out_size > 1u);
|
| 80 |
+
{% elif source.coordinateMode == "tf_crop_and_resize" %}
|
| 81 |
+
let input_span = f32(in_size) - 1.0;
|
| 82 |
+
if (out_size <= 1u) { return 0.5 * input_span; }
|
| 83 |
+
return f32(out_coord) * input_span / (f32(out_size) - 1.0);
|
| 84 |
+
{% elif source.coordinateMode == "half_pixel_symmetric" %}
|
| 85 |
+
let output_width = f32(in_size) * scale;
|
| 86 |
+
let adjustment = f32(out_size) / output_width;
|
| 87 |
+
let center = f32(in_size) * 0.5;
|
| 88 |
+
let offset = center * (1.0 - adjustment);
|
| 89 |
+
return offset + (f32(out_coord) + 0.5) / scale - 0.5;
|
| 90 |
+
{% else %}
|
| 91 |
+
return (f32(out_coord) + 0.5) / scale - 0.5;
|
| 92 |
+
{% endif %}
|
| 93 |
+
}
|
| 94 |
+
{%- endif %}
|
| 95 |
+
{% macro istride(a) %}
|
| 96 |
+
{% set s = namespace(v = 1) %}
|
| 97 |
+
{% for j in range(a + 1, r) %}{% set s.v = s.v * xs[j] %}{% endfor %}
|
| 98 |
+
{{ s.v }}
|
| 99 |
+
{%- endmacro %}
|
| 100 |
+
|
| 101 |
+
{% macro out_coord(a) -%}
|
| 102 |
+
{% if ys[a] == 1 -%}0u{% else -%}oc{{ a }}{% endif -%}
|
| 103 |
+
{% endmacro %}{% macro src_coord(a, oc) %}
|
| 104 |
+
{% set inS = xs[a] %}
|
| 105 |
+
{% set outS = ys[a] %}
|
| 106 |
+
{% if source.roi is defined and source.roi %}
|
| 107 |
+
{% set start = staticRoi[a] %}
|
| 108 |
+
{% set end = staticRoi[a + r] %}
|
| 109 |
+
{% if outS > 1 -%}
|
| 110 |
+
({{ start }} * {{ inS - 1 }}.0 + f32({{ oc }}) * (({{ end }}) - ({{ start }})) * {{ inS - 1 }}.0 / {{ outS - 1 }}.0)
|
| 111 |
+
{%- else -%}
|
| 112 |
+
(0.5 * (({{ start }}) + ({{ end }})) * {{ inS - 1 }}.0)
|
| 113 |
+
{%- endif %}
|
| 114 |
+
{% elif source.coordinateMode == "tf_crop_and_resize" %}
|
| 115 |
+
select(
|
| 116 |
+
(0.5 * {{ inS - 1 }}.0),
|
| 117 |
+
(f32({{ oc }}) * {{ inS - 1 }}.0 / f32(max({{ outS }}u - 1u, 1u))),
|
| 118 |
+
{{ outS }}u > 1u)
|
| 119 |
+
{% elif staticScales[a] > 0 %}
|
| 120 |
+
{% set scale = staticScales[a] %}
|
| 121 |
+
{% if source.coordinateMode == "asymmetric" -%}
|
| 122 |
+
(f32({{ oc }}) / {{ scale }})
|
| 123 |
+
{%- elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 124 |
+
{% if outS > 1 %}((f32({{ oc }}) + 0.5) / {{ scale }} - 0.5){% else %}0.0{% endif %}
|
| 125 |
+
{%- elif source.coordinateMode == "align_corners" %}
|
| 126 |
+
{% if outS == 1 %}0.0{% else %}(f32({{ oc }}) * {{ inS - 1 }}.0 / {{ outS - 1 }}.0){% endif %}
|
| 127 |
+
{% elif source.coordinateMode == "half_pixel_symmetric" %}
|
| 128 |
+
{% set adjustment = outS / (inS * scale) %}
|
| 129 |
+
{% set offset = inS * 0.5 * (1.0 - adjustment) %}
|
| 130 |
+
({{ offset }} + (f32({{ oc }}) + 0.5) / {{ scale }} - 0.5)
|
| 131 |
+
{% else -%}
|
| 132 |
+
((f32({{ oc }}) + 0.5) / {{ scale }} - 0.5)
|
| 133 |
+
{%- endif %}
|
| 134 |
+
{% else %}
|
| 135 |
+
{% if source.coordinateMode == "align_corners" %}
|
| 136 |
+
{% if outS == 1 %}0.0{% else %}(f32({{ oc }}) * {{ inS - 1 }}.0 / {{ outS - 1 }}.0){% endif %}
|
| 137 |
+
{% elif source.coordinateMode == "asymmetric" -%}
|
| 138 |
+
(f32({{ oc }}) * {{ inS }}.0 / {{ outS }}.0)
|
| 139 |
+
{%- elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 140 |
+
{% if outS > 1 %}((f32({{ oc }}) + 0.5) * {{ inS }}.0 / {{ outS }}.0 - 0.5){% else %}0.0{% endif %}
|
| 141 |
+
{%- else -%}
|
| 142 |
+
((f32({{ oc }}) + 0.5) * {{ inS }}.0 / {{ outS }}.0 - 0.5)
|
| 143 |
+
{%- endif %}
|
| 144 |
+
{% endif %}
|
| 145 |
+
{% endmacro %}
|
| 146 |
+
|
| 147 |
+
{% set EXCL = source.excludeOutside is defined and source.excludeOutside == 1 %}
|
| 148 |
+
{% if source.mode == "cubic" %}
|
| 149 |
+
// Keys cubic convolution kernel with coefficient coeffA.
|
| 150 |
+
fn cubic_weight(s_in: f32) -> f32 {
|
| 151 |
+
let a = f32({{ source.coeffA }});
|
| 152 |
+
let s = abs(s_in);
|
| 153 |
+
let s2 = s * s;
|
| 154 |
+
let s3 = s2 * s;
|
| 155 |
+
if (s <= 1.0) { return (a + 2.0) * s3 - (a + 3.0) * s2 + 1.0; }
|
| 156 |
+
if (s < 2.0) { return a * s3 - 5.0 * a * s2 + 8.0 * a * s - 4.0 * a; }
|
| 157 |
+
return 0.0;
|
| 158 |
+
}
|
| 159 |
+
{% endif %}
|
| 160 |
+
|
| 161 |
+
{% if source.mode == "nearest" and k > 0 %}
|
| 162 |
+
/* Backend-stable nearest-neighbor coordinate selection. Inferred align-corners,
|
| 163 |
+
* asymmetric, and half-pixel transforms use exact integer ratios. This matters
|
| 164 |
+
* at nearest-mode ties: reassociating an f32 division into a reciprocal
|
| 165 |
+
* multiply can move an exact 0.5 to either side of the boundary. */{% set ratioMode = source.coordinateMode in ["align_corners", "asymmetric"] %}
|
| 166 |
+
{% set inferredHalfPixelMode = source.coordinateMode in ["pytorch_half_pixel", "half_pixel", "half_pixel_symmetric"] %}
|
| 167 |
+
{% set noRoi = source.roi is not defined or (source.roi | length) == 0 %}
|
| 168 |
+
{% set usesExplicitScale = source.coordinateMode not in ["align_corners", "tf_crop_and_resize"] %}
|
| 169 |
+
{% if source.mode == "nearest" %}
|
| 170 |
+
fn nearest_index(v: f32) -> i32 {
|
| 171 |
+
{% if source.nearestMode == "floor" %}
|
| 172 |
+
return i32(floor(v));
|
| 173 |
+
{% elif source.nearestMode == "ceil" %}
|
| 174 |
+
return i32(ceil(v));
|
| 175 |
+
{% elif source.nearestMode == "round_prefer_ceil" %}
|
| 176 |
+
// round half toward +inf. WGSL round() is round-half-to-even and therefore
|
| 177 |
+
// cannot express ONNX's preferred direction. The lower-side epsilon absorbs
|
| 178 |
+
// backend division error that lands a mathematical half-tie just below .5
|
| 179 |
+
// (mirror of the round_prefer_floor epsilon below).
|
| 180 |
+
let tie_epsilon = 0.0000003;
|
| 181 |
+
let f = floor(v);
|
| 182 |
+
let frac = v - f;
|
| 183 |
+
return i32(select(f, f + 1.0, frac >= 0.5 - tie_epsilon));
|
| 184 |
+
{% else %}
|
| 185 |
+
// Explicit scales still require f32 coordinate math. Keep a narrow epsilon
|
| 186 |
+
// for their backend rounding while inferred scales take the exact path below.
|
| 187 |
+
let tie_epsilon = 0.0000003;
|
| 188 |
+
let f = floor(v);
|
| 189 |
+
let frac = v - f;
|
| 190 |
+
return i32(select(f, f + 1.0, frac > 0.5 + tie_epsilon));
|
| 191 |
+
{% endif %}
|
| 192 |
+
}
|
| 193 |
+
|
| 194 |
+
{% if noRoi and (ratioMode or inferredHalfPixelMode) %}
|
| 195 |
+
fn rational_value_nearest_index(numerator: u32, denominator: u32) -> i32 {
|
| 196 |
+
if (numerator == 0u || denominator == 0u) {
|
| 197 |
+
return 0;
|
| 198 |
+
}
|
| 199 |
+
let lower = numerator / denominator;
|
| 200 |
+
{% if source.nearestMode != "floor" %}
|
| 201 |
+
let remainder = numerator - lower * denominator;
|
| 202 |
+
{% endif %}
|
| 203 |
+
{% if source.nearestMode == "floor" %}
|
| 204 |
+
return i32(lower);
|
| 205 |
+
{% elif source.nearestMode == "ceil" %}
|
| 206 |
+
return i32(lower + select(0u, 1u, remainder != 0u));
|
| 207 |
+
{% elif source.nearestMode == "round_prefer_ceil" %}
|
| 208 |
+
return i32(lower + select(0u, 1u, remainder >= denominator - remainder));
|
| 209 |
+
{% else %}
|
| 210 |
+
return i32(lower + select(0u, 1u, remainder > denominator - remainder));
|
| 211 |
+
{% endif %}
|
| 212 |
+
}
|
| 213 |
+
|
| 214 |
+
{% endif %}
|
| 215 |
+
{% if noRoi and ratioMode %}
|
| 216 |
+
fn rational_nearest_index(coord: u32, numerator: u32, denominator: u32) -> i32 {
|
| 217 |
+
if (coord == 0u || numerator == 0u || denominator == 0u) {
|
| 218 |
+
return 0;
|
| 219 |
+
}
|
| 220 |
+
// Use 0xfffffffe rather than UINT_MAX for the overflow guard: some shader
|
| 221 |
+
// compilers reject division with a UINT_MAX constant dividend. The smaller
|
| 222 |
+
// bound is conservative and sends at most one extra coordinate to the f32 path.
|
| 223 |
+
if (coord > 0xfffffffeu / numerator) {
|
| 224 |
+
return nearest_index(f32(coord) * f32(numerator) / f32(denominator));
|
| 225 |
+
}
|
| 226 |
+
return rational_value_nearest_index(coord * numerator, denominator);
|
| 227 |
+
}
|
| 228 |
+
|
| 229 |
+
{% endif %}
|
| 230 |
+
{% if noRoi and inferredHalfPixelMode %}
|
| 231 |
+
fn inferred_half_pixel_nearest_index(
|
| 232 |
+
coord: u32,
|
| 233 |
+
in_size: u32,
|
| 234 |
+
out_size: u32,
|
| 235 |
+
subtract_half: bool
|
| 236 |
+
) -> i32 {
|
| 237 |
+
// (coord + 1/2) * in/out [- 1/2]
|
| 238 |
+
// = ((2*coord + 1)*in [- out]) / (2*out)
|
| 239 |
+
// Use the f32 definition only when the u32 intermediates would overflow.
|
| 240 |
+
if (in_size == 0u || out_size == 0u || out_size > 0x7fffffffu ||
|
| 241 |
+
coord > (0xffffffffu - 1u) / 2u) {
|
| 242 |
+
return nearest_index(source_coord(coord, in_size, out_size, 0.0));
|
| 243 |
+
}
|
| 244 |
+
let odd_coord = coord * 2u + 1u;
|
| 245 |
+
// 0xfffffffe, not 0xffffffff — see rational_nearest_index (NVVM crash).
|
| 246 |
+
if (odd_coord > 0xfffffffeu / in_size) {
|
| 247 |
+
return nearest_index(source_coord(coord, in_size, out_size, 0.0));
|
| 248 |
+
}
|
| 249 |
+
let scaled = odd_coord * in_size;
|
| 250 |
+
let shift = select(0u, out_size, subtract_half);
|
| 251 |
+
// A negative half-pixel coordinate is in [-0.5, 0). Every nearest mode may
|
| 252 |
+
// choose -1 or 0 there, but Resize clamps both to source index zero.
|
| 253 |
+
if (scaled <= shift) {
|
| 254 |
+
return 0;
|
| 255 |
+
}
|
| 256 |
+
return rational_value_nearest_index(scaled - shift, out_size * 2u);
|
| 257 |
+
}
|
| 258 |
+
|
| 259 |
+
{% endif %}
|
| 260 |
+
{% if noRoi %}
|
| 261 |
+
fn nearest_coord_index(
|
| 262 |
+
out_coord: u32,
|
| 263 |
+
in_size: u32,
|
| 264 |
+
out_size: u32{% if usesExplicitScale %},
|
| 265 |
+
explicit_scale: f32
|
| 266 |
+
{% endif %}
|
| 267 |
+
) -> i32 {
|
| 268 |
+
{% if source.coordinateMode == "align_corners" %}
|
| 269 |
+
if (out_size <= 1u || in_size <= 1u) {
|
| 270 |
+
return 0;
|
| 271 |
+
}
|
| 272 |
+
return rational_nearest_index(out_coord, in_size - 1u, out_size - 1u);
|
| 273 |
+
{% else %}
|
| 274 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 275 |
+
if (explicit_scale <= 0.0) {
|
| 276 |
+
return rational_nearest_index(out_coord, in_size, out_size);
|
| 277 |
+
}
|
| 278 |
+
{% elif source.coordinateMode == "pytorch_half_pixel" %}
|
| 279 |
+
if (out_size <= 1u) {
|
| 280 |
+
return 0;
|
| 281 |
+
}
|
| 282 |
+
if (explicit_scale <= 0.0) {
|
| 283 |
+
return inferred_half_pixel_nearest_index(out_coord, in_size, out_size, true);
|
| 284 |
+
}
|
| 285 |
+
{% elif source.coordinateMode == "tf_crop_and_resize" %}
|
| 286 |
+
{% else %}
|
| 287 |
+
if (explicit_scale <= 0.0) {
|
| 288 |
+
return inferred_half_pixel_nearest_index(out_coord, in_size, out_size, true);
|
| 289 |
+
}
|
| 290 |
+
{% endif %}
|
| 291 |
+
return nearest_index(source_coord(out_coord, in_size, out_size{% if usesExplicitScale %}, explicit_scale{% endif %}));
|
| 292 |
+
{% endif %}
|
| 293 |
+
}
|
| 294 |
+
|
| 295 |
+
{% endif %}
|
| 296 |
+
|
| 297 |
+
{% else %}
|
| 298 |
+
|
| 299 |
+
{% endif -%}
|
| 300 |
+
{% endif %}
|
| 301 |
+
|
| 302 |
+
@compute @workgroup_size({{ wg }})
|
| 303 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
|
| 304 |
+
let stride = nwg.x * WG;
|
| 305 |
+
for (var i = gid.x; i < {{ ocount.v }}u; i += stride) {
|
| 306 |
+
// Decompose flat output index into per-axis output coords.
|
| 307 |
+
{% if ocount.v != 1 %}
|
| 308 |
+
var rem = i;
|
| 309 |
+
{% endif %}
|
| 310 |
+
{% for a in range(r) %}
|
| 311 |
+
{% set os = namespace(v = 1) %}
|
| 312 |
+
{% for j in range(a + 1, r) %}{% set os.v = os.v * ys[j] %}{% endfor %}
|
| 313 |
+
{% if ys[a] != 1 %}
|
| 314 |
+
let oc{{ a }} = rem / {{ os.v }}u;
|
| 315 |
+
rem = rem % {{ os.v }}u;
|
| 316 |
+
{% endif %}
|
| 317 |
+
{% endfor %}
|
| 318 |
+
// Per-axis source coordinates.
|
| 319 |
+
{% for a in range(r) %}
|
| 320 |
+
{% if a in resized.list and (source.mode != "nearest" or (source.roi is defined and source.roi)) %}
|
| 321 |
+
let s{{ a }} = {{ src_coord(a, out_coord(a)) }};
|
| 322 |
+
{% endif %}
|
| 323 |
+
{% endfor %}
|
| 324 |
+
{% if source.roi is defined and source.roi %}
|
| 325 |
+
// tf_crop_and_resize extrapolation: any resized axis out of [0, in-1].
|
| 326 |
+
var oob = false;
|
| 327 |
+
{% for a in resized.list %}
|
| 328 |
+
if (s{{ a }} < 0.0 || s{{ a }} > {{ xs[a] - 1 }}.0) { oob = true; }
|
| 329 |
+
{% endfor %}
|
| 330 |
+
if (oob) { y[i] = {{ scalar }}(f32({{ source.extrapolationValue }})); continue; }
|
| 331 |
+
{% endif %}
|
| 332 |
+
{% if source.mode == "nearest" %}
|
| 333 |
+
{% for a in range(r) %}
|
| 334 |
+
{% if a in resized.list %}
|
| 335 |
+
{% if source.roi is defined and source.roi %}
|
| 336 |
+
let idx{{ a }} = u32(clamp(nearest_index(s{{ a }}), 0, {{ xs[a] - 1 }}));
|
| 337 |
+
{% else %}
|
| 338 |
+
{% if staticScales[a] > 0 %}
|
| 339 |
+
{% set explicitScale = staticScales[a] %}
|
| 340 |
+
{% else %}
|
| 341 |
+
{% set explicitScale = "0.0" %}
|
| 342 |
+
{% endif %}
|
| 343 |
+
let idx{{ a }} = u32(clamp(nearest_coord_index({{ out_coord(a) }}, {{ xs[a] }}u, {{ ys[a] }}u{% if usesExplicitScale %}, {{ explicitScale }}{% endif %}), 0, {{ xs[a] - 1 }}));
|
| 344 |
+
{% endif %}
|
| 345 |
+
{% else %}
|
| 346 |
+
let idx{{ a }} = {{ out_coord(a) }};
|
| 347 |
+
{% endif %}
|
| 348 |
+
{% endfor %}
|
| 349 |
+
y[i] = {{ scalar }}(f32(x[0u{% for a in range(r) %} + idx{{ a }} * {{ istride(a) }}u{% endfor %}]));
|
| 350 |
+
{% elif source.mode == "cubic" %}
|
| 351 |
+
// Separable Keys cubic: 4 taps per resized axis (offsets -1..2), passthrough
|
| 352 |
+
// axes contribute a single tap. 4^k taps total.
|
| 353 |
+
{% for a in resized.list %}
|
| 354 |
+
let lo{{ a }} = i32(floor(s{{ a }}));
|
| 355 |
+
let t{{ a }} = s{{ a }} - f32(lo{{ a }});
|
| 356 |
+
var cwv{{ a }} = array<f32, 4>(cubic_weight(1.0 + t{{ a }}), cubic_weight(t{{ a }}), cubic_weight(1.0 - t{{ a }}), cubic_weight(2.0 - t{{ a }}));
|
| 357 |
+
{% if EXCL %}
|
| 358 |
+
// exclude_outside: zero taps outside the input range, renormalize per axis.
|
| 359 |
+
for (var q = 0; q < 4; q++) {
|
| 360 |
+
if (lo{{ a }} + q - 1 < 0 || lo{{ a }} + q - 1 > {{ xs[a] - 1 }}) { cwv{{ a }}[q] = 0.0; }
|
| 361 |
+
}
|
| 362 |
+
let cws{{ a }} = cwv{{ a }}[0] + cwv{{ a }}[1] + cwv{{ a }}[2] + cwv{{ a }}[3];
|
| 363 |
+
for (var q = 0; q < 4; q++) { cwv{{ a }}[q] = cwv{{ a }}[q] / cws{{ a }}; }
|
| 364 |
+
{% endif %}
|
| 365 |
+
{% endfor %}
|
| 366 |
+
var acc = 0.0;
|
| 367 |
+
{% for combo in range(pow(4, k)) %}
|
| 368 |
+
{
|
| 369 |
+
var cw = 1.0;
|
| 370 |
+
{% for ridx in range(k) %}
|
| 371 |
+
{% set a = resized.list[ridx] %}
|
| 372 |
+
{% set tap = ((combo / pow(4, ridx)) | int) % 4 %}
|
| 373 |
+
let ci{{ a }} = clamp(lo{{ a }} + {{ tap - 1 }}, 0, {{ xs[a] - 1 }});
|
| 374 |
+
cw = cw * cwv{{ a }}[{{ tap }}];
|
| 375 |
+
{% endfor %}
|
| 376 |
+
let off = 0u{% for a in range(r) %}{% if a in resized.list %} + u32(ci{{ a }}) * {{ istride(a) }}u{% else %} + {{ out_coord(a) }} * {{ istride(a) }}u{% endif %}{% endfor %};
|
| 377 |
+
acc = acc + cw * f32(x[off]);
|
| 378 |
+
}
|
| 379 |
+
{% endfor %}
|
| 380 |
+
y[i] = {{ scalar }}(acc);
|
| 381 |
+
{% else %}
|
| 382 |
+
// Linear: interpolate over the resized axes (2^k corners); passthrough axes
|
| 383 |
+
// contribute their fixed index with weight 1.
|
| 384 |
+
{% for a in resized.list %}
|
| 385 |
+
let lo{{ a }} = i32(floor(s{{ a }}));
|
| 386 |
+
let w{{ a }} = s{{ a }} - f32(lo{{ a }});
|
| 387 |
+
{% endfor %}
|
| 388 |
+
var acc = 0.0;
|
| 389 |
+
{% for corner in range(pow(2, k)) %}
|
| 390 |
+
{
|
| 391 |
+
var cw = 1.0;
|
| 392 |
+
{% for ridx in range(k) %}
|
| 393 |
+
{% set a = resized.list[ridx] %}
|
| 394 |
+
{% if ((corner / pow(2, ridx)) | int) % 2 == 1 %}
|
| 395 |
+
let ci{{ a }} = clamp(lo{{ a }} + 1, 0, {{ xs[a] - 1 }});
|
| 396 |
+
cw = cw * w{{ a }};
|
| 397 |
+
{% else %}
|
| 398 |
+
let ci{{ a }} = clamp(lo{{ a }}, 0, {{ xs[a] - 1 }});
|
| 399 |
+
cw = cw * (1.0 - w{{ a }});
|
| 400 |
+
{% endif %}
|
| 401 |
+
{% endfor %}
|
| 402 |
+
let off = 0u{% for a in range(r) %}{% if a in resized.list %} + u32(ci{{ a }}) * {{ istride(a) }}u{% else %} + {{ out_coord(a) }} * {{ istride(a) }}u{% endif %}{% endfor %};
|
| 403 |
+
acc = acc + cw * f32(x[off]);
|
| 404 |
+
}
|
| 405 |
+
{% endfor %}
|
| 406 |
+
y[i] = {{ scalar }}(acc);
|
| 407 |
+
{% endif %}
|
| 408 |
+
}
|
| 409 |
+
}
|
build/webgpu/resize-linear-2x-stencil-x8.wgsl.jinja
ADDED
|
@@ -0,0 +1,81 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 3 |
+
|
| 4 |
+
// Exact-2x f16 bilinear resize with two adjacent vec4 outputs per invocation.
|
| 5 |
+
// Neighboring output quads overlap by two source columns, so an x8 tile needs
|
| 6 |
+
// only six horizontal taps per source row instead of two independent x4
|
| 7 |
+
// stencils' eight. Coordinate/parity work and thread count are also halved.
|
| 8 |
+
const WG: u32 = {{ workgroupSize }}u;
|
| 9 |
+
const IN_H: u32 = {{ source.xShape[2] }}u;
|
| 10 |
+
const IN_W: u32 = {{ source.xShape[3] }}u;
|
| 11 |
+
const IN_HW: u32 = {{ source.xShape[2] * source.xShape[3] }}u;
|
| 12 |
+
const OUT_H: u32 = {{ source.yShape[2] }}u;
|
| 13 |
+
const OUT_W: u32 = {{ source.yShape[3] }}u;
|
| 14 |
+
const COUNT: u32 = {{ source.yShape[0] * source.yShape[1] * source.yShape[2] * source.yShape[3] }}u;
|
| 15 |
+
|
| 16 |
+
fn input_at(base: u32, h: i32, w: i32) -> f32 {
|
| 17 |
+
let hh = u32(clamp(h, 0, i32(IN_H) - 1));
|
| 18 |
+
let ww = u32(clamp(w, 0, i32(IN_W) - 1));
|
| 19 |
+
return f32(x[base + hh * IN_W + ww]);
|
| 20 |
+
}
|
| 21 |
+
|
| 22 |
+
@compute @workgroup_size(WG, 1, 1)
|
| 23 |
+
fn main(
|
| 24 |
+
@builtin(global_invocation_id) gid: vec3<u32>,
|
| 25 |
+
@builtin(num_workgroups) workgroup_count: vec3<u32>
|
| 26 |
+
) {
|
| 27 |
+
let stride = workgroup_count.x * WG;
|
| 28 |
+
let groups_per_row = OUT_W / 8u;
|
| 29 |
+
let group_count = COUNT / 8u;
|
| 30 |
+
for (var group = gid.x; group < group_count; group = group + stride) {
|
| 31 |
+
let group_in_row = group % groups_per_row;
|
| 32 |
+
let row = group / groups_per_row;
|
| 33 |
+
let oh = row % OUT_H;
|
| 34 |
+
let nc = row / OUT_H;
|
| 35 |
+
let base = nc * IN_HW;
|
| 36 |
+
let qx = i32(group_in_row * 4u);
|
| 37 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 38 |
+
let y0 = i32(oh / 2u);
|
| 39 |
+
let wy = select(0.0, 0.5, (oh & 1u) == 1u);
|
| 40 |
+
let a0 = input_at(base, y0, qx + 0);
|
| 41 |
+
let b0 = input_at(base, y0, qx + 1);
|
| 42 |
+
let c0 = input_at(base, y0, qx + 2);
|
| 43 |
+
let d0 = input_at(base, y0, qx + 3);
|
| 44 |
+
let e0 = input_at(base, y0, qx + 4);
|
| 45 |
+
let y1 = y0 + 1;
|
| 46 |
+
let a1 = input_at(base, y1, qx + 0);
|
| 47 |
+
let b1 = input_at(base, y1, qx + 1);
|
| 48 |
+
let c1 = input_at(base, y1, qx + 2);
|
| 49 |
+
let d1 = input_at(base, y1, qx + 3);
|
| 50 |
+
let e1 = input_at(base, y1, qx + 4);
|
| 51 |
+
let top0 = vec4<f32>(a0, 0.5 * a0 + 0.5 * b0, b0, 0.5 * b0 + 0.5 * c0);
|
| 52 |
+
let top1 = vec4<f32>(c0, 0.5 * c0 + 0.5 * d0, d0, 0.5 * d0 + 0.5 * e0);
|
| 53 |
+
let bottom0 = vec4<f32>(a1, 0.5 * a1 + 0.5 * b1, b1, 0.5 * b1 + 0.5 * c1);
|
| 54 |
+
let bottom1 = vec4<f32>(c1, 0.5 * c1 + 0.5 * d1, d1, 0.5 * d1 + 0.5 * e1);
|
| 55 |
+
{% else %}
|
| 56 |
+
let y0 = i32(oh / 2u) - 1 + i32(oh & 1u);
|
| 57 |
+
let wy = select(0.75, 0.25, (oh & 1u) == 1u);
|
| 58 |
+
let a0 = input_at(base, y0, qx - 1);
|
| 59 |
+
let b0 = input_at(base, y0, qx + 0);
|
| 60 |
+
let c0 = input_at(base, y0, qx + 1);
|
| 61 |
+
let d0 = input_at(base, y0, qx + 2);
|
| 62 |
+
let e0 = input_at(base, y0, qx + 3);
|
| 63 |
+
let f0 = input_at(base, y0, qx + 4);
|
| 64 |
+
let y1 = y0 + 1;
|
| 65 |
+
let a1 = input_at(base, y1, qx - 1);
|
| 66 |
+
let b1 = input_at(base, y1, qx + 0);
|
| 67 |
+
let c1 = input_at(base, y1, qx + 1);
|
| 68 |
+
let d1 = input_at(base, y1, qx + 2);
|
| 69 |
+
let e1 = input_at(base, y1, qx + 3);
|
| 70 |
+
let f1 = input_at(base, y1, qx + 4);
|
| 71 |
+
let top0 = vec4<f32>(0.25 * a0 + 0.75 * b0, 0.75 * b0 + 0.25 * c0, 0.25 * b0 + 0.75 * c0, 0.75 * c0 + 0.25 * d0);
|
| 72 |
+
let top1 = vec4<f32>(0.25 * c0 + 0.75 * d0, 0.75 * d0 + 0.25 * e0, 0.25 * d0 + 0.75 * e0, 0.75 * e0 + 0.25 * f0);
|
| 73 |
+
let bottom0 = vec4<f32>(0.25 * a1 + 0.75 * b1, 0.75 * b1 + 0.25 * c1, 0.25 * b1 + 0.75 * c1, 0.75 * c1 + 0.25 * d1);
|
| 74 |
+
let bottom1 = vec4<f32>(0.25 * c1 + 0.75 * d1, 0.75 * d1 + 0.25 * e1, 0.25 * d1 + 0.75 * e1, 0.75 * e1 + 0.25 * f1);
|
| 75 |
+
{% endif %}
|
| 76 |
+
let output0 = vec4<f32>(1.0 - wy) * top0 + vec4<f32>(wy) * bottom0;
|
| 77 |
+
let output1 = vec4<f32>(1.0 - wy) * top1 + vec4<f32>(wy) * bottom1;
|
| 78 |
+
y[group * 2u] = vec4<f16>(output0);
|
| 79 |
+
y[group * 2u + 1u] = vec4<f16>(output1);
|
| 80 |
+
}
|
| 81 |
+
}
|
build/webgpu/resize-linear-2x-stencil.wgsl.jinja
ADDED
|
@@ -0,0 +1,124 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// NCHW bilinear Resize with an exact 2x upscale on both spatial axes. The
|
| 2 |
+
// interpolation weights collapse to fixed per-parity stencils, replacing
|
| 3 |
+
// per-element floating-point coordinate divisions with integer parity tests
|
| 4 |
+
// and constants.
|
| 5 |
+
//
|
| 6 |
+
// For half_pixel and pytorch_half_pixel, src = o / 2 - 0.25: even outputs use
|
| 7 |
+
// floor q - 1 with fraction 0.75, and odd outputs use floor q with fraction
|
| 8 |
+
// 0.25. For asymmetric mode, src = o / 2: even outputs use fraction 0.0 and odd
|
| 9 |
+
// outputs use 0.5. align_corners remains on the generic fallback because its
|
| 10 |
+
// position-dependent scale does not reduce to fixed stencils.
|
| 11 |
+
//
|
| 12 |
+
// Input coordinates clamp at the edges, including the q - 1 and q + 1 border
|
| 13 |
+
// taps. The f32 weight construction is bit-identical to the generic fallback.
|
| 14 |
+
{% if usesF16 %}
|
| 15 |
+
enable f16;
|
| 16 |
+
{% endif %}
|
| 17 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 18 |
+
{% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
|
| 19 |
+
|
| 20 |
+
const WG: u32 = {{ wg }}u;
|
| 21 |
+
const IN_H: u32 = {{ source.xShape[2] }}u;
|
| 22 |
+
const IN_W: u32 = {{ source.xShape[3] }}u;
|
| 23 |
+
const IN_HW: u32 = {{ source.xShape[2] * source.xShape[3] }}u;
|
| 24 |
+
const OUT_H: u32 = {{ source.yShape[2] }}u;
|
| 25 |
+
const OUT_W: u32 = {{ source.yShape[3] }}u;
|
| 26 |
+
const COUNT: u32 = {{ source.yShape[0] * source.yShape[1] * source.yShape[2] * source.yShape[3] }}u;
|
| 27 |
+
|
| 28 |
+
fn input_at(base: u32, h: i32, w: i32) -> f32 {
|
| 29 |
+
let hh = u32(clamp(h, 0, i32(IN_H) - 1));
|
| 30 |
+
let ww = u32(clamp(w, 0, i32(IN_W) - 1));
|
| 31 |
+
return f32(x[base + hh * IN_W + ww]);
|
| 32 |
+
}
|
| 33 |
+
|
| 34 |
+
@compute @workgroup_size({{ wg }})
|
| 35 |
+
fn main(
|
| 36 |
+
@builtin(global_invocation_id) gid: vec3<u32>,
|
| 37 |
+
@builtin(num_workgroups) nwg: vec3<u32>
|
| 38 |
+
) {
|
| 39 |
+
let stride = nwg.x * WG;
|
| 40 |
+
{% if source.vec4 %}
|
| 41 |
+
// Vec4 path: one thread owns 4 contiguous output columns (OUT_W % 4 == 0, so
|
| 42 |
+
// a 4-group never crosses a row). All 4 lanes share the same oh -> same source
|
| 43 |
+
// rows (y0,y1) and the same 4-column source window (qx0-1 .. qx0+2); load that
|
| 44 |
+
// window once, compute the 4 bilinear results with the baked per-lane parity
|
| 45 |
+
// weights, and emit a single 128-bit vec4 store. Bit-identical to the scalar
|
| 46 |
+
// path (same input_at, same (1-w) split, same f32 operation order).
|
| 47 |
+
let OUT_W4 = OUT_W / 4u;
|
| 48 |
+
let COUNT4 = COUNT / 4u;
|
| 49 |
+
for (var v = gid.x; v < COUNT4; v += stride) {
|
| 50 |
+
let ow4 = v % OUT_W4;
|
| 51 |
+
let tt = v / OUT_W4;
|
| 52 |
+
let oh = tt % OUT_H;
|
| 53 |
+
let nc = tt / OUT_H;
|
| 54 |
+
let base = nc * IN_HW;
|
| 55 |
+
let ow0 = ow4 * 4u;
|
| 56 |
+
let qx0 = i32(ow0 / 2u);
|
| 57 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 58 |
+
let y0 = i32(oh / 2u);
|
| 59 |
+
let wy = select(0.0, 0.5, (oh & 1u) == 1u);
|
| 60 |
+
{% else %}
|
| 61 |
+
let y0 = i32(oh / 2u) - 1 + i32(oh & 1u);
|
| 62 |
+
let wy = select(0.75, 0.25, (oh & 1u) == 1u);
|
| 63 |
+
{% endif %}
|
| 64 |
+
let y1 = y0 + 1;
|
| 65 |
+
{% if source.coordinateMode != "asymmetric" %}
|
| 66 |
+
// The q-1 left tap is only weighted by the half-pixel stencils; the
|
| 67 |
+
// asymmetric stencils (weights 0.0/0.5) never read it.
|
| 68 |
+
let a0 = input_at(base, y0, qx0 - 1);
|
| 69 |
+
let a1 = input_at(base, y1, qx0 - 1);
|
| 70 |
+
{% endif %}
|
| 71 |
+
let b0 = input_at(base, y0, qx0);
|
| 72 |
+
let c0 = input_at(base, y0, qx0 + 1);
|
| 73 |
+
let d0 = input_at(base, y0, qx0 + 2);
|
| 74 |
+
let b1 = input_at(base, y1, qx0);
|
| 75 |
+
let c1 = input_at(base, y1, qx0 + 1);
|
| 76 |
+
let d1 = input_at(base, y1, qx0 + 2);
|
| 77 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 78 |
+
// lanes (ow0..ow0+3): wx 0.0,0.5,0.0,0.5 ; (v00,v01) cols (b,c),(b,c),(c,d),(c,d)
|
| 79 |
+
let o0 = (1.0 - wy) * (1.0 * b0 + 0.0 * c0) + wy * (1.0 * b1 + 0.0 * c1);
|
| 80 |
+
let o1 = (1.0 - wy) * (0.5 * b0 + 0.5 * c0) + wy * (0.5 * b1 + 0.5 * c1);
|
| 81 |
+
let o2 = (1.0 - wy) * (1.0 * c0 + 0.0 * d0) + wy * (1.0 * c1 + 0.0 * d1);
|
| 82 |
+
let o3 = (1.0 - wy) * (0.5 * c0 + 0.5 * d0) + wy * (0.5 * c1 + 0.5 * d1);
|
| 83 |
+
{% else %}
|
| 84 |
+
// lanes (ow0..ow0+3): wx 0.75,0.25,0.75,0.25 ; (v00,v01) cols (a,b),(b,c),(b,c),(c,d)
|
| 85 |
+
let o0 = (1.0 - wy) * (0.25 * a0 + 0.75 * b0) + wy * (0.25 * a1 + 0.75 * b1);
|
| 86 |
+
let o1 = (1.0 - wy) * (0.75 * b0 + 0.25 * c0) + wy * (0.75 * b1 + 0.25 * c1);
|
| 87 |
+
let o2 = (1.0 - wy) * (0.25 * b0 + 0.75 * c0) + wy * (0.25 * b1 + 0.75 * c1);
|
| 88 |
+
let o3 = (1.0 - wy) * (0.75 * c0 + 0.25 * d0) + wy * (0.75 * c1 + 0.25 * d1);
|
| 89 |
+
{% endif %}
|
| 90 |
+
y[v] = vec4<{{ scalar }}>(vec4<f32>(o0, o1, o2, o3));
|
| 91 |
+
}
|
| 92 |
+
{% else %}
|
| 93 |
+
for (var i = gid.x; i < COUNT; i += stride) {
|
| 94 |
+
let ow = i % OUT_W;
|
| 95 |
+
let t = i / OUT_W;
|
| 96 |
+
let oh = t % OUT_H;
|
| 97 |
+
let nc = t / OUT_H;
|
| 98 |
+
let base = nc * IN_HW;
|
| 99 |
+
let qy = i32(oh / 2u);
|
| 100 |
+
let qx = i32(ow / 2u);
|
| 101 |
+
let oddY = (oh & 1u) == 1u;
|
| 102 |
+
let oddX = (ow & 1u) == 1u;
|
| 103 |
+
{% if source.coordinateMode == "asymmetric" %}
|
| 104 |
+
let y0 = qy;
|
| 105 |
+
let x0 = qx;
|
| 106 |
+
let wy = select(0.0, 0.5, oddY);
|
| 107 |
+
let wx = select(0.0, 0.5, oddX);
|
| 108 |
+
{% else %}
|
| 109 |
+
let y0 = qy - 1 + i32(oh & 1u);
|
| 110 |
+
let x0 = qx - 1 + i32(ow & 1u);
|
| 111 |
+
let wy = select(0.75, 0.25, oddY);
|
| 112 |
+
let wx = select(0.75, 0.25, oddX);
|
| 113 |
+
{% endif %}
|
| 114 |
+
let y1 = y0 + 1;
|
| 115 |
+
let x1 = x0 + 1;
|
| 116 |
+
let v00 = input_at(base, y0, x0);
|
| 117 |
+
let v01 = input_at(base, y0, x1);
|
| 118 |
+
let v10 = input_at(base, y1, x0);
|
| 119 |
+
let v11 = input_at(base, y1, x1);
|
| 120 |
+
let result = (1.0 - wy) * ((1.0 - wx) * v00 + wx * v01) + wy * ((1.0 - wx) * v10 + wx * v11);
|
| 121 |
+
y[i] = {{ scalar }}(result);
|
| 122 |
+
}
|
| 123 |
+
{% endif %}
|
| 124 |
+
}
|
build/webgpu/resize-nearest-integer-scale.wgsl.jinja
ADDED
|
@@ -0,0 +1,73 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// NCHW nearest-neighbor resize with compile-time integer upscale factors. The
|
| 2 |
+
// source index is exactly out_coord / scale using u32 division, with no floating
|
| 3 |
+
// point coordinate calculation.
|
| 4 |
+
//
|
| 5 |
+
// This is valid for asymmetric + floor because floor(ow / s) is u32 division.
|
| 6 |
+
// It is also valid for half_pixel with either round-prefer rule: writing
|
| 7 |
+
// ow = q*s + r gives an offset whose magnitude is strictly below one half, so
|
| 8 |
+
// ties cannot occur and every nearest rounding rule returns q. Downsampling and
|
| 9 |
+
// other coordinate-mode combinations require the generic coordinate transform.
|
| 10 |
+
{% if usesF16 %}
|
| 11 |
+
enable f16;
|
| 12 |
+
{% endif %}
|
| 13 |
+
{{ env.wgsl.resourceDeclarations }}
|
| 14 |
+
{% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
|
| 15 |
+
{%- set inH = source.xShape[2] %}
|
| 16 |
+
{% set inW = source.xShape[3] %}
|
| 17 |
+
{% set outH = source.yShape[2] %}
|
| 18 |
+
{% set outW = source.yShape[3] %}
|
| 19 |
+
{% set scaleW = ((outW / inW) | int) %}
|
| 20 |
+
{% set rows = source.yShape[0] * source.yShape[1] * outH %}
|
| 21 |
+
const WG: u32 = {{ wg }}u;
|
| 22 |
+
const IN_W: u32 = {{ inW }}u;
|
| 23 |
+
const IN_HW: u32 = {{ inH * inW }}u;
|
| 24 |
+
const OUT_H: u32 = {{ outH }}u;
|
| 25 |
+
const SCALE_H: u32 = {{ ((outH / inH) | int) }}u;
|
| 26 |
+
{% if not source.vec4 or scaleW == 3 %}const SCALE_W: u32 = {{ scaleW }}u;
|
| 27 |
+
{% endif %}
|
| 28 |
+
{% if source.vec4 %}
|
| 29 |
+
const OUT_W_VECS: u32 = {{ ((outW / 4) | int) }}u;
|
| 30 |
+
const COUNT: u32 = {{ (((rows * outW) / 4) | int) }}u;
|
| 31 |
+
{% else %}
|
| 32 |
+
const OUT_W: u32 = {{ outW }}u;
|
| 33 |
+
const COUNT: u32 = {{ rows * outW }}u;
|
| 34 |
+
{% endif %}
|
| 35 |
+
|
| 36 |
+
@compute @workgroup_size({{ wg }})
|
| 37 |
+
fn main(
|
| 38 |
+
@builtin(global_invocation_id) gid: vec3<u32>,
|
| 39 |
+
@builtin(num_workgroups) nwg: vec3<u32>
|
| 40 |
+
) {
|
| 41 |
+
let stride = nwg.x * WG;
|
| 42 |
+
{% if source.vec4 %}
|
| 43 |
+
for (var v = gid.x; v < COUNT; v += stride) {
|
| 44 |
+
let ow0 = (v % OUT_W_VECS) * 4u;
|
| 45 |
+
let t = v / OUT_W_VECS;
|
| 46 |
+
let oh = t % OUT_H;
|
| 47 |
+
let nc = t / OUT_H;
|
| 48 |
+
let rowBase = nc * IN_HW + (oh / SCALE_H) * IN_W;
|
| 49 |
+
{% if scaleW == 4 %}
|
| 50 |
+
y[v] = {{ vectorScalar }}(x[rowBase + ow0 / 4u]);
|
| 51 |
+
{% elif scaleW == 2 %}
|
| 52 |
+
let a = x[rowBase + ow0 / 2u];
|
| 53 |
+
let b = x[rowBase + ow0 / 2u + 1u];
|
| 54 |
+
y[v] = {{ vectorScalar }}(a, a, b, b);
|
| 55 |
+
{% else %}
|
| 56 |
+
y[v] = {{ vectorScalar }}(
|
| 57 |
+
x[rowBase + ow0 / SCALE_W],
|
| 58 |
+
x[rowBase + (ow0 + 1u) / SCALE_W],
|
| 59 |
+
x[rowBase + (ow0 + 2u) / SCALE_W],
|
| 60 |
+
x[rowBase + (ow0 + 3u) / SCALE_W]
|
| 61 |
+
);
|
| 62 |
+
{% endif %}
|
| 63 |
+
}
|
| 64 |
+
{% else %}
|
| 65 |
+
for (var i = gid.x; i < COUNT; i += stride) {
|
| 66 |
+
let ow = i % OUT_W;
|
| 67 |
+
let t = i / OUT_W;
|
| 68 |
+
let oh = t % OUT_H;
|
| 69 |
+
let nc = t / OUT_H;
|
| 70 |
+
y[i] = x[nc * IN_HW + (oh / SCALE_H) * IN_W + ow / SCALE_W];
|
| 71 |
+
}
|
| 72 |
+
{% endif %}
|
| 73 |
+
}
|
build/webgpu/test.json
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|