【Bug已解决】GatherBlockQuantized: invalid dispatch group size (0,1,1) on macOS Metal WebGPU 解决方案

一、现象长什么样

在macOS 的 Metal 后端（WebGPU）上跑GatherBlockQuantized算子时，提交计算着色器（compute shader）直接报错：

Metal: invalid dispatch threadgroups size (0, 1, 1) # 或 Validation layer: dispatchThreadgroups called with width=0

具体表现：

• 只在macOS / Metal上出现；Windows 的 D3D12 / 其它 WebGPU 后端可能不报（有的后端对 0 维更宽松或静默处理）。
• 只在特定量化块大小（block size）或权重维度组合下触发——通常是某维算出 0 个线程组时。
• 报错明确指向“dispatch group size 的某个维度是 0”，而合法值必须>= 1。
• 推理在该算子处中断，后续全废。

关键特征：着色器派发（dispatch）的线程组数量被算成了 0，而 Metal 要求每个维度至少为 1。这是“根据输入/块大小推算 dispatch 尺寸时，没有对 0 做保护/向上取整”的经典 bug。

二、背景

GPU 计算是“分块并行”的：你把一个大任务切成很多线程组（threadgroup / workgroup），每组再含若干线程。GatherBlockQuantized的 WGSL 计算着色器要把“按块量化的权重”拆解到线程组上处理。派发时，CPU 侧（或 WGSL 的dispatch_workgroups）要算出每个维度需要多少个线程组：

groups_x = ceil(total_items_x / workgroup_size_x) groups_y = ceil(total_items_y / workgroup_size_y) groups_z = ceil(total_items_z / workgroup_size_z)

(0,1,1)这种错误意味着groups_x被算成了 0。为什么会 0？通常因为：

• 某个维度上的“总项数”为 0（比如权重张量该维长度为 0，或量化块大小配置让有效项数变成 0）。
• 或者ceil没做对：当total_items_x为 0 时，0 / workgroup_size在某些整数除法下得到 0，而派发维度不允许 0。
• Metal 特别严格：D3D12 或 Vulkan 某些实现可能对 0 维有不同处理，但Metal 的dispatchThreadgroups要求每个维度 >= 1，于是直接校验失败报invalid dispatch group size (0,1,1)。

GatherBlockQuantized因为是“按块”的，块大小（block size）、量化位宽、权重维度三者共同决定线程组数；当其中某个组合让某维项数变成 0 或极小，整数除法就给出 0，触发 Metal 的硬校验。

三、根因

根因是GatherBlockQuantized在派发计算着色器时，没有保证每个 dispatch 维度至少为 1，当某维总项数为 0 / 极小导致整数除法得 0 时，Metal 直接拒绝：

1. dispatch 维度可能为 0：groups = total / workgroup_size用整数除法，当total == 0时结果为 0；Metal 要求>= 1，于是校验失败。正确的做法是用ceil并max(1, ...)，且对total == 0这种退化情况特殊处理（跳过派发或至少给 1）。
2. 块大小/维度组合未校验：量化块大小与权重维度组合可能产出一个“0 项”的退化维度，算子没在派发前断言/修正。
3. 跨后端不一致：Metal 严格校验 0，其它后端宽松，导致“只在 macOS 崩”的迷惑现象，容易让人以为是 Metal 驱动问题而非自己算错。
4. 错误定位难：报错只说(0,1,1)，不告诉你哪维、为什么 0，需要回到派发计算处查。

一句话：GatherBlockQuantized 的 dispatch 尺寸计算没对 0 维做max(1,...)保护，退化维度下得出 0，被严格的 Metal 拒掉。

四、最小可运行复现

下面用 Python 模拟“dispatch 组数计算”，复现(0,1,1)的来历与修复：

import math def dispatch_buggy(total_x: int, total_y: int, wg: int = 64) -> tuple: """错误：整数除法，total=0 时得 0。""" gx = total_x // wg gy = total_y // wg gz = 1 return (gx, gy, gz) def dispatch_fixed(total_x: int, total_y: int, wg: int = 64) -> tuple: """修复：ceil 且每维至少 1；极端退化(total=0)直接不派发。""" if total_x == 0 or total_y == 0: return (0, 0, 0) # 调用方应跳过派发 gx = max(1, math.ceil(total_x / wg)) gy = max(1, math.ceil(total_y / wg)) return (gx, gy, 1) # 退化场景：某维总项数为 0 print("buggy :", dispatch_buggy(total_x=0, total_y=128)) # (0, 2, 1) -> Metal 拒 print("fixed :", dispatch_fixed(total_x=0, total_y=128)) # (0, 0, 0) 调用方跳过 # 正常场景 print("buggy normal:", dispatch_buggy(200, 128)) # (3, 2, 1) print("fixed normal:", dispatch_fixed(200, 128)) # (4, 2, 1) ceil 正确

buggy在total_x=0时给出gx=0，正是 Metal 报的(0,1,1)形态；fixed用ceil+max(1,...)，且退化时返回全 0 让调用方跳过派发，避免非法 dispatch。

五、解决方案（第一层：最小直接修复）

最小修复是在派发GatherBlockQuantized着色器前，对每个维度用ceil并max(1, ...)，且对全 0 的退化情况跳过派发：

