Hugging Face's logo

Kernels:
webgpu-kernels
/
com.microsoft.GatherBlockQuantized

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
e18dfab
·
verified ·
1 Parent(s): 6dbed36

sync 6fdf6301e2bb

Browse files
Files changed (6) hide show
  1. README.md +2 -2
  2. build/webgpu/bench.json +2 -2
  3. build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja +6 -3
  4. build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja +6 -3
  5. build/webgpu/manifest.json +35 -78
  6. build/webgpu/metadata.json +8 -8
README.md CHANGED
@@ -55,14 +55,14 @@ Default values (overridable per request):
55
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
56
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
57
  - [`test.json`](build/webgpu/test.json) — correctness cases
58
- - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
59
  - [`gather-block-quantized-q4-pair.wgsl.jinja`](build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja)
60
  - [`gather-block-quantized-q8-vec4.wgsl.jinja`](build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja)
61
 
62
  ## Use with `@huggingface/kernels`
63
 
64
  ```sh
65
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
66
  ```
67
 
68
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 
55
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
56
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
57
  - [`test.json`](build/webgpu/test.json) — correctness cases
58
+ - [`bench.json`](build/webgpu/bench.json) — benchmark cases
59
  - [`gather-block-quantized-q4-pair.wgsl.jinja`](build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja)
60
  - [`gather-block-quantized-q8-vec4.wgsl.jinja`](build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja)
61
 
62
  ## Use with `@huggingface/kernels`
63
 
64
  ```sh
65
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
66
  ```
67
 
68
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
build/webgpu/bench.json CHANGED
@@ -174,7 +174,7 @@
174
  }
175
  },
176
  {
177
- "name": "gather-block-q8-alignment-healthy-cols1024-vec4",
178
  "preset": "smoke",
179
  "vars": { "rows": 4096, "cols": 1024, "indexCount": 1024, "bits": 8, "blockSize": 32 },
180
  "attrs": { "bits": 8, "block_size": 32 },
@@ -233,7 +233,7 @@
233
  }
234
  },
235
  {
236
- "name": "gather-block-q8-dispatch-healthy-idx1024-cols4096-1d",
237
  "preset": "stress",
238
  "provenance": {
239
  "notes": "Widened uint8 GPU storage gives this capacity stress case a declared footprint of 280 MiB."
 
174
  }
175
  },
176
  {
177
+ "name": "gather-block-q8-alignment-control-cols1024-vec4",
178
  "preset": "smoke",
179
  "vars": { "rows": 4096, "cols": 1024, "indexCount": 1024, "bits": 8, "blockSize": 32 },
180
  "attrs": { "bits": 8, "block_size": 32 },
 
233
  }
234
  },
