sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +14 -10
build/webgpu/bench.json +1 -2
build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja +4 -5
build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja +4 -5
build/webgpu/manifest.json +112 -453
build/webgpu/metadata.json +18 -8
build/webgpu/test.json +2 -3

README.md CHANGED Viewed

@@ -18,18 +18,18 @@ See the [ONNX Runtime `GatherBlockQuantized` contrib-operator spec](https://gith
 ## Inputs
-| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- | --- |
-| `data` | `dataT` | `T1` | runtime-selected; narrow integers and bool use 32-bit slots | `2` | — | Constant uint8 weight matrix. With `bits = 4`, each byte stores two values low-nibble first; with `bits = 8`, each byte stores one value. | required |
-| `indices` | `indicesT` | `Tind` | `uint32` | `1` | — | Non-negative logical int64 indices selecting rows from axis 0 of `data`. Every index must be less than the row count; values use checked uint32 WebGPU storage. | required |
-| `scales` | `scalesT` | `T2` | same as logical dtype | `2` | — | Per-block dequantization scale factors of shape `(rows, ceil(output_columns / block_size))`. | required |
-| `zero_points` | `zeroPointsT` | `T1` | runtime-selected; narrow integers and bool use 32-bit slots | `2` | — | Optional uint8 zero points. At 4 bits two zero points are packed per byte along the quantized axis, low-nibble first; at 8 bits the shape matches `scales`. If absent, uint8 data uses 2^(bits-1). | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `output` | `outputT` | `T2` | `2` | derived; see description | Dequantized floating-point output rows corresponding to the gathered indices. | required |
 ## Attributes
@@ -52,7 +52,7 @@ 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
@@ -61,10 +61,14 @@ 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 | WebGPU storage | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- | --- |
+| `dataT` | `data` | `T1` | runtime-selected; narrow integers and bool use 32-bit slots | `2` | — | Constant uint8 weight matrix. With `bits = 4`, each byte stores two values low-nibble first; with `bits = 8`, each byte stores one value. | required |
+| `indicesT` | `indices` | `Tind` | `uint32` | `1` | — | Non-negative logical int64 indices selecting rows from axis 0 of `data`. Every index must be less than the row count; values use checked uint32 WebGPU storage. | required |
+| `scalesT` | `scales` | `T2` | same as logical dtype | `2` | — | Per-block dequantization scale factors of shape `(rows, ceil(output_columns / block_size))`. | required |
+| `zeroPointsT` | `zero_points` | `T1` | runtime-selected; narrow integers and bool use 32-bit slots | `2` | — | Optional uint8 zero points. At 4 bits two zero points are packed per byte along the quantized axis, low-nibble first; at 8 bits the shape matches `scales`. If absent, uint8 data uses 2^(bits-1). | optional |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `outputT` | `output` | `T2` | `2` | derived | Dequantized floating-point output rows corresponding to the gathered indices. | required |
 ## Attributes
 ## 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": "com.microsoft.GatherBlockQuantized",
   "tunableSpace": { "workgroupSize": [64, 128, 256] },
   "cases": [
     {
@@ -237,7 +236,7 @@
       "name": "gather-block-q8-dispatch-healthy-idx1024-cols4096-1d",
       "preset": "stress",
       "provenance": {
-        "notes": "Stress-only capacity sibling: widened uint8 GPU storage brings the declared footprint to 280 MiB."
       },
       "vars": { "rows": 16384, "cols": 4096, "indexCount": 1024, "bits": 8, "blockSize": 32 },
       "attrs": { "bits": 8, "block_size": 32 },

 {
   "tunableSpace": { "workgroupSize": [64, 128, 256] },
   "cases": [
     {
       "name": "gather-block-q8-dispatch-healthy-idx1024-cols4096-1d",
       "preset": "stress",
       "provenance": {
+        "notes": "Widened uint8 GPU storage gives this capacity stress case a declared footprint of 280 MiB."
       },
       "vars": { "rows": 16384, "cols": 4096, "indexCount": 1024, "bits": 8, "blockSize": 32 },
       "attrs": { "bits": 8, "block_size": 32 },

build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja CHANGED Viewed

@@ -3,11 +3,10 @@
 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>,
-        @builtin(num_workgroups) nwg: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits past the maxComputeWorkgroupsPerDimension
-  // workgroup-per-dimension dispatch limit. Reduces to gid.x when nwg.y == 1.
-  let pair_index = gid.x + gid.y * nwg.x * WG;
   let total = params.indexCount * params.packedCols;
   if (pair_index >= total) {
     return;

 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the per-axis dispatch fold width.
+  // Reduces to gid.x when the dispatch does not fold.
+  let pair_index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   let total = params.indexCount * params.packedCols;
   if (pair_index >= total) {
     return;

build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja CHANGED Viewed

@@ -3,11 +3,10 @@
 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>,
-        @builtin(num_workgroups) nwg: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits past the maxComputeWorkgroupsPerDimension
-  // workgroup-per-dimension dispatch limit. Reduces to gid.x when nwg.y == 1.
-  let vec_index = gid.x + gid.y * nwg.x * WG;
 {% if scalarTail %}
   let row_vecs = (params.cols + 3u) / 4u;
 {% else %}

 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the per-axis dispatch fold width.
+  // Reduces to gid.x when the dispatch does not fold.
+  let vec_index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
 {% if scalarTail %}
   let row_vecs = (params.cols + 3u) / 4u;
 {% else %}

build/webgpu/manifest.json CHANGED Viewed

@@ -2,453 +2,111 @@
   "domain": "com.microsoft",
   "name": "GatherBlockQuantized",
   "sinceVersion": 1,
-  "description": "Gathers rows from a block-wise quantized weight matrix and dequantizes them. This inference implementation supports the standard `gather_axis = 0`, `quantize_axis = 1` matrix subset with uint8 `data`, 4-bit packed or 8-bit values, rank-1 non-negative in-bounds int64 `indices` projected to uint32 WebGPU storage, and float32 scales/output. Higher-rank gathers, negative indices, int32 indices, int4/uint4 data, 2-bit data, float16/bfloat16 output, and non-default axes are not implemented.",
-  "inputs": [
-    {
-      "role": "data",
-      "dtype": "T1",
-      "rank": 2,
-      "description": "Constant uint8 weight matrix. With `bits = 4`, each byte stores two values low-nibble first; with `bits = 8`, each byte stores one value."
-    },
-    {
-      "role": "indices",
-      "dtype": "Tind",
-      "rank": 1,
-      "description": "Non-negative logical int64 indices selecting rows from axis 0 of `data`. Every index must be less than the row count; values use checked uint32 WebGPU storage."
-    },
-    {
-      "role": "scales",
-      "dtype": "T2",
-      "rank": 2,
-      "description": "Per-block dequantization scale factors of shape `(rows, ceil(output_columns / block_size))`."
-    },
-    {
-      "role": "zero_points",
-      "dtype": "T1",
-      "rank": 2,
-      "optional": true,
-      "description": "Optional uint8 zero points. At 4 bits two zero points are packed per byte along the quantized axis, low-nibble first; at 8 bits the shape matches `scales`. If absent, uint8 data uses 2^(bits-1)."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "output",
       "dtype": "T2",
       "rank": 2,
-      "shape": "[dim(shapes.indices, 0), dim(shapes.data, 1) * (8 / attrs.bits)]",
-      "description": "Dequantized floating-point output rows corresponding to the gathered indices."
     }
-  ],
-  "attributes": { "bits": 4, "block_size": 128, "gather_axis": 0, "quantize_axis": 1 },
   "attributeConstraints": {
     "bits": { "values": [4, 8] },
     "gather_axis": { "values": [0] },
     "quantize_axis": { "values": [1] }
   },
-  "attributeDescriptions": {
-    "bits": "Bits per quantized value. The schema default is 4; this implementation supports 4 or 8.",
-    "block_size": "Number of values sharing a scale. Defaults to 128 and must be a power of two at least 16.",
-    "gather_axis": "Axis from which values are gathered. This matrix implementation supports the standard default, axis 0.",
-    "quantize_axis": "Axis split into quantization blocks. This matrix implementation supports the standard default, axis 1."
-  },
   "typeConstraints": { "T1": ["uint8"], "T2": ["float32"], "Tind": ["int64"] },
-  "args": {
-    "dataT": { "kind": "tensor", "semantic": "data", "role": "input" },
-    "indicesT": { "kind": "tensor", "semantic": "indices", "role": "input", "dtype": "uint32", "narrowing": "checked" },
-    "scalesT": { "kind": "tensor", "semantic": "scales", "role": "input" },
-    "zeroPointsT": { "kind": "tensor", "semantic": "zero_points", "role": "input", "required": false },
-    "outputT": { "kind": "tensor", "semantic": "output", "role": "output" }
-  },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "bits": "attrs.bits",
     "blockSize": "attrs.block_size",
     "blockSizeOk": "blockSize >= 16 and pow2ceil(blockSize) == blockSize",
-    "outBlocks": "ceilDiv(dim(shapes.output, 1), blockSize)",
     "zeroPointCols": "ceilDiv(outBlocks, 2) if bits == 4 else outBlocks",
     "workgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
-    "commonShapeValid": "blockSizeOk and ranks.data == 2 and ranks.indices == 1 and ranks.scales == 2 and ranks.output == 2 and tensorDtypes.indices == \"uint32\" and tensorDtypes.scales == \"float32\" and tensorDtypes.output == \"float32\" and dim(shapes.output, 0) == dim(shapes.indices, 0) and dim(shapes.scales, 0) == dim(shapes.data, 0) and dim(shapes.scales, 1) == outBlocks",
-    "q4ShapeValid": "commonShapeValid and tensorDtypes.data == \"uint8\" and dim(shapes.output, 1) == dim(shapes.data, 1) * 2",
-    "q8ShapeValid": "commonShapeValid and tensorDtypes.data == \"uint8\" and dim(shapes.output, 1) == dim(shapes.data, 1)",
-    "zeroPointsValid": "present.zeroPointsT and ranks.zero_points == 2 and tensorDtypes.zero_points == \"uint8\" and dim(shapes.zero_points, 0) == dim(shapes.data, 0) and dim(shapes.zero_points, 1) == zeroPointCols",
     "noZeroMode": "not present.zeroPointsT",
     "zeroMode": "zeroPointsValid",
     "workgroupFits": "workgroupSize > 0",
-    "foldedDispatchFits": "ceil(ceil(numel(shapes.output) / device.limits.maxComputeWorkgroupsPerDimension) / workgroupSize) <= device.limits.maxComputeWorkgroupsPerDimension"
   },
-  "tunables": { "WORKGROUP_SIZE": 64 },
-  "bindingSets": {
-    "noZero": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "indexCount", "type": "u32", "value": "dim(shapes.indices, 0)" },
-            { "name": "cols", "type": "u32", "value": "dim(shapes.output, 1)" },
-            { "name": "packedCols", "type": "u32", "value": "dim(shapes.data, 1)" },
-            { "name": "blocks", "type": "u32", "value": "dim(shapes.scales, 1)" },
-            { "name": "blockSize", "type": "u32", "value": "blockSize" },
-            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" }
-          ]
-        }
-      }
-    ],
-    "zero": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "zero_points",
-        "arg": "zeroPointsT",
-        "semantic": "zero_points",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$zeroPointElement"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "indexCount", "type": "u32", "value": "dim(shapes.indices, 0)" },
-            { "name": "cols", "type": "u32", "value": "dim(shapes.output, 1)" },
-            { "name": "packedCols", "type": "u32", "value": "dim(shapes.data, 1)" },
-            { "name": "blocks", "type": "u32", "value": "dim(shapes.scales, 1)" },
-            { "name": "blockSize", "type": "u32", "value": "blockSize" },
-            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" }
-          ]
-        }
-      }
-    ],
-    "noZeroIo": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      }
-    ],
-    "zeroIo": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "zero_points",
-        "arg": "zeroPointsT",
-        "semantic": "zero_points",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$zeroPointElement"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      }
-    ],
-    "q4NoZero": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "indexCount", "type": "u32", "value": "dim(shapes.indices, 0)" },
-            { "name": "packedCols", "type": "u32", "value": "dim(shapes.data, 1)" },
-            { "name": "blocks", "type": "u32", "value": "dim(shapes.scales, 1)" },
-            { "name": "blockSize", "type": "u32", "value": "blockSize" },
-            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" }
-          ]
-        }
-      }
-    ],
-    "q4Zero": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "zero_points",
-        "arg": "zeroPointsT",
-        "semantic": "zero_points",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$zeroPointElement"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "indexCount", "type": "u32", "value": "dim(shapes.indices, 0)" },
-            { "name": "packedCols", "type": "u32", "value": "dim(shapes.data, 1)" },
-            { "name": "blocks", "type": "u32", "value": "dim(shapes.scales, 1)" },
-            { "name": "zeroPointCols", "type": "u32", "value": "zeroPointCols" },
-            { "name": "blockSize", "type": "u32", "value": "blockSize" },
-            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" }
-          ]
-        }
-      }
-    ],
-    "q8Vec4NoZero": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "indexCount", "type": "u32", "value": "dim(shapes.indices, 0)" },
-            { "name": "cols", "type": "u32", "value": "dim(shapes.output, 1)" },
-            { "name": "blocks", "type": "u32", "value": "dim(shapes.scales, 1)" },
-            { "name": "blockSize", "type": "u32", "value": "blockSize" },
-            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" }
-          ]
-        }
-      }
-    ],
-    "q8Vec4Zero": [
-      {
-        "name": "data",
-        "arg": "dataT",
-        "semantic": "data",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$dataElement"
-      },
-      {
-        "name": "indices",
-        "arg": "indicesT",
-        "semantic": "indices",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$indexScalar"
-      },
-      {
-        "name": "scales",
-        "arg": "scalesT",
-        "semantic": "scales",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "zero_points",
-        "arg": "zeroPointsT",
-        "semantic": "zero_points",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$zeroPointElement"
-      },
-      {
-        "name": "output",
-        "arg": "outputT",
-        "semantic": "output",
-        "buffer": { "type": "storage" },
-        "elementType": "$outputElement"
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "indexCount", "type": "u32", "value": "dim(shapes.indices, 0)" },
-            { "name": "cols", "type": "u32", "value": "dim(shapes.output, 1)" },
-            { "name": "blocks", "type": "u32", "value": "dim(shapes.scales, 1)" },
-            { "name": "blockSize", "type": "u32", "value": "blockSize" },
-            { "name": "rows", "type": "u32", "value": "dim(shapes.data, 0)" }
-          ]
-        }
-      }
-    ]
   },
   "variants": [
     {
       "id": "q8_no_zero_vec4",
       "priority": 10,
-      "when": ["q8ShapeValid", "noZeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.output, 1) % 4 == 0", "workgroupFits", "foldedDispatchFits"],
-      "constants": {
         "hasZero": false,
         "scalarTail": false,
-        "workgroupSize": "workgroupSize",
         "dataElement": "\"vec4<u32>\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
@@ -458,10 +116,11 @@
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
-          "bindings": "q8Vec4NoZero",
           "dispatch": {
-            "threads": "dim(shapes.indices, 0) * (dim(shapes.output, 1) / 4)",
-            "workgroupSize": "constants.workgroupSize"
           }
         }
       ]
@@ -469,10 +128,9 @@
     {
       "id": "q4_no_zero_pair",
       "priority": 10,
-      "when": ["q4ShapeValid", "noZeroMode", "bits == 4", "workgroupFits", "foldedDispatchFits"],
-      "constants": {
         "hasZero": false,
-        "workgroupSize": "workgroupSize",
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
@@ -482,10 +140,11 @@
         {
           "id": "main",
           "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
-          "bindings": "q4NoZero",
           "dispatch": {
-            "threads": "dim(shapes.indices, 0) * dim(shapes.data, 1)",
-            "workgroupSize": "constants.workgroupSize"
           }
         }
       ]
@@ -493,10 +152,9 @@
     {
       "id": "q4_zero_pair",
       "priority": 10,
-      "when": ["q4ShapeValid", "zeroMode", "bits == 4", "workgroupFits", "foldedDispatchFits"],
-      "constants": {
         "hasZero": true,
-        "workgroupSize": "workgroupSize",
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
@@ -507,10 +165,11 @@
         {
           "id": "main",
           "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
-          "bindings": "q4Zero",
           "dispatch": {
-            "threads": "dim(shapes.indices, 0) * dim(shapes.data, 1)",
-            "workgroupSize": "constants.workgroupSize"
           }
         }
       ]
@@ -518,11 +177,10 @@
     {
       "id": "q8_zero_vec4",
       "priority": 10,
-      "when": ["q8ShapeValid", "zeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.output, 1) % 4 == 0", "workgroupFits", "foldedDispatchFits"],
-      "constants": {
         "hasZero": true,
         "scalarTail": false,
-        "workgroupSize": "workgroupSize",
         "dataElement": "\"vec4<u32>\"",
         "zeroPointElement": "\"u32\"",
         "indexScalar": "\"u32\"",
@@ -533,10 +191,11 @@
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
-          "bindings": "q8Vec4Zero",
           "dispatch": {
-            "threads": "dim(shapes.indices, 0) * (dim(shapes.output, 1) / 4)",
-            "workgroupSize": "constants.workgroupSize"
           }
         }
       ]
@@ -544,11 +203,10 @@
     {
       "id": "q8_no_zero_tail4",
       "priority": 5,
-      "when": ["q8ShapeValid", "noZeroMode", "bits == 8", "workgroupFits", "foldedDispatchFits"],
-      "constants": {
         "hasZero": false,
         "scalarTail": true,
-        "workgroupSize": "workgroupSize",
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
@@ -558,10 +216,11 @@
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
-          "bindings": "noZero",
           "dispatch": {
-            "threads": "dim(shapes.indices, 0) * ceilDiv(dim(shapes.output, 1), 4)",
-            "workgroupSize": "constants.workgroupSize"
           }
         }
       ]
@@ -569,11 +228,10 @@
     {
       "id": "q8_zero_tail4",
       "priority": 5,
-      "when": ["q8ShapeValid", "zeroMode", "bits == 8", "workgroupFits", "foldedDispatchFits"],
-      "constants": {
         "hasZero": true,
         "scalarTail": true,
-        "workgroupSize": "workgroupSize",
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
@@ -584,10 +242,11 @@
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
-          "bindings": "zero",
           "dispatch": {
-            "threads": "dim(shapes.indices, 0) * ceilDiv(dim(shapes.output, 1), 4)",
-            "workgroupSize": "constants.workgroupSize"
           }
         }
       ]

   "domain": "com.microsoft",
   "name": "GatherBlockQuantized",
   "sinceVersion": 1,
+  "inputs": {
+    "dataT": { "onnx": "data", "dtype": "T1", "rank": 2 },
+    "indicesT": { "onnx": "indices", "dtype": "Tind", "rank": 1, "storage": "uint32", "narrowing": "checked" },
+    "scalesT": { "onnx": "scales", "dtype": "T2", "rank": 2 },
+    "zeroPointsT": { "onnx": "zero_points", "dtype": "T1", "rank": 2, "optional": true }
+  },
+  "outputs": {
+    "outputT": {
+      "onnx": "output",
       "dtype": "T2",
       "rank": 2,
+      "shape": "[dim(shapes.indicesT, 0), dim(shapes.dataT, 1) * (8 / attrs.bits)]"
     }
+  },
+  "attributes": {
+    "bits": { "default": 4 },
+    "block_size": { "default": 128 },
+    "gather_axis": { "default": 0 },
+    "quantize_axis": { "default": 1 }
+  },
   "attributeConstraints": {
     "bits": { "values": [4, 8] },
     "gather_axis": { "values": [0] },
     "quantize_axis": { "values": [1] }
   },
   "typeConstraints": { "T1": ["uint8"], "T2": ["float32"], "Tind": ["int64"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 64 } },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "bits": "attrs.bits",
     "blockSize": "attrs.block_size",
     "blockSizeOk": "blockSize >= 16 and pow2ceil(blockSize) == blockSize",
+    "outBlocks": "ceilDiv(dim(shapes.outputT, 1), blockSize)",
     "zeroPointCols": "ceilDiv(outBlocks, 2) if bits == 4 else outBlocks",
     "workgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "commonShapeValid": "blockSizeOk and ranks.dataT == 2 and ranks.indicesT == 1 and ranks.scalesT == 2 and ranks.outputT == 2 and tensorDtypes.indicesT == \"uint32\" and tensorDtypes.scalesT == \"float32\" and tensorDtypes.outputT == \"float32\" and dim(shapes.outputT, 0) == dim(shapes.indicesT, 0) and dim(shapes.scalesT, 0) == dim(shapes.dataT, 0) and dim(shapes.scalesT, 1) == outBlocks",
+    "q4ShapeValid": "commonShapeValid and tensorDtypes.dataT == \"uint8\" and dim(shapes.outputT, 1) == dim(shapes.dataT, 1) * 2",
+    "q8ShapeValid": "commonShapeValid and tensorDtypes.dataT == \"uint8\" and dim(shapes.outputT, 1) == dim(shapes.dataT, 1)",
+    "zeroPointsValid": "present.zeroPointsT and ranks.zeroPointsT == 2 and tensorDtypes.zeroPointsT == \"uint8\" and dim(shapes.zeroPointsT, 0) == dim(shapes.dataT, 0) and dim(shapes.zeroPointsT, 1) == zeroPointCols",
     "noZeroMode": "not present.zeroPointsT",
     "zeroMode": "zeroPointsValid",
     "workgroupFits": "workgroupSize > 0",
+    "foldedDispatchFits": "ceil(ceil(numel(shapes.outputT) / min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) / workgroupSize) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"
   },
+  "when": ["foldedDispatchFits", "workgroupFits"],
+  "bindings": {
+    "data": { "arg": "dataT", "buffer": "read-only-storage", "elementType": "$dataElement" },
+    "indices": { "arg": "indicesT", "buffer": "read-only-storage", "elementType": "$indexScalar" },
+    "scales": { "arg": "scalesT", "buffer": "read-only-storage", "elementType": "$scaleScalar" },
+    "output": { "arg": "outputT", "buffer": "storage", "elementType": "$outputElement" },
+    "params": {
+      "buffer": "uniform",
+      "struct": [
+        { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
+        { "name": "cols", "type": "u32", "value": "dim(shapes.outputT, 1)" },
+        { "name": "blocks", "type": "u32", "value": "dim(shapes.scalesT, 1)" },
+        { "name": "blockSize", "type": "u32", "value": "blockSize" },
+        { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
+      ]
+    },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
+        { "name": "packedCols", "type": "u32", "value": "dim(shapes.dataT, 1)" },
+        { "name": "blocks", "type": "u32", "value": "dim(shapes.scalesT, 1)" },
+        { "name": "blockSize", "type": "u32", "value": "blockSize" },
+        { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
+      ]
+    },
+    "zero_points": { "arg": "zeroPointsT", "buffer": "read-only-storage", "elementType": "$zeroPointElement" },
+    "params_3": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
+        { "name": "packedCols", "type": "u32", "value": "dim(shapes.dataT, 1)" },
+        { "name": "blocks", "type": "u32", "value": "dim(shapes.scalesT, 1)" },
+        { "name": "zeroPointCols", "type": "u32", "value": "zeroPointCols" },
+        { "name": "blockSize", "type": "u32", "value": "blockSize" },
+        { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
+      ]
+    },
+    "params_4": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
+        { "name": "cols", "type": "u32", "value": "dim(shapes.outputT, 1)" },
+        { "name": "packedCols", "type": "u32", "value": "dim(shapes.dataT, 1)" },
+        { "name": "blocks", "type": "u32", "value": "dim(shapes.scalesT, 1)" },
+        { "name": "blockSize", "type": "u32", "value": "blockSize" },
+        { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
+      ]
+    }
   },
   "variants": [
     {
       "id": "q8_no_zero_vec4",
       "priority": 10,
+      "when": ["q8ShapeValid", "noZeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.outputT, 1) % 4 == 0"],
+      "derive": {
         "hasZero": false,
         "scalarTail": false,
         "dataElement": "\"vec4<u32>\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
+          "bindings": ["data", "indices", "scales", "output", "params"],
           "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
+            "z": 1
           }
         }
       ]
     {
       "id": "q4_no_zero_pair",
       "priority": 10,
+      "when": ["q4ShapeValid", "noZeroMode", "bits == 4"],
+      "derive": {
         "hasZero": false,
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
         {
           "id": "main",
           "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
+          "bindings": ["data", "indices", "scales", "output", "params_2"],
           "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
+            "z": 1
           }
         }
       ]
     {
       "id": "q4_zero_pair",
       "priority": 10,
+      "when": ["q4ShapeValid", "zeroMode", "bits == 4"],
+      "derive": {
         "hasZero": true,
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
         {
           "id": "main",
           "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
+          "bindings": ["data", "indices", "scales", "zero_points", "output", "params_3"],
           "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
+            "z": 1
           }
         }
       ]
     {
       "id": "q8_zero_vec4",
       "priority": 10,
+      "when": ["q8ShapeValid", "zeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.outputT, 1) % 4 == 0"],
+      "derive": {
         "hasZero": true,
         "scalarTail": false,
         "dataElement": "\"vec4<u32>\"",
         "zeroPointElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
+          "bindings": ["data", "indices", "scales", "zero_points", "output", "params"],
           "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
+            "z": 1
           }
         }
       ]
     {
       "id": "q8_no_zero_tail4",
       "priority": 5,
+      "when": ["q8ShapeValid", "noZeroMode", "bits == 8"],
+      "derive": {
         "hasZero": false,
         "scalarTail": true,
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
+          "bindings": ["data", "indices", "scales", "output", "params_4"],
           "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
+            "z": 1
           }
         }
       ]
     {
       "id": "q8_zero_tail4",
       "priority": 5,
+      "when": ["q8ShapeValid", "zeroMode", "bits == 8"],
+      "derive": {
         "hasZero": true,
         "scalarTail": true,
         "dataElement": "\"u32\"",
         "indexScalar": "\"u32\"",
         "scaleScalar": "\"f32\"",
         {
           "id": "main",
           "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
+          "bindings": ["data", "indices", "scales", "zero_points", "output", "params_4"],
           "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
+            "z": 1
           }
         }
       ]

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,29 @@
 {
   "name": "com.microsoft.GatherBlockQuantized",
-  "id": "_com_microsoft_gatherblockquantized_webgpu_88e1761",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "Pi6+cheiVelf+nINEm8o7zLFULpK/8jCmL8+6Jlp9I8=",
-      "gather-block-quantized-q4-pair.wgsl.jinja": "U+OSY+xymDWP4cMqTJneiOhIv+Sfazov3A7s3gplCXI=",
-      "gather-block-quantized-q8-vec4.wgsl.jinja": "YKplb0sTVxSB/jvuYiQ3HFxNR6sTBp206dp5EEi4TNM=",
-      "manifest.json": "On/e3F5KW0FA0mt8SUJXckCZKr4EZrxPKZS9JrQdvKU=",
-      "test.json": "XEZF0vZ5ZbO6reeoRXjDL1hLXosLscWnzfZjtHIr9aU="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.GatherBlockQuantized" }
 }

 {
   "name": "com.microsoft.GatherBlockQuantized",
+  "id": "_com_microsoft_gatherblockquantized_webgpu_dceac49",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "iQN5nfNOefh08J/5z4Vc63omCfclP0JZggomZracQjI=",
+      "gather-block-quantized-q4-pair.wgsl.jinja": "9LBelD9NWvia8DuYSL/FCXreyJ6mLtpiiddFdbronSM=",
+      "gather-block-quantized-q8-vec4.wgsl.jinja": "3RQDkgbVbh8TiZZzROkXVWxlFTYBgz4IqY27DOvuVyY=",
+      "manifest.json": "3+GjaCq6H0K41bzVOxcw9rg9vf0MVERgqWjUqBgopq4=",
+      "test.json": "+2tkFsiI9bP4xaQI+hw2wX3CbMkCEqUzgTHYrTAL6bg="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "q8_no_zero_vec4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
+      "q4_no_zero_pair": ["gather-block-quantized-q4-pair.wgsl.jinja"],
+      "q4_zero_pair": ["gather-block-quantized-q4-pair.wgsl.jinja"],
+      "q8_zero_vec4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
+      "q8_no_zero_tail4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
+      "q8_zero_tail4": ["gather-block-quantized-q8-vec4.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.GatherBlockQuantized",
   "fixtureArrays": {
     "ort_q8_no_zero_input_dataT": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63],
     "q8_no_zero_multiblock_vec4_cols32_input_dataT": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127]
@@ -32,7 +31,7 @@
       "provenance": {
         "source": "onnxruntime/test/python/transformers/test_cuda_plugin_ep.py",
         "test": "TestCudaPluginEP.test_op_gather_block_quantized",
-        "notes": "Pinned. With zero_points omitted the default zero point is 2^(bits-1) = 128 at 8 bits, because `data` is unsigned storage for signed values offset by the midpoint; the cited test states this and ORT's kernel applies it in the uint8 branch of contrib_ops/cpu/quantization/gather_block_quantized.cc. Values are chosen so every expected element is exact in f32 and the two rows use different scales, so a wrong default shifts each row by a different amount rather than by a common offset."
       },
       "attrs": { "bits": 8, "block_size": 16 },
       "inputs": {
@@ -124,7 +123,7 @@
     {
       "name": "q8_no_zero_wide_vec4_cols256_idx64",
       "provenance": {
-        "notes": "Compact sibling for the q8 wide-column gather benchmark; preserves no-zero-point q8, blockSize=32, vec4-aligned columns, many gathered rows, and multiple scale blocks per source row."
       },
       "attrs": { "bits": 8, "block_size": 32 },
       "inputs": {

 {
   "fixtureArrays": {
     "ort_q8_no_zero_input_dataT": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63],
     "q8_no_zero_multiblock_vec4_cols32_input_dataT": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127]
       "provenance": {
         "source": "onnxruntime/test/python/transformers/test_cuda_plugin_ep.py",
         "test": "TestCudaPluginEP.test_op_gather_block_quantized",
+        "notes": "With zero_points omitted, the default is 2^(bits-1) = 128 at 8 bits: unsigned storage represents signed values offset by the midpoint. Expected values are exact in float32; different row scales make an incorrect zero point shift the two rows by different amounts."
       },
       "attrs": { "bits": 8, "block_size": 16 },
       "inputs": {
     {
       "name": "q8_no_zero_wide_vec4_cols256_idx64",
       "provenance": {
+        "notes": "A compact q8 gather uses block size 32, the default zero point, four-wide aligned columns, many gathered rows, and multiple scale blocks per source row."
       },
       "attrs": { "bits": 8, "block_size": 32 },
       "inputs": {