sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +77 -0
build/webgpu/bench.json +386 -0
build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja +36 -0
build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja +60 -0
build/webgpu/manifest.json +596 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +299 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.GatherBlockQuantized
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 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.
+See the [ONNX Runtime `GatherBlockQuantized` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.GatherBlockQuantized) for the reference semantics.
+## 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
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `bits` | `4` | Bits per quantized value. The schema default is 4; this implementation supports 4 or 8. |
+| `block_size` | `128` | Number of values sharing a scale. Defaults to 128 and must be a power of two at least 16. |
+| `gather_axis` | `0` | Axis from which values are gathered. This matrix implementation supports the standard default, axis 0. |
+| `quantize_axis` | `1` | Axis split into quantization blocks. This matrix implementation supports the standard default, axis 1. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T1` | `uint8` |
+| `T2` | `float32` |
+| `Tind` | `int64` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`gather-block-quantized-q4-pair.wgsl.jinja`](build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja)
+- [`gather-block-quantized-q8-vec4.wgsl.jinja`](build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja)
+## 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.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.GatherBlockQuantized", { version: 1 });
+const { outputT } = await kernel({
+  dataT: { data: dataTData, shape: [4, 8] },
+  indicesT: { data: indicesTData, shape: [2] },
+  scalesT: { data: scalesTData, shape: [4, 1] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,386 @@

+{
+  "op": "com.microsoft.GatherBlockQuantized",
+  "tunableSpace": { "workgroupSize": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "gather-block-q8-4096x1024-idx1024",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 1024, "indexCount": 1024, "bits": 8, "blockSize": 32 },
+      "attrs": { "bits": 8, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [4096, 1024], "dtype": "uint8", "dist": "randint", "seed": 205, "min": 0, "max": 255 },
+        "indicesT": { "shape": [1024], "dtype": "uint32", "dist": "linearMod", "seed": 205, "step": 37, "mod": 4096 },
+        "scalesT": {
+          "shape": [4096, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 206,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "outputT": { "shape": [1024, 1024], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q4-512x128-idx512",
+      "preset": "smoke",
+      "vars": { "rows": 512, "cols": 128, "indexCount": 512, "bits": 4, "blockSize": 32 },
+      "attrs": { "bits": 4, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [512, 64], "dtype": "uint8", "dist": "q4pair", "seed": 207 },
+        "indicesT": { "shape": [512], "dtype": "uint32", "dist": "linearMod", "seed": 207, "step": 17, "mod": 512 },
+        "scalesT": {
+          "shape": [512, 4],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 208,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "outputT": { "shape": [512, 128], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q4-zero-512x128-idx512",
+      "preset": "smoke",
+      "vars": { "rows": 512, "cols": 128, "indexCount": 512, "bits": 4, "blockSize": 32 },
+      "attrs": { "bits": 4, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [512, 64], "dtype": "uint8", "dist": "q4pair", "seed": 208 },
+        "indicesT": { "shape": [512], "dtype": "uint32", "dist": "linearMod", "seed": 208, "step": 17, "mod": 512 },
+        "scalesT": {
+          "shape": [512, 4],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 209,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": { "shape": [512, 2], "dtype": "uint8", "dist": "q4pair", "seed": 210 }
+      },
+      "outputs": { "outputT": { "shape": [512, 128], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + dim(shapes.zeroPointsT, 1) + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q4-4096x1024-idx1024",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 1024, "indexCount": 1024, "bits": 4, "blockSize": 32 },
+      "attrs": { "bits": 4, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [4096, 512], "dtype": "uint8", "dist": "q4pair", "seed": 211 },
+        "indicesT": { "shape": [1024], "dtype": "uint32", "dist": "linearMod", "seed": 211, "step": 37, "mod": 4096 },
+        "scalesT": {
+          "shape": [4096, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 212,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "outputT": { "shape": [1024, 1024], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q4-zero-4096x1024-idx1024",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 1024, "indexCount": 1024, "bits": 4, "blockSize": 32 },
+      "attrs": { "bits": 4, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [4096, 512], "dtype": "uint8", "dist": "q4pair", "seed": 212 },
+        "indicesT": { "shape": [1024], "dtype": "uint32", "dist": "linearMod", "seed": 212, "step": 37, "mod": 4096 },
+        "scalesT": {
+          "shape": [4096, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 213,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": { "shape": [4096, 16], "dtype": "uint8", "dist": "q4pair", "seed": 214 }
+      },
+      "outputs": { "outputT": { "shape": [1024, 1024], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + dim(shapes.zeroPointsT, 1) + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q8-alignment-cliff-cols1023-scalar-fallback",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 1023, "indexCount": 1024, "bits": 8, "blockSize": 32 },
+      "attrs": { "bits": 8, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [4096, 1023], "dtype": "uint8", "dist": "randint", "seed": 215, "min": 0, "max": 255 },
+        "indicesT": { "shape": [1024], "dtype": "uint32", "dist": "linearMod", "seed": 215, "step": 37, "mod": 4096 },
+        "scalesT": {
+          "shape": [4096, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 216,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "outputT": { "shape": [1024, 1023], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q8-alignment-healthy-cols1024-vec4",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "cols": 1024, "indexCount": 1024, "bits": 8, "blockSize": 32 },
+      "attrs": { "bits": 8, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [4096, 1024], "dtype": "uint8", "dist": "randint", "seed": 215, "min": 0, "max": 255 },
+        "indicesT": { "shape": [1024], "dtype": "uint32", "dist": "linearMod", "seed": 215, "step": 37, "mod": 4096 },
+        "scalesT": {
+          "shape": [4096, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 216,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "outputT": { "shape": [1024, 1024], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q8-dispatch-cliff-vocab-idx8192-cols4096-2dfold",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Stress-only dispatch/capacity case: widened uint8 GPU storage brings the declared footprint to 392 MiB."
+      },
+      "vars": { "rows": 16384, "cols": 4096, "indexCount": 8192, "bits": 8, "blockSize": 32 },
+      "attrs": { "bits": 8, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [16384, 4096], "dtype": "uint8", "dist": "randint", "seed": 217, "min": 0, "max": 255 },
+        "indicesT": { "shape": [8192], "dtype": "uint32", "dist": "linearMod", "seed": 217, "step": 37, "mod": 16384 },
+        "scalesT": {
+          "shape": [16384, 128],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 218,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "outputT": { "shape": [8192, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "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 },
+      "inputs": {
+        "dataT": { "shape": [16384, 4096], "dtype": "uint8", "dist": "randint", "seed": 217, "min": 0, "max": 255 },
+        "indicesT": { "shape": [1024], "dtype": "uint32", "dist": "linearMod", "seed": 217, "step": 37, "mod": 16384 },
+        "scalesT": {
+          "shape": [16384, 128],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 218,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "outputT": { "shape": [1024, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q4-llama-v32000-h4096-idx1",
+      "preset": "model",
+      "provenance": {
+        "notes": "Llama class defaults (vocab_size 32000, hidden_size 4096) with a 4-bit block-quantized embedding table, gathering the single token of a decode step."
+      },
+      "vars": { "rows": 32000, "cols": 4096, "indexCount": 1, "bits": 4, "blockSize": 128 },
+      "attrs": { "bits": 4, "block_size": 128 },
+      "inputs": {
+        "dataT": { "shape": [32000, 2048], "dtype": "uint8", "dist": "uniform", "seed": 6200, "min": 0, "max": 255 },
+        "indicesT": { "shape": [1], "dtype": "uint32", "dist": "randint", "seed": 6201, "min": 0, "max": 31999 },
+        "scalesT": {
+          "shape": [32000, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 6202,
+          "min": 0.005,
+          "max": 0.05
+        }
+      },
+      "outputs": { "outputT": { "shape": [1, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q4-llama-v32000-h4096-idx512",
+      "preset": "model",
+      "provenance": { "notes": "Llama class defaults gathering a 512-token prefill chunk." },
+      "vars": { "rows": 32000, "cols": 4096, "indexCount": 512, "bits": 4, "blockSize": 128 },
+      "attrs": { "bits": 4, "block_size": 128 },
+      "inputs": {
+        "dataT": { "shape": [32000, 2048], "dtype": "uint8", "dist": "uniform", "seed": 6300, "min": 0, "max": 255 },
+        "indicesT": { "shape": [512], "dtype": "uint32", "dist": "randint", "seed": 6301, "min": 0, "max": 31999 },
+        "scalesT": {
+          "shape": [32000, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 6302,
+          "min": 0.005,
+          "max": 0.05
+        }
+      },
+      "outputs": { "outputT": { "shape": [512, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q4-qwen3-v151936-h2048-idx512",
+      "preset": "model",
+      "provenance": {
+        "notes": "Qwen3-MoE class defaults (vocab_size 151936, hidden_size 2048) -- the large-vocabulary case, where the gather is scattered over a much taller table."
+      },
+      "vars": { "rows": 151936, "cols": 2048, "indexCount": 512, "bits": 4, "blockSize": 128 },
+      "attrs": { "bits": 4, "block_size": 128 },
+      "inputs": {
+        "dataT": { "shape": [151936, 1024], "dtype": "uint8", "dist": "uniform", "seed": 6400, "min": 0, "max": 255 },
+        "indicesT": { "shape": [512], "dtype": "uint32", "dist": "randint", "seed": 6401, "min": 0, "max": 151935 },
+        "scalesT": {
+          "shape": [151936, 16],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 6402,
+          "min": 0.005,
+          "max": 0.05
+        }
+      },
+      "outputs": { "outputT": { "shape": [512, 2048], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "gather-block-q8-whisper-v51865-h384-idx448",
+      "preset": "model",
+      "provenance": {
+        "notes": "Whisper class defaults (vocab_size 51865, d_model 384, max_target_positions 448) with an 8-bit table, gathering a full decoder context."
+      },
+      "vars": { "rows": 51865, "cols": 384, "indexCount": 448, "bits": 8, "blockSize": 32 },
+      "attrs": { "bits": 8, "block_size": 32 },
+      "inputs": {
+        "dataT": { "shape": [51865, 384], "dtype": "uint8", "dist": "uniform", "seed": 6500, "min": 0, "max": 255 },
+        "indicesT": { "shape": [448], "dtype": "uint32", "dist": "randint", "seed": 6501, "min": 0, "max": 51864 },
+        "scalesT": {
+          "shape": [51865, 12],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 6502,
+          "min": 0.005,
+          "max": 0.05
+        }
+      },
+      "outputs": { "outputT": { "shape": [448, 384], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.indexCount * (args.cols * args.bits / 8 + dim(shapes.scalesT, 1) * 4 + args.cols * 4 + 4)"
+          }
+        ]
+      }
+    }
+  ]
+}

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

	@@ -0,0 +1,36 @@

+{{ env.wgsl.resourceDeclarations }}
+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;
+  }
+  let out_row = pair_index / params.packedCols;
+  let packed_col = pair_index % params.packedCols;
+  let data_row = indices[out_row];
+  if (data_row >= params.rows) {
+    return;
+  }
+  let packed = data[data_row * params.packedCols + packed_col];
+  let col0 = packed_col * 2u;
+  let block = col0 / params.blockSize;
+  let scale = scales[data_row * params.blocks + block];
+{% if hasZero %}
+  let packed_zero = zero_points[data_row * params.zeroPointCols + block / 2u];
+  let zero = f32((packed_zero >> ((block % 2u) * 4u)) & 15u);
+{% else %}
+  // uint8 data with bits=4 defaults to zero point 2^(bits-1) = 8.
+  let zero = 8.0;
+{% endif %}
+  let q0 = f32(packed & 15u);
+  let q1 = f32((packed >> 4u) & 15u);
+  output[pair_index] = vec2<f32>((q0 - zero) * scale, (q1 - zero) * scale);
+}

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

	@@ -0,0 +1,60 @@

+{{ env.wgsl.resourceDeclarations }}
+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 %}
+  let row_vecs = params.cols / 4u;
+{% endif %}
+  let total = params.indexCount * row_vecs;
+  if (vec_index >= total) {
+    return;
+  }
+  let out_row = vec_index / row_vecs;
+  let vec_col = vec_index % row_vecs;
+  let data_row = indices[out_row];
+  let col0 = vec_col * 4u;
+  if (data_row >= params.rows) {
+    return;
+  }
+{% if scalarTail %}
+{% for lane in range(4) %}
+  if (col0 + {{ lane }}u < params.cols) {
+    let col{{ lane }} = col0 + {{ lane }}u;
+    let block{{ lane }} = col{{ lane }} / params.blockSize;
+    let scale{{ lane }} = scales[data_row * params.blocks + block{{ lane }}];
+    let q{{ lane }} = data[data_row * params.packedCols + col{{ lane }}] & 255u;
+{% if hasZero %}
+    let zero{{ lane }} = f32(zero_points[data_row * params.blocks + block{{ lane }}]);
+    output[out_row * params.cols + col{{ lane }}] = (f32(q{{ lane }}) - zero{{ lane }}) * scale{{ lane }};
+{% else %}
+    // zero_points omitted: the default is 2^(bits-1) = 128 at 8 bits, because `data`
+    // is unsigned storage for signed values offset by the midpoint. Dropping the term
+    // would shift every dequantized value by 128 * scale.
+    output[out_row * params.cols + col{{ lane }}] = (f32(q{{ lane }}) - 128.0) * scale{{ lane }};
+{% endif %}
+  }
+{% endfor %}
+{% else %}
+  let block = col0 / params.blockSize;
+  let scale = scales[data_row * params.blocks + block];
+  let q = data[data_row * row_vecs + vec_col] & vec4<u32>(255u);
+  // blockSize % 4 == 0 (gated) guarantees the 4 packed columns share one block,
+  // so one scale and one zero point apply to the whole vec4.
+{% if hasZero %}
+  let zero = f32(zero_points[data_row * params.blocks + block]);
+  output[vec_index] = (vec4<f32>(q) - vec4<f32>(zero)) * vec4<f32>(scale);
+{% else %}
+  // zero_points omitted: default 2^(bits-1) = 128 at 8 bits (see the scalar arm).
+  output[vec_index] = (vec4<f32>(q) - vec4<f32>(128.0)) * vec4<f32>(scale);
+{% endif %}
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,596 @@

+{
+  "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\"",
+        "outputElement": "\"vec4<f32>\""
+      },
+      "passes": [
+        {
+          "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"
+          }
+        }
+      ]
+    },
+    {
+      "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\"",
+        "outputElement": "\"vec2<f32>\""
+      },
+      "passes": [
+        {
+          "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"
+          }
+        }
+      ]
+    },
+    {
+      "id": "q4_zero_pair",
+      "priority": 10,
+      "when": ["q4ShapeValid", "zeroMode", "bits == 4", "workgroupFits", "foldedDispatchFits"],
+      "constants": {
+        "hasZero": true,
+        "workgroupSize": "workgroupSize",
+        "dataElement": "\"u32\"",
+        "indexScalar": "\"u32\"",
+        "scaleScalar": "\"f32\"",
+        "outputElement": "\"vec2<f32>\"",
+        "zeroPointElement": "\"u32\""
+      },
+      "passes": [
+        {
+          "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"
+          }
+        }
+      ]
+    },
+    {
+      "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\"",
+        "scaleScalar": "\"f32\"",
+        "outputElement": "\"vec4<f32>\""
+      },
+      "passes": [
+        {
+          "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"
+          }
+        }
+      ]
+    },
+    {
+      "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\"",
+        "outputElement": "\"f32\""
+      },
+      "passes": [
+        {
+          "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"
+          }
+        }
+      ]
+    },
+    {
+      "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\"",
+        "outputElement": "\"f32\"",
+        "zeroPointElement": "\"u32\""
+      },
+      "passes": [
+        {
+          "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"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "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" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,299 @@

+{
+  "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]
+  },
+  "cases": [
+    {
+      "name": "ort_q8_no_zero",
+      "provenance": {
+        "source": "onnxruntime/test/python/transformers/test_cuda_plugin_ep.py",
+        "test": "TestCudaPluginEP.test_op_gather_block_quantized"
+      },
+      "attrs": { "bits": 8, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [4, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_q8_no_zero_input_dataT" } }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [0, 2] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4, 1],
+          "data": { "kind": "values", "values": [0.01, 0.02, 0.03, 0.04] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 16], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_q8_no_zero_points_default_midpoint",
+      "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": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [2, 16],
+          "data": {
+            "kind": "values",
+            "values": [120, 121, 122, 123, 124, 125, 126, 127, 128, 129, 130, 131, 132, 133, 134, 135, 128, 130, 132, 134, 136, 138, 140, 142, 144, 146, 148, 150, 152, 154, 156, 158]
+          }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "scalesT": { "dtype": "float32", "shape": [2, 1], "data": { "kind": "values", "values": [0.5, 0.25] } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [2, 16],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [-4.0, -3.5, -3.0, -2.5, -2.0, -1.5, -1.0, -0.5, 0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0, 5.5, 6.0, 6.5, 7.0, 7.5]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_projection_q4_zero_points_block16_tail_rows",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc",
+        "test": "GatherBlockQuantizedOpTest.GatherAxis0WithZeroPoints_4Bits",
+        "notes": "Rank-2 projection of ORT's uint8 q4 gather-axis-0 case with block_size=16, packed tail block, and three gathered rows."
+      },
+      "attrs": { "bits": 4, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [6, 9],
+          "data": {
+            "kind": "values",
+            "values": [16, 50, 16, 50, 16, 50, 16, 50, 128, 84, 118, 84, 118, 84, 118, 84, 118, 132, 152, 186, 152, 186, 152, 186, 152, 186, 136, 220, 254, 220, 254, 220, 254, 220, 254, 140, 220, 254, 220, 254, 220, 254, 220, 254, 140, 84, 118, 84, 118, 84, 118, 84, 118, 132]
+          }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [3, 4, 5] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 1.0, 2.0, 1.0, 2.0, 2.0, 2.0, 1.0, 1.0, 2.0, 1.0] }
+        },
+        "zeroPointsT": {
+          "dtype": "uint8",
+          "shape": [6, 1],
+          "data": { "kind": "values", "values": [151, 136, 121, 121, 137, 151] }
+        }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [3, 18],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [6.0, 8.0, 10.0, 12.0, 6.0, 8.0, 10.0, 12.0, 6.0, 8.0, 10.0, 12.0, 6.0, 8.0, 10.0, 12.0, 10.0, 2.0, 3.0, 4.0, 5.0, 6.0, 3.0, 4.0, 5.0, 6.0, 3.0, 4.0, 5.0, 6.0, 3.0, 4.0, 5.0, 6.0, 4.0, 0.0, -6.0, -4.0, -2.0, 0.0, -6.0, -4.0, -2.0, 0.0, -6.0, -4.0, -2.0, 0.0, -6.0, -4.0, -2.0, 0.0, -5.0, -1.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "q8_no_zero_multiblock_vec4_cols32",
+      "attrs": { "bits": 8, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [4, 32],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/q8_no_zero_multiblock_vec4_cols32_input_dataT" }
+          }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 3, 1] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.01, 0.02, 0.03, 0.04, 0.05, 0.06, 0.07, 0.08] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 32], "tolerance": 0.000001 } }
+    },
+    {
+      "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": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [256, 256],
+          "data": { "kind": "cycle", "values": [0, 1, 2, 3, 4, 5, 31, 63, 127, 191, 255] }
+        },
+        "indicesT": {
+          "dtype": "uint32",
+          "shape": [64],
+          "data": { "kind": "cycle", "values": [0, 17, 63, 128, 255, 3, 42, 191] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [256, 8],
+          "data": { "kind": "cycle", "values": [0.01, 0.02, 0.03, 0.04, 0.05, 0.06, 0.07, 0.08] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [64, 256], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "q4_zero_points_block16_ort_valid",
+      "attrs": { "bits": 4, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [4, 16],
+          "data": {
+            "kind": "values",
+            "values": [16, 50, 84, 118, 152, 186, 220, 254, 16, 33, 136, 119, 102, 85, 17, 34, 200, 152, 104, 56, 8, 216, 168, 120, 72, 24, 232, 184, 136, 88, 40, 248, 35, 70, 105, 140, 175, 210, 245, 21, 56, 91, 126, 161, 196, 231, 11, 46, 81, 116, 151, 186, 221, 0, 35, 70, 105, 140, 175, 210, 245, 21, 56, 91]
+          }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [3, 0, 2] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.05, 0.125, 0.25, 0.15, 0.075, 0.3] }
+        },
+        "zeroPointsT": {
+          "dtype": "uint8",
+          "shape": [4, 1],
+          "data": { "kind": "values", "values": [120, 105, 135, 149] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 32], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "q8_with_zero_points_vec4_cols32",
+      "attrs": { "bits": 8, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [4, 32],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/q8_no_zero_multiblock_vec4_cols32_input_dataT" }
+          }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 3, 1] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [0.01, 0.02, 0.03, 0.04, 0.05, 0.06, 0.07, 0.08] }
+        },
+        "zeroPointsT": {
+          "dtype": "uint8",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [128, 100, 64, 90, 110, 130, 70, 120] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 32], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "attrs": { "bits": 8, "block_size": 16 },
+      "inputs": {
+        "dataT": { "dtype": "uint8", "shape": [0, 16], "data": { "kind": "values", "values": [] } },
+        "indicesT": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "scalesT": { "dtype": "float32", "shape": [0, 1], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [0, 16], "tolerance": 0 } }
+    },
+    {
+      "name": "q8_empty_indices_populated_data",
+      "attrs": { "bits": 8, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [4, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_q8_no_zero_input_dataT" } }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4, 1],
+          "data": { "kind": "values", "values": [0.01, 0.02, 0.03, 0.04] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [0, 16], "tolerance": 0 } }
+    },
+    {
+      "name": "q8_max_valid_index_and_duplicates_vec4",
+      "attrs": { "bits": 8, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [4, 8],
+          "data": {
+            "kind": "values",
+            "values": [0, 1, 2, 3, 4, 5, 6, 7, 10, 20, 30, 40, 50, 60, 70, 80, 255, 128, 64, 32, 16, 8, 4, 2, 90, 100, 110, 120, 130, 140, 150, 160]
+          }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [3, 3, 0, 3, 1] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4, 1],
+          "data": { "kind": "values", "values": [0.01, 0.02, 0.03, 0.04] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [5, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_default_bits_and_block_size_q4",
+      "provenance": {
+        "source": "onnxruntime/core/graph/contrib_ops/contrib_defs.cc",
+        "test": "GatherBlockQuantized schema attribute defaults",
+        "notes": "Omitting `bits` and `block_size` exercises their standard defaults of 4 and 128."
+      },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [4, 8],
+          "data": { "kind": "cycle", "values": [7, 12, 3, 9, 14, 1, 5, 11] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4, 1],
+          "data": { "kind": "values", "values": [0.5, 0.25, 1.5, 2.0] }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 0] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 16], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "q8_zero_points_tail_cols18",
+      "provenance": {
+        "notes": "Explicit q8 zero points with 18 output columns exercise the scalar tail after four complete vec4 groups; the final quantization block is partial and uses its own scale and zero point."
+      },
+      "attrs": { "bits": 8, "block_size": 16 },
+      "inputs": {
+        "dataT": {
+          "dtype": "uint8",
+          "shape": [3, 18],
+          "data": { "kind": "cycle", "values": [0, 17, 64, 99, 128, 143, 191, 255, 37, 211] }
+        },
+        "indicesT": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 0, 2] } },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [0.05, 0.2, 0.025, 0.125, 0.075, 0.3] }
+        },
+        "zeroPointsT": {
+          "dtype": "uint8",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [120, 130, 100, 140, 110, 150] }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 18], "tolerance": 0.000001 } }
+    }
+  ]
+}