235
  {
236
+ "name": "gather-block-q8-dispatch-control-idx1024-cols4096-1d",
237
  "preset": "stress",
238
  "provenance": {
239
  "notes": "Widened uint8 GPU storage gives this capacity stress case a declared footprint of 280 MiB."
build/webgpu/gather-block-quantized-q4-pair.wgsl.jinja CHANGED
@@ -1,12 +1,15 @@
 
 
 
 
 
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
  const WG: u32 = {{ workgroupSize }}u;
4
 
5
  @compute @workgroup_size(WG, 1, 1)
6
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
- // 2D-folded flat index: gid.y carries the high bits past the per-axis dispatch fold width.
8
- // Reduces to gid.x when the dispatch does not fold.
9
- let pair_index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
10
  let total = params.indexCount * params.packedCols;
11
  if (pair_index >= total) {
12
  return;
 
1
+ {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
2
+ {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
3
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
4
+ // per-axis workgroup fold width.
5
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% endmacro %}
6
  {{ env.wgsl.resourceDeclarations }}
7
 
8
  const WG: u32 = {{ workgroupSize }}u;
9
 
10
  @compute @workgroup_size(WG, 1, 1)
11
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
12
+ {{ flat_index_2d("WG", "pair_index", "") }}
 
 
13
  let total = params.indexCount * params.packedCols;
14
  if (pair_index >= total) {
15
  return;
build/webgpu/gather-block-quantized-q8-vec4.wgsl.jinja CHANGED
@@ -1,12 +1,15 @@
 
 
 
 
 
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
  const WG: u32 = {{ workgroupSize }}u;
4
 
5
  @compute @workgroup_size(WG, 1, 1)
6
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
- // 2D-folded flat index: gid.y carries the high bits past the per-axis dispatch fold width.
8
- // Reduces to gid.x when the dispatch does not fold.
9
- let vec_index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
10
  {% if scalarTail %}
11
  let row_vecs = (params.cols + 3u) / 4u;
12
  {% else %}
 
1
+ {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
2
+ {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
3
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
4
+ // per-axis workgroup fold width.
5
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% endmacro %}
6
  {{ env.wgsl.resourceDeclarations }}
7
 
8
  const WG: u32 = {{ workgroupSize }}u;
9
 
10
  @compute @workgroup_size(WG, 1, 1)
11
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
12
+ {{ flat_index_2d("WG", "vec_index", "") }}
 
 
13
  {% if scalarTail %}
14
  let row_vecs = (params.cols + 3u) / 4u;
15
  {% else %}
build/webgpu/manifest.json CHANGED
@@ -43,17 +43,20 @@
43
  "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",
44
  "noZeroMode": "not present.zeroPointsT",
45
  "zeroMode": "zeroPointsValid",
 
 
 
 
46
  "workgroupFits": "workgroupSize > 0",
47
  "foldedDispatchFits": "ceil(ceil(numel(shapes.outputT) / min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) / workgroupSize) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"
48
  },
49
  "when": ["foldedDispatchFits", "workgroupFits"],
50
  "bindings": {
51
- "data": { "arg": "dataT", "buffer": "read-only-storage", "elementType": "$dataElement" },
52
- "indices": { "arg": "indicesT", "buffer": "read-only-storage", "elementType": "$indexScalar" },
53
- "scales": { "arg": "scalesT", "buffer": "read-only-storage", "elementType": "$scaleScalar" },
54
- "output": { "arg": "outputT", "buffer": "storage", "elementType": "$outputElement" },
55
  "params": {
56
- "buffer": "uniform",
57
  "struct": [
58
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
59
  { "name": "cols", "type": "u32", "value": "dim(shapes.outputT, 1)" },
@@ -62,9 +65,8 @@
62
  { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
63
  ]
64
  },
65
- "params_2": {
66
  "name": "params",
67
- "buffer": "uniform",
68
  "struct": [
69
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
70
  { "name": "packedCols", "type": "u32", "value": "dim(shapes.dataT, 1)" },
@@ -73,10 +75,9 @@
73
  { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
74
  ]
75
  },
76
- "zero_points": { "arg": "zeroPointsT", "buffer": "read-only-storage", "elementType": "$zeroPointElement" },
77
- "params_3": {
78
  "name": "params",
79
- "buffer": "uniform",
80
  "struct": [
81
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
82
  { "name": "packedCols", "type": "u32", "value": "dim(shapes.dataT, 1)" },
@@ -86,9 +87,8 @@
86
  { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
87
  ]
88
  },
89
- "params_4": {
90
  "name": "params",
91
- "buffer": "uniform",
92
  "struct": [
93
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
94
  { "name": "cols", "type": "u32", "value": "dim(shapes.outputT, 1)" },
@@ -104,14 +104,7 @@
104
  "id": "q8_no_zero_vec4",
105
  "priority": 10,
106
  "when": ["q8ShapeValid", "noZeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.outputT, 1) % 4 == 0"],
107
- "derive": {
108
- "hasZero": false,
109
- "scalarTail": false,
110
- "dataElement": "\"vec4<u32>\"",
111
- "indexScalar": "\"u32\"",
112
- "scaleScalar": "\"f32\"",
113
- "outputElement": "\"vec4<f32>\""
114
- },
115
  "passes": [
116
  {
117
  "id": "main",
@@ -126,46 +119,33 @@
126
  ]
127
  },
128
  {
129
- "id": "q4_no_zero_pair",
130
  "priority": 10,
131
- "when": ["q4ShapeValid", "noZeroMode", "bits == 4"],
132
- "derive": {
133
- "hasZero": false,
134
- "dataElement": "\"u32\"",
135
- "indexScalar": "\"u32\"",
136
- "scaleScalar": "\"f32\"",
137
- "outputElement": "\"vec2<f32>\""
138
- },
139
  "passes": [
140
  {
141
  "id": "main",
142
- "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
143
- "bindings": ["data", "indices", "scales", "output", "params_2"],
144
  "dispatch": {
145
- "x": "min(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
146
- "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
147
  "z": 1
148
  }
149
  }
150
  ]
151
  },
152
  {
153
- "id": "q4_zero_pair",
154
  "priority": 10,
155
- "when": ["q4ShapeValid", "zeroMode", "bits == 4"],
156
- "derive": {
157
- "hasZero": true,
158
- "dataElement": "\"u32\"",
159
- "indexScalar": "\"u32\"",
160
- "scaleScalar": "\"f32\"",
161
- "outputElement": "\"vec2<f32>\"",
162
- "zeroPointElement": "\"u32\""
163
- },
164
  "passes": [
165
  {
166
  "id": "main",
167
  "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
168
- "bindings": ["data", "indices", "scales", "zero_points", "output", "params_3"],
169
  "dispatch": {
170
  "x": "min(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
171
  "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
@@ -175,26 +155,18 @@
175
  ]
176
  },
177
  {
178
- "id": "q8_zero_vec4",
179
  "priority": 10,
180
- "when": ["q8ShapeValid", "zeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.outputT, 1) % 4 == 0"],
181
- "derive": {
182
- "hasZero": true,
183
- "scalarTail": false,
184
- "dataElement": "\"vec4<u32>\"",
185
- "zeroPointElement": "\"u32\"",
186
- "indexScalar": "\"u32\"",
187
- "scaleScalar": "\"f32\"",
188
- "outputElement": "\"vec4<f32>\""
189
- },
190
  "passes": [
191
  {
192
  "id": "main",
193
- "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
194
- "bindings": ["data", "indices", "scales", "zero_points", "output", "params"],
195
  "dispatch": {
196
- "x": "min(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
197
- "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
198
  "z": 1
199
  }
200
  }
@@ -204,19 +176,12 @@
204
  "id": "q8_no_zero_tail4",
205
  "priority": 5,
206
  "when": ["q8ShapeValid", "noZeroMode", "bits == 8"],
207
- "derive": {
208
- "hasZero": false,
209
- "scalarTail": true,
210
- "dataElement": "\"u32\"",
211
- "indexScalar": "\"u32\"",
212
- "scaleScalar": "\"f32\"",
213
- "outputElement": "\"f32\""
214
- },
215
  "passes": [
216
  {
217
  "id": "main",
218
  "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
219
- "bindings": ["data", "indices", "scales", "output", "params_4"],
220
  "dispatch": {
221
  "x": "min(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
222
  "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
@@ -229,20 +194,12 @@
229
  "id": "q8_zero_tail4",
230
  "priority": 5,
231
  "when": ["q8ShapeValid", "zeroMode", "bits == 8"],
232
- "derive": {
233
- "hasZero": true,
234
- "scalarTail": true,
235
- "dataElement": "\"u32\"",
236
- "indexScalar": "\"u32\"",
237
- "scaleScalar": "\"f32\"",
238
- "outputElement": "\"f32\"",
239
- "zeroPointElement": "\"u32\""
240
- },
241
  "passes": [
242
  {
243
  "id": "main",
244
  "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
245
- "bindings": ["data", "indices", "scales", "zero_points", "output", "params_4"],
246
  "dispatch": {
247
  "x": "min(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
248
  "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
 
43
  "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",
44
  "noZeroMode": "not present.zeroPointsT",
45
  "zeroMode": "zeroPointsValid",
46
+ "hasZero": "zeroMode",
47
+ "indexScalar": "\"u32\"",
48
+ "scaleScalar": "\"f32\"",
49
+ "zeroPointElement": "\"u32\"",
50
  "workgroupFits": "workgroupSize > 0",
51
  "foldedDispatchFits": "ceil(ceil(numel(shapes.outputT) / min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) / workgroupSize) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"
52
  },
53
  "when": ["foldedDispatchFits", "workgroupFits"],
54
  "bindings": {
55
+ "data": { "arg": "dataT", "elementType": "$dataElement" },
56
+ "indices": { "arg": "indicesT", "elementType": "$indexScalar" },
57
+ "scales": { "arg": "scalesT", "elementType": "$scaleScalar" },
58
+ "output": { "arg": "outputT", "elementType": "$outputElement" },
59
  "params": {
 
60
  "struct": [
61
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
62
  { "name": "cols", "type": "u32", "value": "dim(shapes.outputT, 1)" },
 
65
  { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
66
  ]
67
  },
68
+ "params_main": {
69
  "name": "params",
 
70
  "struct": [
71
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
72
  { "name": "packedCols", "type": "u32", "value": "dim(shapes.dataT, 1)" },
 
75
  { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
76
  ]
77
  },
78
+ "zero_points": { "arg": "zeroPointsT", "elementType": "$zeroPointElement" },
79
+ "params__uniform": {
80
  "name": "params",
 
81
  "struct": [
82
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
83
  { "name": "packedCols", "type": "u32", "value": "dim(shapes.dataT, 1)" },
 
87
  { "name": "rows", "type": "u32", "value": "dim(shapes.dataT, 0)" }
88
  ]
89
  },
90
+ "params_q8_zero_tail4": {
91
  "name": "params",
 
92
  "struct": [
93
  { "name": "indexCount", "type": "u32", "value": "dim(shapes.indicesT, 0)" },
94
  { "name": "cols", "type": "u32", "value": "dim(shapes.outputT, 1)" },
 
104
  "id": "q8_no_zero_vec4",
105
  "priority": 10,
106
  "when": ["q8ShapeValid", "noZeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.outputT, 1) % 4 == 0"],
107
+ "derive": { "scalarTail": false, "dataElement": "\"vec4<u32>\"", "outputElement": "\"vec4<f32>\"" },
 
 
 
 
 
 
 
108
  "passes": [
109
  {
110
  "id": "main",
 
119
  ]
120
  },
121
  {
122
+ "id": "q8_zero_vec4",
123
  "priority": 10,
124
+ "when": ["q8ShapeValid", "zeroMode", "bits == 8", "blockSize % 4 == 0", "dim(shapes.outputT, 1) % 4 == 0"],
125
+ "derive": { "scalarTail": false, "dataElement": "\"vec4<u32>\"", "outputElement": "\"vec4<f32>\"" },
 
 
 
 
 
 
126
  "passes": [
127
  {
128
  "id": "main",
129
+ "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
130
+ "bindings": ["data", "indices", "scales", "zero_points", "output", "params"],
131
  "dispatch": {
132
+ "x": "min(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
133
+ "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * (dim(shapes.outputT, 1) / 4)), (workgroupSize)), 65535)",
134
  "z": 1
135
  }
136
  }
137
  ]
138
  },
139
  {
140
+ "id": "q4_no_zero_pair",
141
  "priority": 10,
142
+ "when": ["q4ShapeValid", "noZeroMode", "bits == 4"],
143
+ "derive": { "dataElement": "\"u32\"", "outputElement": "\"vec2<f32>\"" },
 
 
 
 
 
 
 
144
  "passes": [
145
  {
146
  "id": "main",
147
  "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
148
+ "bindings": ["data", "indices", "scales", "output", "params_main"],
149
  "dispatch": {
150
  "x": "min(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
151
  "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
 
155
  ]
156
  },
157
  {
158
+ "id": "q4_zero_pair",
159
  "priority": 10,
160
+ "when": ["q4ShapeValid", "zeroMode", "bits == 4"],
161
+ "derive": { "dataElement": "\"u32\"", "outputElement": "\"vec2<f32>\"" },
 
 
 
 
 
 
 
 
162
  "passes": [
163
  {
164
  "id": "main",
165
+ "shader": "gather-block-quantized-q4-pair.wgsl.jinja",
166
+ "bindings": ["data", "indices", "scales", "zero_points", "output", "params__uniform"],
167
  "dispatch": {
168
+ "x": "min(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
169
+ "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * dim(shapes.dataT, 1)), (workgroupSize)), 65535)",
170
  "z": 1
171
  }
172
  }
 
176
  "id": "q8_no_zero_tail4",
177
  "priority": 5,
178
  "when": ["q8ShapeValid", "noZeroMode", "bits == 8"],
179
+ "derive": { "scalarTail": true, "dataElement": "\"u32\"", "outputElement": "\"f32\"" },
 
 
 
 
 
 
 
180
  "passes": [
181
  {
182
  "id": "main",
183
  "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
184
+ "bindings": ["data", "indices", "scales", "output", "params_q8_zero_tail4"],
185
  "dispatch": {
186
  "x": "min(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
187
  "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
 
194
  "id": "q8_zero_tail4",
195
  "priority": 5,
196
  "when": ["q8ShapeValid", "zeroMode", "bits == 8"],
197
+ "derive": { "scalarTail": true, "dataElement": "\"u32\"", "outputElement": "\"f32\"" },
 
 
 
 
 
 
 
 
198
  "passes": [
199
  {
200
  "id": "main",
201
  "shader": "gather-block-quantized-q8-vec4.wgsl.jinja",
202
+ "bindings": ["data", "indices", "scales", "zero_points", "output", "params_q8_zero_tail4"],
203
  "dispatch": {
204
  "x": "min(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
205
  "y": "ceilDiv(ceilDiv((dim(shapes.indicesT, 0) * ceilDiv(dim(shapes.outputT, 1), 4)), (workgroupSize)), 65535)",
build/webgpu/metadata.json CHANGED
@@ -1,27 +1,27 @@
1
  {
2
  "name": "com.microsoft.GatherBlockQuantized",
3
- "id": "_com_microsoft_gatherblockquantized_webgpu_dceac49",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "iQN5nfNOefh08J/5z4Vc63omCfclP0JZggomZracQjI=",
11
- "gather-block-quantized-q4-pair.wgsl.jinja": "9LBelD9NWvia8DuYSL/FCXreyJ6mLtpiiddFdbronSM=",
12
- "gather-block-quantized-q8-vec4.wgsl.jinja": "3RQDkgbVbh8TiZZzROkXVWxlFTYBgz4IqY27DOvuVyY=",
13
- "manifest.json": "3+GjaCq6H0K41bzVOxcw9rg9vf0MVERgqWjUqBgopq4=",
14
  "test.json": "+2tkFsiI9bP4xaQI+hw2wX3CbMkCEqUzgTHYrTAL6bg="
15
  }
16
  },
17
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
18
  "webgpu": {
19
- "manifestSpec": "2.0",
20
  "variants": {
21
  "q8_no_zero_vec4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
 
22
  "q4_no_zero_pair": ["gather-block-quantized-q4-pair.wgsl.jinja"],
23
  "q4_zero_pair": ["gather-block-quantized-q4-pair.wgsl.jinja"],
24
- "q8_zero_vec4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
25
  "q8_no_zero_tail4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
26
  "q8_zero_tail4": ["gather-block-quantized-q8-vec4.wgsl.jinja"]
27
  }
 
1
  {
2
  "name": "com.microsoft.GatherBlockQuantized",
3
+ "id": "_com_microsoft_gatherblockquantized_webgpu_7698e56",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "jKDPiApwDuwewTrNn1DOYD7gMVHY+vME4qUAe2kUbsw=",
11
+ "gather-block-quantized-q4-pair.wgsl.jinja": "szEkOKfzt9XFHiIsEiQY0g5efbe36fXYBCE3v4SVo6s=",
12
+ "gather-block-quantized-q8-vec4.wgsl.jinja": "akD5Z/+o9Q/HRnbH8yPrD8Y6Pd0SWPgiDRwgHEqt/YQ=",
13
+ "manifest.json": "1/OD0zneUAu2ddg/neSEu+dpmmWqZiMNA9AVO6AI1W8=",
14
  "test.json": "+2tkFsiI9bP4xaQI+hw2wX3CbMkCEqUzgTHYrTAL6bg="
15
  }
16
  },
17
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
18
  "webgpu": {
19
+ "manifestSpec": "2.1",
20
  "variants": {
21
  "q8_no_zero_vec4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
22
+ "q8_zero_vec4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
23
  "q4_no_zero_pair": ["gather-block-quantized-q4-pair.wgsl.jinja"],
24
  "q4_zero_pair": ["gather-block-quantized-q4-pair.wgsl.jinja"],
 
25
  "q8_no_zero_tail4": ["gather-block-quantized-q8-vec4.wgsl.jinja"],
26
  "q8_zero_tail4": ["gather-block-quantized-q8-vec4.wgsl.jinja"]
27
  }