// webgpu_gather_block_quant_dispatch.cpp（修复片段） void DispatchGatherBlockQuant(CommandEncoder& enc, uint32_t total_x, uint32_t total_y, uint32_t wg_size) { // 退化：没有可处理项，不派发（避免 (0,1,1)） if (total_x == 0 || total_y == 0) { return; } const uint32_t gx = std::max<uint32_t>(1, (total_x + wg_size - 1) / wg_size); const uint32_t gy = std::max<uint32_t>(1, (total_y + wg_size - 1) / wg_size); // Metal 要求每维 >= 1，这里已经保证 enc.DispatchWorkgroups(gx, gy, 1); }

配套：在算子入口校验量化块大小与权重维度，确保不会产出“0 项”退化维度（或明确接受并跳过）。

这一层让 macOS Metal 不再收到(0,1,1)的非法 dispatch，退化情况安全跳过。

六、解决方案（第二层：结构性改进）

把“WebGPU 着色器派发尺寸如何计算、退化如何处理”收口成唯一的配置对象OrtWebGpuDispatchSizePolicy，所有 WebGPU 算子派发读它：

from dataclasses import dataclass from typing import Tuple @dataclass(frozen=True) class OrtWebGpuDispatchSizePolicy: """WebGPU dispatch 尺寸计算的单一事实来源。""" # 每个 dispatch 维度必须 >= 1（Metal 硬要求） min_group_per_dim: int = 1 # 用 ceil 计算组数，不用整数截断除法 use_ceil: bool = True # 全 0 退化维度：跳过派发，不提交非法 dispatch skip_on_zero_total: bool = True # 禁止任何维度为 0 的 dispatch forbid_zero_dim_dispatch: bool = True # 代码评审卡点 forbidden_patterns: Tuple[str, ...] = ( "groups = total // wg_size", "dispatch with dim 0", ) def compute(self, totals: Tuple[int, ...], wg: int) -> Tuple[int, ...]: if self.skip_on_zero_total and any(t == 0 for t in totals): return tuple(0 for _ in totals) # 调用方跳过 out = [] for t in totals: g = (t + wg - 1) // wg if self.use_ceil else t // wg out.append(max(self.min_group_per_dim, g)) return tuple(out) def describe(self) -> str: return "dispatch 每维 ceil 且 >=1，退化为 0 时跳过派发" POLICY = OrtWebGpuDispatchSizePolicy() def plan_dispatch(totals: tuple, wg: int = 64, policy: OrtWebGpuDispatchSizePolicy = POLICY) -> tuple: return policy.compute(totals, wg)

所有 WebGPU 算子的dispatchWorkgroups都走plan_dispatch，0 维问题被统一根除，Metal 不再报错。

七、解决方案（第三层：断言 / CI 守护）

把“每维>=1、ceil、退化跳过”做成断言。下面用 pytest 守护：

import pytest def test_no_zero_dim(policy): g = policy.compute((200, 128), wg=64) assert all(d >= 1 for d in g if d != 0) # 非跳过时每维>=1 assert policy.forbid_zero_dim_dispatch is True def test_ceil_used(policy): assert policy.use_ceil is True assert policy.compute((200, 128), wg=64) == (4, 2, 1) def test_skip_on_zero(policy): assert policy.skip_on_zero_total is True assert policy.compute((0, 128), wg=64) == (0, 0, 1) # 全 0 -> 跳过 def test_min_group_one(policy): # 极小 total 也应至少 1 组（非退化时） assert policy.compute((1, 1), wg=64) == (1, 1, 1) def test_no_truncation_div(policy): assert "groups = total // wg_size" in policy.forbidden_patterns

这五组断言锁住：(1) 无 0 维；(2) 用 ceil；(3) 退化跳过；(4) 每维至少 1；(5) 禁止截断除法。CI 跑通即代表 WebGPU 派发不会再触发 Metal 的(0,1,1)错误。

八、排查清单

遇到invalid dispatch group size (0,1,1)on macOS Metal：

1. 看是不是 Metal 专属：其它 WebGPU 后端不报 → 是 Metal 严格校验 0 维。
2. 查 dispatch 计算：是不是用整数截断除法，total=0 时得 0。
3. 查退化维度：量化块大小/权重维度组合是否让某维项数为 0。
4. 改 ceil + max(1)：每维ceil且>=1；全 0 退化跳过派发。
5. 统一到OrtWebGpuDispatchSizePolicy：CI 断言禁止 0 维 dispatch。
6. 端到端：macOS Metal 上跑多种块大小/维度组合，确认不报 (0,1,1)。
7. 跨后端：确认 Windows/Linux 下行为一致（不要靠其它后端宽松掩盖）。

九、小结

GatherBlockQuantized: invalid dispatch group size (0,1,1) on macOS Metal WebGPU的根因是：GatherBlockQuantized在派发计算着色器时，用整数截断除法算线程组数，当某个维度的总项数为 0（量化块大小与权重维度的退化组合）时得到 0，而Metal 的dispatchThreadgroups要求每个维度至少为 1，于是硬校验失败报invalid dispatch group size (0,1,1)；其它后端对 0 维较宽松，于是只在 macOS 上崩。

最小修复是对每个 dispatch 维度用ceil并max(1, ...)，且全 0 退化时跳过派发；结构性改进是用唯一的OrtWebGpuDispatchSizePolicy固化派发尺寸计算；CI 用五组断言守护“每维>=1、ceil、退化跳过”。记住：GPU dispatch 的组数永远不能为 0，任何“总项数/块大小”的整除都必须ceil并兜底到 1，否则严格的后端（如 Metal）会直接拒。