sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +17 -13
build/webgpu/bench.json +1 -2
build/webgpu/deform-conv2d-aligned-groups.wgsl.jinja +38 -38
build/webgpu/deform-conv2d-nchw.wgsl.jinja +15 -15
build/webgpu/manifest.json +149 -400
build/webgpu/metadata.json +18 -8
build/webgpu/test.json +4 -5

README.md CHANGED Viewed

@@ -18,19 +18,19 @@ See the [ONNX `DeformConv` spec](https://onnx.ai/onnx/operators/onnx__DeformConv
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `X` | `x` | `T` | `4` | — | Input data tensor of shape `(N, C, H, W)` for 2D or `(N, C, D1, ..., Dn)` in general. | required |
-| `W` | `w` | `T` | `4` | — | Convolution weight tensor of shape `(oC, C/group, kH, kW)`. | required |
-| `offset` | `offset` | `T` | `4` | — | Per-output-position sampling offsets of shape `(N, offset_group * kH * kW * 2, oH, oW)` for 2D data. | required |
-| `B` | `bias` | `T` | `1` | — | Optional 1D bias of length oC added to the convolution output. | optional |
-| `mask` | `mask` | `T` | `4` | — | Optional modulation mask of shape `(N, offset_group * kH * kW, oH, oW)` scaling each sampled kernel point; defaults to ones. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `T` | `4` | derived; see description | Output tensor of shape `(N, oC, oH, oW)` containing the deformable convolution result. | required |
 ## Attributes
@@ -38,10 +38,10 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
-| `group` | `1` | Number of groups the input channels C and output channels oC are divided into; both must be divisible by group. |
-| `offset_group` | `1` | Number of offset groups; input channels `C` must be divisible by `offset_group`. |
 | `dilations` | — | Dilation factors for the spatial axes in height-width order; defaults to `[1, 1]`. |
 | `kernel_shape` | — | Kernel dimensions in height-width order. When omitted, they are inferred from the spatial dimensions of `W`. |
 | `pads` | — | Padding at the beginning and end of each spatial axis in `[top, left, bottom, right]` order; defaults to zeros. |
 | `strides` | — | Stride factors for the spatial axes in height-width order; defaults to `[1, 1]`. |
@@ -53,7 +53,7 @@ Attributes and default values (overridable per request):
 ## Files
-- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
@@ -62,10 +62,14 @@ Attributes and default values (overridable per request):
 ## Use with `@huggingface/kernels`
-The loader derives every required output's shape and logical dtype from the manifest contract and this call.
-It then allocates the result tensors automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `x` | `X` | `T` | `4` | — | Input data tensor of shape `(N, C, H, W)` for 2D or `(N, C, D1, ..., Dn)` in general. | required |
+| `w` | `W` | `T` | `4` | — | Convolution weight tensor of shape `(oC, C/group, kH, kW)`. | required |
+| `offset` | — | `T` | `4` | — | Per-output-position sampling offsets of shape `(N, offset_group * kH * kW * 2, oH, oW)` for 2D data. | required |
+| `bias` | `B` | `T` | `1` | — | Optional 1D bias of length oC added to the convolution output. | optional |
+| `mask` | — | `T` | `4` | — | Optional modulation mask of shape `(N, offset_group * kH * kW, oH, oW)` scaling each sampled kernel point; defaults to ones. | optional |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `T` | `4` | derived | Output tensor of shape `(N, oC, oH, oW)` containing the deformable convolution result. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `dilations` | — | Dilation factors for the spatial axes in height-width order; defaults to `[1, 1]`. |
+| `group` | `1` | Number of groups the input channels C and output channels oC are divided into; both must be divisible by group. |
 | `kernel_shape` | — | Kernel dimensions in height-width order. When omitted, they are inferred from the spatial dimensions of `W`. |
+| `offset_group` | `1` | Number of offset groups; input channels `C` must be divisible by `offset_group`. |
 | `pads` | — | Padding at the beginning and end of each spatial axis in `[top, left, bottom, right]` order; defaults to zeros. |
 | `strides` | — | Stride factors for the spatial axes in height-width order; defaults to `[1, 1]`. |
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.DeformConv",
   "cases": [
     {
       "name": "nchw_1x8x32x32_k3",
@@ -75,7 +74,7 @@
       "name": "deform-conv-f32-unmodulated-b1c64m64-128x128-k3-g4-pathology",
       "preset": "stress",
       "provenance": {
-        "source": "authored for variant coverage",
         "notes": "Unmodulated grouped DeformConv exercises the nchw2d route; the paired bias-and-mask case selects the blocked kernel and isolates the optional-input route boundary."
       },
       "attrs": { "group": 4, "offset_group": 4, "pads": [1, 1, 1, 1] },

 {
   "cases": [
     {
       "name": "nchw_1x8x32x32_k3",
       "name": "deform-conv-f32-unmodulated-b1c64m64-128x128-k3-g4-pathology",
       "preset": "stress",
       "provenance": {
+        "source": "synthetic benchmark",
         "notes": "Unmodulated grouped DeformConv exercises the nchw2d route; the paired bias-and-mask case selects the blocked kernel and isolates the optional-input route boundary."
       },
       "attrs": { "group": 4, "offset_group": 4, "pads": [1, 1, 1, 1] },

build/webgpu/deform-conv2d-aligned-groups.wgsl.jinja CHANGED Viewed

@@ -4,26 +4,26 @@
 // output group shares the same tap coordinates and modulation mask. Hoist the
 // offset fetch, floor, bounds and interpolation coefficients out of the channel
 // loop instead of repeating them C/group times for every output channel.
-const COUNT: u32 = {{ source.count }}u;
-const IN_C: u32 = {{ source.inC }}u;
-const IN_CPG: u32 = {{ source.inCPerGroup }}u;
-const IN_H: i32 = {{ source.inH }};
-const IN_W: i32 = {{ source.inW }};
-const IN_H_U: u32 = {{ source.inH }}u;
-const IN_W_U: u32 = {{ source.inW }}u;
-const OUT_C: u32 = {{ source.outC }}u;
-const OUT_CPG: u32 = {{ source.outCPerGroup }}u;
-{% if source.outChannelBlock != 1 %}const OC_BLOCK: u32 = {{ source.outChannelBlock }}u;
 {% endif %}
-const OUT_H: u32 = {{ source.outH }}u;
-const OUT_W: u32 = {{ source.outW }}u;
-const KAREA: u32 = {{ source.kernelH * source.kernelW }}u;
-{% if source.outChannelBlock != 1 %}const W_OC_STRIDE: u32 = IN_CPG * KAREA;
 {% endif %}
-const STRIDE_H: u32 = {{ source.strideH }}u;
-const STRIDE_W: u32 = {{ source.strideW }}u;
-const PAD_H: i32 = {{ source.padH }};
-const PAD_W: i32 = {{ source.padW }};
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 fn pixel(n: u32, c: u32, h: i32, wc: i32) -> f32 {
@@ -32,10 +32,10 @@ fn pixel(n: u32, c: u32, h: i32, wc: i32) -> f32 {
 }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
-  let index = gid.x + gid.y * nwg.x * WG;
   if (index >= COUNT) { return; }
-{% if source.outChannelBlock != 1 %}
   // Output channels in the same convolution group use identical deformable
   // coordinates, masks, and sampled input values. Accumulate a compile-time
   // output-channel block together so those values are shared across the block.
@@ -45,11 +45,11 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   t /= OUT_H;
   let ocBlock = t % (OUT_CPG / OC_BLOCK);
   t /= OUT_CPG / OC_BLOCK;
-  let group = t % {{ source.groups }}u;
-  let n = t / {{ source.groups }}u;
   let oc = group * OUT_CPG + ocBlock * OC_BLOCK;
-{% for v in range(source.outChannelVecs) %}
-{% if source.hasBias %}
   var acc{{ v }} = vec4<f32>(
     bias[oc + {{ v * 4 }}u], bias[oc + {{ v * 4 + 1 }}u],
     bias[oc + {{ v * 4 + 2 }}u], bias[oc + {{ v * 4 + 3 }}u]);
@@ -65,7 +65,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let oc = t % OUT_C;
   let n = t / OUT_C;
   let group = oc / OUT_CPG;
-{% if source.hasBias %}
   var acc = bias[oc];
 {% else %}
   var acc = 0.0;
@@ -75,24 +75,24 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let spatial = oh * OUT_W + ow;
   let offsetPlane = OUT_H * OUT_W;
-{% for kh in range(source.kernelH) %}
-{% for kw in range(source.kernelW) %}
   {
-    const TAP: u32 = {{ kh * source.kernelW + kw }}u;
     let offsetChannel = group * KAREA * 2u + TAP * 2u;
-    let offsetBase = n * ({{ source.groups }}u * KAREA * 2u) * offsetPlane;
     let dy = offset[offsetBase + offsetChannel * offsetPlane + spatial];
     let dx = offset[offsetBase + (offsetChannel + 1u) * offsetPlane + spatial];
-    let sy = f32(i32(oh * STRIDE_H) + {{ kh * source.dilationH }} - PAD_H) + dy;
-    let sx = f32(i32(ow * STRIDE_W) + {{ kw * source.dilationW }} - PAD_W) + dx;
     if (sy > -1.0 && sy < f32(IN_H) && sx > -1.0 && sx < f32(IN_W)) {
       let h0 = i32(floor(sy)); let h1 = h0 + 1;
       let w0 = i32(floor(sx)); let w1 = w0 + 1;
       let lh = sy - f32(h0); let lw = sx - f32(w0);
       let hh = 1.0 - lh; let hw = 1.0 - lw;
-{% if source.hasMask %}
       let maskChannel = group * KAREA + TAP;
-      let maskBase = n * ({{ source.groups }}u * KAREA) * offsetPlane;
       let mv = mask[maskBase + maskChannel * offsetPlane + spatial];
 {% else %}
       let mv = 1.0;
@@ -104,8 +104,8 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
                    + lh * hw * pixel(n, ic, h1, w0)
                    + lh * lw * pixel(n, ic, h1, w1);
         let wi = ((oc * IN_CPG + localIc) * KAREA) + TAP;
-{% if source.outChannelBlock != 1 %}
-{% for v in range(source.outChannelVecs) %}
         let weights{{ v }} = vec4<f32>(
           w[wi + {{ v * 4 }}u * W_OC_STRIDE], w[wi + {{ v * 4 + 1 }}u * W_OC_STRIDE],
           w[wi + {{ v * 4 + 2 }}u * W_OC_STRIDE], w[wi + {{ v * 4 + 3 }}u * W_OC_STRIDE]);
@@ -119,10 +119,10 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   }
 {% endfor %}
 {% endfor %}
-{% if source.outChannelBlock != 1 %}
   let y0 = ((n * OUT_C + oc) * OUT_H + oh) * OUT_W + ow;
   let yPlane = OUT_H * OUT_W;
-{% for v in range(source.outChannelVecs) %}
   y[y0 + {{ v * 4 }}u * yPlane] = acc{{ v }}.x;
   y[y0 + {{ v * 4 + 1 }}u * yPlane] = acc{{ v }}.y;
   y[y0 + {{ v * 4 + 2 }}u * yPlane] = acc{{ v }}.z;

 // output group shares the same tap coordinates and modulation mask. Hoist the
 // offset fetch, floor, bounds and interpolation coefficients out of the channel
 // loop instead of repeating them C/group times for every output channel.
+const COUNT: u32 = {{ count }}u;
+const IN_C: u32 = {{ inC }}u;
+const IN_CPG: u32 = {{ inCPerGroup }}u;
+const IN_H: i32 = {{ inH }};
+const IN_W: i32 = {{ inW }};
+const IN_H_U: u32 = {{ inH }}u;
+const IN_W_U: u32 = {{ inW }}u;
+const OUT_C: u32 = {{ outC }}u;
+const OUT_CPG: u32 = {{ outCPerGroup }}u;
+{% if outChannelBlock != 1 %}const OC_BLOCK: u32 = {{ outChannelBlock }}u;
 {% endif %}
+const OUT_H: u32 = {{ outH }}u;
+const OUT_W: u32 = {{ outW }}u;
+const KAREA: u32 = {{ kernelHSpec * kernelWSpec }}u;
+{% if outChannelBlock != 1 %}const W_OC_STRIDE: u32 = IN_CPG * KAREA;
 {% endif %}
+const STRIDE_H: u32 = {{ strideHSpec }}u;
+const STRIDE_W: u32 = {{ strideWSpec }}u;
+const PAD_H: i32 = {{ padH }};
+const PAD_W: i32 = {{ padW }};
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 fn pixel(n: u32, c: u32, h: i32, wc: i32) -> f32 {
 }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   if (index >= COUNT) { return; }
+{% if outChannelBlock != 1 %}
   // Output channels in the same convolution group use identical deformable
   // coordinates, masks, and sampled input values. Accumulate a compile-time
   // output-channel block together so those values are shared across the block.
   t /= OUT_H;
   let ocBlock = t % (OUT_CPG / OC_BLOCK);
   t /= OUT_CPG / OC_BLOCK;
+  let group = t % {{ groups }}u;
+  let n = t / {{ groups }}u;
   let oc = group * OUT_CPG + ocBlock * OC_BLOCK;
+{% for v in range(outChannelVecs) %}
+{% if hasBias %}
   var acc{{ v }} = vec4<f32>(
     bias[oc + {{ v * 4 }}u], bias[oc + {{ v * 4 + 1 }}u],
     bias[oc + {{ v * 4 + 2 }}u], bias[oc + {{ v * 4 + 3 }}u]);
   let oc = t % OUT_C;
   let n = t / OUT_C;
   let group = oc / OUT_CPG;
+{% if hasBias %}
   var acc = bias[oc];
 {% else %}
   var acc = 0.0;
   let spatial = oh * OUT_W + ow;
   let offsetPlane = OUT_H * OUT_W;
+{% for kh in range(kernelHSpec) %}
+{% for kw in range(kernelWSpec) %}
   {
+    const TAP: u32 = {{ kh * kernelWSpec + kw }}u;
     let offsetChannel = group * KAREA * 2u + TAP * 2u;
+    let offsetBase = n * ({{ groups }}u * KAREA * 2u) * offsetPlane;
     let dy = offset[offsetBase + offsetChannel * offsetPlane + spatial];
     let dx = offset[offsetBase + (offsetChannel + 1u) * offsetPlane + spatial];
+    let sy = f32(i32(oh * STRIDE_H) + {{ kh * dilationHSpec }} - PAD_H) + dy;
+    let sx = f32(i32(ow * STRIDE_W) + {{ kw * dilationWSpec }} - PAD_W) + dx;
     if (sy > -1.0 && sy < f32(IN_H) && sx > -1.0 && sx < f32(IN_W)) {
       let h0 = i32(floor(sy)); let h1 = h0 + 1;
       let w0 = i32(floor(sx)); let w1 = w0 + 1;
       let lh = sy - f32(h0); let lw = sx - f32(w0);
       let hh = 1.0 - lh; let hw = 1.0 - lw;
+{% if hasMask %}
       let maskChannel = group * KAREA + TAP;
+      let maskBase = n * ({{ groups }}u * KAREA) * offsetPlane;
       let mv = mask[maskBase + maskChannel * offsetPlane + spatial];
 {% else %}
       let mv = 1.0;
                    + lh * hw * pixel(n, ic, h1, w0)
                    + lh * lw * pixel(n, ic, h1, w1);
         let wi = ((oc * IN_CPG + localIc) * KAREA) + TAP;
+{% if outChannelBlock != 1 %}
+{% for v in range(outChannelVecs) %}
         let weights{{ v }} = vec4<f32>(
           w[wi + {{ v * 4 }}u * W_OC_STRIDE], w[wi + {{ v * 4 + 1 }}u * W_OC_STRIDE],
           w[wi + {{ v * 4 + 2 }}u * W_OC_STRIDE], w[wi + {{ v * 4 + 3 }}u * W_OC_STRIDE]);
   }
 {% endfor %}
 {% endfor %}
+{% if outChannelBlock != 1 %}
   let y0 = ((n * OUT_C + oc) * OUT_H + oh) * OUT_W + ow;
   let yPlane = OUT_H * OUT_W;
+{% for v in range(outChannelVecs) %}
   y[y0 + {{ v * 4 }}u * yPlane] = acc{{ v }}.x;
   y[y0 + {{ v * 4 + 1 }}u * yPlane] = acc{{ v }}.y;
   y[y0 + {{ v * 4 + 2 }}u * yPlane] = acc{{ v }}.z;

build/webgpu/deform-conv2d-nchw.wgsl.jinja CHANGED Viewed

@@ -1,30 +1,30 @@
 {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
 {% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
 {% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at the device's per-axis workgroup
-  // limit; gid.y carries the high portion of the output index.
 {% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
 {% elif note == "element-limit" %}
   // 2D-folded flat element index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension limit.
 {% elif note == "dispatch" %}
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit.
 {% endif %}
 {% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
-  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
@@ -69,7 +69,7 @@ fn mask_at(n: u32, channel: u32, oh: u32, ow: u32) -> f32 {
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
 {{ flat_index_2d("index", guardInline=true) }}
   let ow = index % params.outW;

 {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
 {% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; gid.y carries the high portion of the output index.
 {% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (the dispatch caps x and spills into y).
 {% elif note == "element-limit" %}
   // 2D-folded flat element index: gid.y carries the high bits past the
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d("index", guardInline=true) }}
   let ow = index % params.outW;

build/webgpu/manifest.json CHANGED Viewed

@@ -2,69 +2,31 @@
   "domain": "ai.onnx",
   "name": "DeformConv",
   "sinceVersion": 22,
-  "description": "Performs deformable convolution by sampling input at spatially offset locations specified per output position, enabling the kernel to adapt its receptive field shape. The `offset` tensor provides `(y, x)` offsets for each kernel point and output location; fractional offsets are resolved via bilinear interpolation and out-of-bounds locations contribute zero. An optional `mask` tensor modulates each sampling point's contribution.",
-  "inputs": [
-    {
-      "role": "X",
-      "dtype": "T",
-      "description": "Input data tensor of shape `(N, C, H, W)` for 2D or `(N, C, D1, ..., Dn)` in general.",
-      "rank": 4
-    },
-    {
-      "role": "W",
-      "dtype": "T",
-      "description": "Convolution weight tensor of shape `(oC, C/group, kH, kW)`.",
-      "rank": 4
-    },
-    {
-      "role": "offset",
-      "dtype": "T",
-      "description": "Per-output-position sampling offsets of shape `(N, offset_group * kH * kW * 2, oH, oW)` for 2D data.",
-      "rank": 4
-    },
-    {
-      "role": "B",
-      "dtype": "T",
-      "optional": true,
-      "description": "Optional 1D bias of length oC added to the convolution output.",
-      "rank": 1
-    },
-    {
-      "role": "mask",
-      "dtype": "T",
-      "optional": true,
-      "description": "Optional modulation mask of shape `(N, offset_group * kH * kW, oH, oW)` scaling each sampled kernel point; defaults to ones.",
-      "rank": 4
-    }
-  ],
-  "outputs": [
-    {
-      "role": "Y",
       "dtype": "T",
-      "description": "Output tensor of shape `(N, oC, oH, oW)` containing the deformable convolution result.",
       "rank": 4,
-      "shape": ["dim(shapes.X, 0)", "dim(shapes.W, 0)", "floor((dim(shapes.X, 2) + padTop + padBottom - ((dim(shapes.W, 2) - 1) * dilationH + 1)) / strideH) + 1", "floor((dim(shapes.X, 3) + padLeft + padRight - ((dim(shapes.W, 3) - 1) * dilationW + 1)) / strideW) + 1"]
     }
-  ],
-  "attributes": { "group": 1, "offset_group": 1 },
-  "attributeDescriptions": {
-    "dilations": "Dilation factors for the spatial axes in height-width order; defaults to `[1, 1]`.",
-    "group": "Number of groups the input channels C and output channels oC are divided into; both must be divisible by group.",
-    "kernel_shape": "Kernel dimensions in height-width order. When omitted, they are inferred from the spatial dimensions of `W`.",
-    "offset_group": "Number of offset groups; input channels `C` must be divisible by `offset_group`.",
-    "pads": "Padding at the beginning and end of each spatial axis in `[top, left, bottom, right]` order; defaults to zeros.",
-    "strides": "Stride factors for the spatial axes in height-width order; defaults to `[1, 1]`."
   },
-  "typeConstraints": { "T": ["float32"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "w": { "kind": "tensor", "semantic": "W", "role": "input" },
-    "offset": { "kind": "tensor", "semantic": "offset", "role": "input" },
-    "bias": { "kind": "tensor", "semantic": "B", "role": "input", "required": false },
-    "mask": { "kind": "tensor", "semantic": "mask", "role": "input", "required": false },
-    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
   },
-  "tunables": { "WORKGROUP_SIZE": 128 },
   "derive": {
     "spatialRank": 2,
     "kernelShapeLengthOk": "not has(attrs, \"kernel_shape\") or (attrs.kernel_shape | length) == spatialRank",
@@ -81,157 +43,52 @@
     "padLeft": "attrs.pads[spatialRank - 1] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 1 else 0",
     "padBottom": "attrs.pads[2 * spatialRank - 2] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 2 else 0",
     "padRight": "attrs.pads[2 * spatialRank - 1] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 1 else 0",
-    "spatialAttributesOk": "kernelShapeLengthOk and stridesLengthOk and dilationsLengthOk and padsLengthOk and strideH >= 1 and strideW >= 1 and dilationH >= 1 and dilationW >= 1 and padTop >= 0 and padLeft >= 0 and padBottom >= 0 and padRight >= 0 and (not has(attrs, \"kernel_shape\") or (kernelH == dim(shapes.W, 2) and kernelW == dim(shapes.W, 3)))",
-    "alignedGroupsCoreOk": "spatialAttributesOk and ranks.X == 4 and ranks.W == 4 and ranks.offset == 4 and ranks.Y == 4 and attrs.group >= 1 and attrs.offset_group == attrs.group and dim(shapes.W, 2) >= 1 and dim(shapes.W, 2) <= 7 and dim(shapes.W, 3) >= 1 and dim(shapes.W, 3) <= 7 and dim(shapes.W, 1) * attrs.group == dim(shapes.X, 1) and dim(shapes.W, 0) % attrs.group == 0 and dim(shapes.Y, 0) == dim(shapes.X, 0) and dim(shapes.Y, 1) == dim(shapes.W, 0) and dim(shapes.Y, 2) == floor((dim(shapes.X, 2) + padTop + padBottom - ((dim(shapes.W, 2) - 1) * dilationH + 1)) / strideH) + 1 and dim(shapes.Y, 3) == floor((dim(shapes.X, 3) + padLeft + padRight - ((dim(shapes.W, 3) - 1) * dilationW + 1)) / strideW) + 1 and dim(shapes.offset, 0) == dim(shapes.X, 0) and dim(shapes.offset, 1) == attrs.offset_group * dim(shapes.W, 2) * dim(shapes.W, 3) * 2 and dim(shapes.offset, 2) == dim(shapes.Y, 2) and dim(shapes.offset, 3) == dim(shapes.Y, 3)",
-    "alignedBiasMaskOk": "alignedGroupsCoreOk and present.bias and present.mask and ranks.B == 1 and ranks.mask == 4 and dim(shapes.B, 0) == dim(shapes.W, 0) and dim(shapes.mask, 0) == dim(shapes.X, 0) and dim(shapes.mask, 1) == attrs.offset_group * dim(shapes.W, 2) * dim(shapes.W, 3) and dim(shapes.mask, 2) == dim(shapes.Y, 2) and dim(shapes.mask, 3) == dim(shapes.Y, 3)",
     "alignedPlainOk": "alignedGroupsCoreOk and not present.bias and not present.mask",
-    "genericCoreOk": "spatialAttributesOk and ranks.X == 4 and ranks.W == 4 and ranks.offset == 4 and ranks.Y == 4 and attrs.group >= 1 and attrs.offset_group >= 1 and dim(shapes.W, 1) * attrs.group == dim(shapes.X, 1) and dim(shapes.X, 1) % attrs.offset_group == 0 and dim(shapes.W, 0) % attrs.group == 0 and dim(shapes.Y, 0) == dim(shapes.X, 0) and dim(shapes.Y, 1) == dim(shapes.W, 0) and dim(shapes.Y, 2) == floor((dim(shapes.X, 2) + padTop + padBottom - ((dim(shapes.W, 2) - 1) * dilationH + 1)) / strideH) + 1 and dim(shapes.Y, 3) == floor((dim(shapes.X, 3) + padLeft + padRight - ((dim(shapes.W, 3) - 1) * dilationW + 1)) / strideW) + 1 and dim(shapes.offset, 0) == dim(shapes.X, 0) and dim(shapes.offset, 1) == attrs.offset_group * dim(shapes.W, 2) * dim(shapes.W, 3) * 2 and dim(shapes.offset, 2) == dim(shapes.Y, 2) and dim(shapes.offset, 3) == dim(shapes.Y, 3)",
-    "genericBiasShapeOk": "present.bias and ranks.B == 1 and dim(shapes.B, 0) == dim(shapes.W, 0)",
-    "genericMaskShapeOk": "present.mask and ranks.mask == 4 and dim(shapes.mask, 0) == dim(shapes.X, 0) and dim(shapes.mask, 1) == attrs.offset_group * dim(shapes.W, 2) * dim(shapes.W, 3) and dim(shapes.mask, 2) == dim(shapes.Y, 2) and dim(shapes.mask, 3) == dim(shapes.Y, 3)",
     "genericPlainOk": "genericCoreOk and not present.bias and not present.mask",
     "genericBiasOk": "genericCoreOk and genericBiasShapeOk and not present.mask",
     "genericMaskOk": "genericCoreOk and not present.bias and genericMaskShapeOk",
     "genericBiasMaskOk": "genericCoreOk and genericBiasShapeOk and genericMaskShapeOk",
-    "alignedOcBlock": "16 if (dim(shapes.W, 0) / attrs.group) % 16 == 0 else (8 if (dim(shapes.W, 0) / attrs.group) % 8 == 0 else (4 if (dim(shapes.W, 0) / attrs.group) % 4 == 0 else 1))",
     "alignedOcVecs": "4 if alignedOcBlock == 16 else (2 if alignedOcBlock == 8 else (1 if alignedOcBlock == 4 else 0))"
   },
-  "bindingSets": {
-    "alignedBiasMask": [
-      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      { "name": "w", "arg": "w", "semantic": "W", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      {
-        "name": "offset",
-        "arg": "offset",
-        "semantic": "offset",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "bias",
-        "arg": "bias",
-        "semantic": "B",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "mask",
-        "arg": "mask",
-        "semantic": "mask",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "f32" }
-    ],
-    "alignedPlain": [
-      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      { "name": "w", "arg": "w", "semantic": "W", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      {
-        "name": "offset",
-        "arg": "offset",
-        "semantic": "offset",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "f32" }
-    ],
-    "nchw2dPlain": [
-      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      { "name": "w", "arg": "w", "semantic": "W", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      {
-        "name": "offset",
-        "arg": "offset",
-        "semantic": "offset",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "f32" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
-            { "name": "inChannels", "type": "u32", "value": "dim(shapes.X, 1)" },
-            { "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
-            { "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
-            { "name": "outChannels", "type": "u32", "value": "dim(shapes.W, 0)" },
-            { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.W, 1)" },
-            { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.X, 1) / attrs.group" },
-            { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.W, 0) / attrs.group" },
-            { "name": "inChannelsPerOffsetGroup", "type": "u32", "value": "dim(shapes.X, 1) / attrs.offset_group" },
-            { "name": "kernelH", "type": "u32", "value": "dim(shapes.W, 2)" },
-            { "name": "kernelW", "type": "u32", "value": "dim(shapes.W, 3)" },
-            { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
-            { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" },
-            { "name": "strideH", "type": "u32", "value": "strideH" },
-            { "name": "strideW", "type": "u32", "value": "strideW" },
-            { "name": "dilationH", "type": "u32", "value": "dilationH" },
-            { "name": "dilationW", "type": "u32", "value": "dilationW" },
-            { "name": "padH", "type": "i32", "value": "padTop" },
-            { "name": "padW", "type": "i32", "value": "padLeft" },
-            { "name": "offsetGroup", "type": "u32", "value": "attrs.offset_group" }
-          ]
-        }
-      }
-    ],
-    "nchw2dBiasMask": [
-      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      { "name": "w", "arg": "w", "semantic": "W", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      {
-        "name": "offset",
-        "arg": "offset",
-        "semantic": "offset",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "bias",
-        "arg": "bias",
-        "semantic": "B",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "mask",
-        "arg": "mask",
-        "semantic": "mask",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "f32" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
-            { "name": "inChannels", "type": "u32", "value": "dim(shapes.X, 1)" },
-            { "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
-            { "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
-            { "name": "outChannels", "type": "u32", "value": "dim(shapes.W, 0)" },
-            { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.W, 1)" },
-            { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.X, 1) / attrs.group" },
-            { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.W, 0) / attrs.group" },
-            { "name": "inChannelsPerOffsetGroup", "type": "u32", "value": "dim(shapes.X, 1) / attrs.offset_group" },
-            { "name": "kernelH", "type": "u32", "value": "dim(shapes.W, 2)" },
-            { "name": "kernelW", "type": "u32", "value": "dim(shapes.W, 3)" },
-            { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
-            { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" },
-            { "name": "strideH", "type": "u32", "value": "strideH" },
-            { "name": "strideW", "type": "u32", "value": "strideW" },
-            { "name": "dilationH", "type": "u32", "value": "dilationH" },
-            { "name": "dilationW", "type": "u32", "value": "dilationW" },
-            { "name": "padH", "type": "i32", "value": "padTop" },
-            { "name": "padW", "type": "i32", "value": "padLeft" },
-            { "name": "offsetGroup", "type": "u32", "value": "attrs.offset_group" }
-          ]
-        }
-      }
-    ]
   },
   "variants": [
     {
@@ -242,35 +99,37 @@
         {
           "id": "main",
           "name": "DeformConv.AlignedGroupsBlocked",
-          "source": {
-            "shader": "deform-conv2d-aligned-groups.wgsl.jinja",
-            "inputs": {
-              "count": "numel(shapes.Y) / alignedOcBlock",
-              "inC": "dim(shapes.X, 1)",
-              "inCPerGroup": "dim(shapes.X, 1) / attrs.group",
-              "inH": "dim(shapes.X, 2)",
-              "inW": "dim(shapes.X, 3)",
-              "outC": "dim(shapes.Y, 1)",
-              "outCPerGroup": "dim(shapes.Y, 1) / attrs.group",
-              "outChannelBlock": "alignedOcBlock",
-              "outChannelVecs": "alignedOcVecs",
-              "groups": "attrs.group",
-              "outH": "dim(shapes.Y, 2)",
-              "outW": "dim(shapes.Y, 3)",
-              "kernelH": "dim(shapes.W, 2)",
-              "kernelW": "dim(shapes.W, 3)",
-              "strideH": "strideH",
-              "strideW": "strideW",
-              "dilationH": "dilationH",
-              "dilationW": "dilationW",
-              "padH": "padTop",
-              "padW": "padLeft",
-              "hasBias": true,
-              "hasMask": true
-            }
           },
-          "bindings": "alignedBiasMask",
-          "dispatch": { "threads": "numel(shapes.Y) / alignedOcBlock", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -282,143 +141,93 @@
         {
           "id": "main",
           "name": "DeformConv.AlignedGroupsBlocked",
-          "source": {
-            "shader": "deform-conv2d-aligned-groups.wgsl.jinja",
-            "inputs": {
-              "count": "numel(shapes.Y) / alignedOcBlock",
-              "inC": "dim(shapes.X, 1)",
-              "inCPerGroup": "dim(shapes.X, 1) / attrs.group",
-              "inH": "dim(shapes.X, 2)",
-              "inW": "dim(shapes.X, 3)",
-              "outC": "dim(shapes.Y, 1)",
-              "outCPerGroup": "dim(shapes.Y, 1) / attrs.group",
-              "outChannelBlock": "alignedOcBlock",
-              "outChannelVecs": "alignedOcVecs",
-              "groups": "attrs.group",
-              "outH": "dim(shapes.Y, 2)",
-              "outW": "dim(shapes.Y, 3)",
-              "kernelH": "dim(shapes.W, 2)",
-              "kernelW": "dim(shapes.W, 3)",
-              "strideH": "strideH",
-              "strideW": "strideW",
-              "dilationH": "dilationH",
-              "dilationW": "dilationW",
-              "padH": "padTop",
-              "padW": "padLeft",
-              "hasBias": false,
-              "hasMask": false
-            }
           },
-          "bindings": "alignedPlain",
-          "dispatch": { "threads": "numel(shapes.Y) / alignedOcBlock", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "nchw2d",
       "when": ["genericPlainOk"],
-      "constants": { "hasBias": false, "hasMask": false },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
-          "bindings": "nchw2dPlain",
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "nchw2d_bias_mask",
       "when": ["genericBiasMaskOk"],
-      "constants": { "hasBias": true, "hasMask": true },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
-          "bindings": "nchw2dBiasMask",
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
-      ],
-      "priority": 10
     },
     {
       "id": "nchw2d_bias",
       "priority": 5,
       "when": ["genericBiasOk"],
-      "constants": { "hasBias": true, "hasMask": false },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
-          "bindings": [
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            {
-              "name": "w",
-              "arg": "w",
-              "semantic": "W",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            {
-              "name": "offset",
-              "arg": "offset",
-              "semantic": "offset",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            {
-              "name": "bias",
-              "arg": "bias",
-              "semantic": "B",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "f32" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [
-                  { "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
-                  { "name": "inChannels", "type": "u32", "value": "dim(shapes.X, 1)" },
-                  { "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
-                  { "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
-                  { "name": "outChannels", "type": "u32", "value": "dim(shapes.W, 0)" },
-                  { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.W, 1)" },
-                  { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.X, 1) / attrs.group" },
-                  { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.W, 0) / attrs.group" },
-                  {
-                    "name": "inChannelsPerOffsetGroup",
-                    "type": "u32",
-                    "value": "dim(shapes.X, 1) / attrs.offset_group"
-                  },
-                  { "name": "kernelH", "type": "u32", "value": "dim(shapes.W, 2)" },
-                  { "name": "kernelW", "type": "u32", "value": "dim(shapes.W, 3)" },
-                  { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
-                  { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" },
-                  { "name": "strideH", "type": "u32", "value": "strideH" },
-                  { "name": "strideW", "type": "u32", "value": "strideW" },
-                  { "name": "dilationH", "type": "u32", "value": "dilationH" },
-                  { "name": "dilationW", "type": "u32", "value": "dilationW" },
-                  { "name": "padH", "type": "i32", "value": "padTop" },
-                  { "name": "padW", "type": "i32", "value": "padLeft" },
-                  { "name": "offsetGroup", "type": "u32", "value": "attrs.offset_group" }
-                ]
-              }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
@@ -426,78 +235,18 @@
       "id": "nchw2d_mask",
       "priority": 5,
       "when": ["genericMaskOk"],
-      "constants": { "hasBias": false, "hasMask": true },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
-          "bindings": [
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            {
-              "name": "w",
-              "arg": "w",
-              "semantic": "W",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            {
-              "name": "offset",
-              "arg": "offset",
-              "semantic": "offset",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            {
-              "name": "mask",
-              "arg": "mask",
-              "semantic": "mask",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "f32"
-            },
-            { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "f32" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [
-                  { "name": "count", "type": "u32", "value": "numel(shapes.Y)" },
-                  { "name": "inChannels", "type": "u32", "value": "dim(shapes.X, 1)" },
-                  { "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
-                  { "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
-                  { "name": "outChannels", "type": "u32", "value": "dim(shapes.W, 0)" },
-                  { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.W, 1)" },
-                  { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.X, 1) / attrs.group" },
-                  { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.W, 0) / attrs.group" },
-                  {
-                    "name": "inChannelsPerOffsetGroup",
-                    "type": "u32",
-                    "value": "dim(shapes.X, 1) / attrs.offset_group"
-                  },
-                  { "name": "kernelH", "type": "u32", "value": "dim(shapes.W, 2)" },
-                  { "name": "kernelW", "type": "u32", "value": "dim(shapes.W, 3)" },
-                  { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
-                  { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" },
-                  { "name": "strideH", "type": "u32", "value": "strideH" },
-                  { "name": "strideW", "type": "u32", "value": "strideW" },
-                  { "name": "dilationH", "type": "u32", "value": "dilationH" },
-                  { "name": "dilationW", "type": "u32", "value": "dilationW" },
-                  { "name": "padH", "type": "i32", "value": "padTop" },
-                  { "name": "padW", "type": "i32", "value": "padLeft" },
-                  { "name": "offsetGroup", "type": "u32", "value": "attrs.offset_group" }
-                ]
-              }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "DeformConv",
   "sinceVersion": 22,
+  "inputs": {
+    "x": { "onnx": "X", "dtype": "T", "rank": 4 },
+    "w": { "onnx": "W", "dtype": "T", "rank": 4 },
+    "offset": { "dtype": "T", "rank": 4 },
+    "bias": { "onnx": "B", "dtype": "T", "rank": 1, "optional": true },
+    "mask": { "dtype": "T", "rank": 4, "optional": true }
+  },
+  "outputs": {
+    "y": {
+      "onnx": "Y",
       "dtype": "T",
       "rank": 4,
+      "shape": ["dim(shapes.x, 0)", "dim(shapes.w, 0)", "floor((dim(shapes.x, 2) + padTop + padBottom - ((dim(shapes.w, 2) - 1) * dilationH + 1)) / strideH) + 1", "floor((dim(shapes.x, 3) + padLeft + padRight - ((dim(shapes.w, 3) - 1) * dilationW + 1)) / strideW) + 1"]
     }
   },
+  "attributes": {
+    "group": { "default": 1 },
+    "offset_group": { "default": 1 },
+    "dilations": {},
+    "kernel_shape": {},
+    "pads": {},
+    "strides": {}
   },
+  "typeConstraints": { "T": ["float32"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 128 } },
   "derive": {
     "spatialRank": 2,
     "kernelShapeLengthOk": "not has(attrs, \"kernel_shape\") or (attrs.kernel_shape | length) == spatialRank",
     "padLeft": "attrs.pads[spatialRank - 1] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 1 else 0",
     "padBottom": "attrs.pads[2 * spatialRank - 2] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 2 else 0",
     "padRight": "attrs.pads[2 * spatialRank - 1] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 1 else 0",
+    "spatialAttributesOk": "kernelShapeLengthOk and stridesLengthOk and dilationsLengthOk and padsLengthOk and strideH >= 1 and strideW >= 1 and dilationH >= 1 and dilationW >= 1 and padTop >= 0 and padLeft >= 0 and padBottom >= 0 and padRight >= 0 and (not has(attrs, \"kernel_shape\") or (kernelH == dim(shapes.w, 2) and kernelW == dim(shapes.w, 3)))",
+    "alignedGroupsCoreOk": "spatialAttributesOk and ranks.x == 4 and ranks.w == 4 and ranks.offset == 4 and ranks.y == 4 and attrs.group >= 1 and attrs.offset_group == attrs.group and dim(shapes.w, 2) >= 1 and dim(shapes.w, 2) <= 7 and dim(shapes.w, 3) >= 1 and dim(shapes.w, 3) <= 7 and dim(shapes.w, 1) * attrs.group == dim(shapes.x, 1) and dim(shapes.w, 0) % attrs.group == 0 and dim(shapes.y, 0) == dim(shapes.x, 0) and dim(shapes.y, 1) == dim(shapes.w, 0) and dim(shapes.y, 2) == floor((dim(shapes.x, 2) + padTop + padBottom - ((dim(shapes.w, 2) - 1) * dilationH + 1)) / strideH) + 1 and dim(shapes.y, 3) == floor((dim(shapes.x, 3) + padLeft + padRight - ((dim(shapes.w, 3) - 1) * dilationW + 1)) / strideW) + 1 and dim(shapes.offset, 0) == dim(shapes.x, 0) and dim(shapes.offset, 1) == attrs.offset_group * dim(shapes.w, 2) * dim(shapes.w, 3) * 2 and dim(shapes.offset, 2) == dim(shapes.y, 2) and dim(shapes.offset, 3) == dim(shapes.y, 3)",
+    "alignedBiasMaskOk": "alignedGroupsCoreOk and present.bias and present.mask and ranks.bias == 1 and ranks.mask == 4 and dim(shapes.bias, 0) == dim(shapes.w, 0) and dim(shapes.mask, 0) == dim(shapes.x, 0) and dim(shapes.mask, 1) == attrs.offset_group * dim(shapes.w, 2) * dim(shapes.w, 3) and dim(shapes.mask, 2) == dim(shapes.y, 2) and dim(shapes.mask, 3) == dim(shapes.y, 3)",
     "alignedPlainOk": "alignedGroupsCoreOk and not present.bias and not present.mask",
+    "genericCoreOk": "spatialAttributesOk and ranks.x == 4 and ranks.w == 4 and ranks.offset == 4 and ranks.y == 4 and attrs.group >= 1 and attrs.offset_group >= 1 and dim(shapes.w, 1) * attrs.group == dim(shapes.x, 1) and dim(shapes.x, 1) % attrs.offset_group == 0 and dim(shapes.w, 0) % attrs.group == 0 and dim(shapes.y, 0) == dim(shapes.x, 0) and dim(shapes.y, 1) == dim(shapes.w, 0) and dim(shapes.y, 2) == floor((dim(shapes.x, 2) + padTop + padBottom - ((dim(shapes.w, 2) - 1) * dilationH + 1)) / strideH) + 1 and dim(shapes.y, 3) == floor((dim(shapes.x, 3) + padLeft + padRight - ((dim(shapes.w, 3) - 1) * dilationW + 1)) / strideW) + 1 and dim(shapes.offset, 0) == dim(shapes.x, 0) and dim(shapes.offset, 1) == attrs.offset_group * dim(shapes.w, 2) * dim(shapes.w, 3) * 2 and dim(shapes.offset, 2) == dim(shapes.y, 2) and dim(shapes.offset, 3) == dim(shapes.y, 3)",
+    "genericBiasShapeOk": "present.bias and ranks.bias == 1 and dim(shapes.bias, 0) == dim(shapes.w, 0)",
+    "genericMaskShapeOk": "present.mask and ranks.mask == 4 and dim(shapes.mask, 0) == dim(shapes.x, 0) and dim(shapes.mask, 1) == attrs.offset_group * dim(shapes.w, 2) * dim(shapes.w, 3) and dim(shapes.mask, 2) == dim(shapes.y, 2) and dim(shapes.mask, 3) == dim(shapes.y, 3)",
     "genericPlainOk": "genericCoreOk and not present.bias and not present.mask",
     "genericBiasOk": "genericCoreOk and genericBiasShapeOk and not present.mask",
     "genericMaskOk": "genericCoreOk and not present.bias and genericMaskShapeOk",
     "genericBiasMaskOk": "genericCoreOk and genericBiasShapeOk and genericMaskShapeOk",
+    "alignedOcBlock": "16 if (dim(shapes.w, 0) / attrs.group) % 16 == 0 else (8 if (dim(shapes.w, 0) / attrs.group) % 8 == 0 else (4 if (dim(shapes.w, 0) / attrs.group) % 4 == 0 else 1))",
     "alignedOcVecs": "4 if alignedOcBlock == 16 else (2 if alignedOcBlock == 8 else (1 if alignedOcBlock == 4 else 0))"
   },
+  "bindings": {
+    "x": { "buffer": "read-only-storage", "elementType": "f32" },
+    "w": { "buffer": "read-only-storage", "elementType": "f32" },
+    "offset": { "buffer": "read-only-storage", "elementType": "f32" },
+    "bias": { "buffer": "read-only-storage", "elementType": "f32" },
+    "mask": { "buffer": "read-only-storage", "elementType": "f32" },
+    "y": { "buffer": "storage", "elementType": "f32" },
+    "params": {
+      "buffer": "uniform",
+      "struct": [
+        { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
+        { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
+        { "name": "inH", "type": "u32", "value": "dim(shapes.x, 2)" },
+        { "name": "inW", "type": "u32", "value": "dim(shapes.x, 3)" },
+        { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
+        { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
+        { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
+        { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.w, 0) / attrs.group" },
+        { "name": "inChannelsPerOffsetGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.offset_group" },
+        { "name": "kernelH", "type": "u32", "value": "dim(shapes.w, 2)" },
+        { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 3)" },
+        { "name": "outH", "type": "u32", "value": "dim(shapes.y, 2)" },
+        { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" },
+        { "name": "strideH", "type": "u32", "value": "strideH" },
+        { "name": "strideW", "type": "u32", "value": "strideW" },
+        { "name": "dilationH", "type": "u32", "value": "dilationH" },
+        { "name": "dilationW", "type": "u32", "value": "dilationW" },
+        { "name": "padH", "type": "i32", "value": "padTop" },
+        { "name": "padW", "type": "i32", "value": "padLeft" },
+        { "name": "offsetGroup", "type": "u32", "value": "attrs.offset_group" }
+      ]
+    }
   },
   "variants": [
     {
         {
           "id": "main",
           "name": "DeformConv.AlignedGroupsBlocked",
+          "shader": "deform-conv2d-aligned-groups.wgsl.jinja",
+          "derive": {
+            "count": "numel(shapes.y) / alignedOcBlock",
+            "inC": "dim(shapes.x, 1)",
+            "inCPerGroup": "dim(shapes.x, 1) / attrs.group",
+            "inH": "dim(shapes.x, 2)",
+            "inW": "dim(shapes.x, 3)",
+            "outC": "dim(shapes.y, 1)",
+            "outCPerGroup": "dim(shapes.y, 1) / attrs.group",
+            "outChannelBlock": "alignedOcBlock",
+            "outChannelVecs": "alignedOcVecs",
+            "groups": "attrs.group",
+            "outH": "dim(shapes.y, 2)",
+            "outW": "dim(shapes.y, 3)",
+            "kernelHSpec": "dim(shapes.w, 2)",
+            "kernelWSpec": "dim(shapes.w, 3)",
+            "strideHSpec": "strideH",
+            "strideWSpec": "strideW",
+            "dilationHSpec": "dilationH",
+            "dilationWSpec": "dilationW",
+            "padH": "padTop",
+            "padW": "padLeft",
+            "hasBias": true,
+            "hasMask": true
           },
+          "bindings": ["x", "w", "offset", "bias", "mask", "y"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / alignedOcBlock), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / alignedOcBlock), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
         {
           "id": "main",
           "name": "DeformConv.AlignedGroupsBlocked",
+          "shader": "deform-conv2d-aligned-groups.wgsl.jinja",
+          "derive": {
+            "count": "numel(shapes.y) / alignedOcBlock",
+            "inC": "dim(shapes.x, 1)",
+            "inCPerGroup": "dim(shapes.x, 1) / attrs.group",
+            "inH": "dim(shapes.x, 2)",
+            "inW": "dim(shapes.x, 3)",
+            "outC": "dim(shapes.y, 1)",
+            "outCPerGroup": "dim(shapes.y, 1) / attrs.group",
+            "outChannelBlock": "alignedOcBlock",
+            "outChannelVecs": "alignedOcVecs",
+            "groups": "attrs.group",
+            "outH": "dim(shapes.y, 2)",
+            "outW": "dim(shapes.y, 3)",
+            "kernelHSpec": "dim(shapes.w, 2)",
+            "kernelWSpec": "dim(shapes.w, 3)",
+            "strideHSpec": "strideH",
+            "strideWSpec": "strideW",
+            "dilationHSpec": "dilationH",
+            "dilationWSpec": "dilationW",
+            "padH": "padTop",
+            "padW": "padLeft",
+            "hasBias": false,
+            "hasMask": false
           },
+          "bindings": ["x", "w", "offset", "y"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / alignedOcBlock), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / alignedOcBlock), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "nchw2d",
       "when": ["genericPlainOk"],
+      "derive": { "hasBias": false, "hasMask": false },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w", "offset", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "nchw2d_bias_mask",
+      "priority": 10,
       "when": ["genericBiasMaskOk"],
+      "derive": { "hasBias": true, "hasMask": true },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w", "offset", "bias", "mask", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
+      ]
     },
     {
       "id": "nchw2d_bias",
       "priority": 5,
       "when": ["genericBiasOk"],
+      "derive": { "hasBias": true, "hasMask": false },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w", "offset", "bias", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
       "id": "nchw2d_mask",
       "priority": 5,
       "when": ["genericMaskOk"],
+      "derive": { "hasBias": false, "hasMask": true },
       "passes": [
         {
           "id": "main",
           "name": "DeformConv",
           "shader": "deform-conv2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w", "offset", "mask", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,29 @@
 {
   "name": "ai.onnx.DeformConv",
-  "id": "_ai_onnx_deformconv_webgpu_3b61063",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "LJ4HwKvQ977QEVGUZvzpFnDdxEAfvPP7zF1Dll/uke0=",
-      "deform-conv2d-aligned-groups.wgsl.jinja": "8Wli1NLpk4sMZgYYbCntOUWTq8HJjsxZVBRfi6mBofY=",
-      "deform-conv2d-nchw.wgsl.jinja": "T7Q8v4p8H6Gpg9DCjp7TN4hzhAWyu6qHGI7URjr0Ri4=",
-      "manifest.json": "KTQl1tyOvIadWeHKJFZs6b0bcYG2SH7m+zhqbGuYHgw=",
-      "test.json": "evZiy5ZGtO0ZxwhNcN0uDHUwi27QlN+r5ghmsDXmUZw="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.DeformConv" }
 }

 {
   "name": "ai.onnx.DeformConv",
+  "id": "_ai_onnx_deformconv_webgpu_bd95a88",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "QuUDCU5THdFkh2/r2/k4/Uo61MTtbMbE8R5uE3XZagU=",
+      "deform-conv2d-aligned-groups.wgsl.jinja": "dtmKq+oPJLAziEgWffUCuNR05TCDSWuFTpIeKrkfhXA=",
+      "deform-conv2d-nchw.wgsl.jinja": "hjYa940wbG47JZU6xEB1f+ocmLfCTdKYd2bZWwZ+8VU=",
+      "manifest.json": "p+fNwbHM9W6A5oca8kgfytxZOvic1x+qOPu44gH6Kts=",
+      "test.json": "868t0+sfE0Szo3ySh3cdEN+OTQ+77X7pAUXkgxxZqrc="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "nchw2d_bias_mask_aligned_groups_blocked": ["deform-conv2d-aligned-groups.wgsl.jinja"],
+      "nchw2d_aligned_groups_blocked": ["deform-conv2d-aligned-groups.wgsl.jinja"],
+      "nchw2d": ["deform-conv2d-nchw.wgsl.jinja"],
+      "nchw2d_bias_mask": ["deform-conv2d-nchw.wgsl.jinja"],
+      "nchw2d_bias": ["deform-conv2d-nchw.wgsl.jinja"],
+      "nchw2d_mask": ["deform-conv2d-nchw.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.DeformConv",
   "fixtureArrays": {
     "onnx_backend_basic_deform_conv_without_padding_input_offset": [0.5, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -0.10000000149011612, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]
   },
@@ -128,7 +127,7 @@
         "w": {
           "dtype": "float32",
           "shape": [1, 1, 2, 2],
-          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, -1.0] }
         },
         "offset": { "dtype": "float32", "shape": [1, 8, 2, 2], "data": { "kind": "constant", "value": 0.0 } }
       },
@@ -211,7 +210,7 @@
     {
       "name": "mask_zeros_preserve_bias",
       "inputs": {
-        "x": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
         "w": { "dtype": "float32", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [7.0] } },
         "offset": { "dtype": "float32", "shape": [1, 2, 2, 2], "data": { "kind": "constant", "value": 0.0 } },
         "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
@@ -788,7 +787,7 @@
         "offset": {
           "dtype": "float32",
           "shape": [1, 2, 2, 2],
-          "data": { "kind": "values", "values": [2.0, 2.0, 2.0, 2.0, 0.0, 0.0, 0.0, 0.0] }
         }
       },
       "outputs": {
@@ -796,7 +795,7 @@
           "dtype": "float32",
           "shape": [1, 1, 2, 2],
           "tolerance": 0,
-          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] }
         }
       }
     },

 {
   "fixtureArrays": {
     "onnx_backend_basic_deform_conv_without_padding_input_offset": [0.5, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -0.10000000149011612, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]
   },
         "w": {
           "dtype": "float32",
           "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, -2.0] }
         },
         "offset": { "dtype": "float32", "shape": [1, 8, 2, 2], "data": { "kind": "constant", "value": 0.0 } }
       },
     {
       "name": "mask_zeros_preserve_bias",
       "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "constant", "value": 2.0 } },
         "w": { "dtype": "float32", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [7.0] } },
         "offset": { "dtype": "float32", "shape": [1, 2, 2, 2], "data": { "kind": "constant", "value": 0.0 } },
         "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
         "offset": {
           "dtype": "float32",
           "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [2.0, 2.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0] }
         }
       },
       "outputs": {
           "dtype": "float32",
           "shape": [1, 1, 2, 2],
           "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 0.0, 3.0, 4.0] }
         }
       }
     },