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

【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 width0具体表现只在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 直接拒绝dispatch 维度可能为 0groups total / workgroup_size用整数除法当total 0时结果为 0Metal 要求 1于是校验失败。正确的做法是用ceil并max(1, ...)且对total 0这种退化情况特殊处理跳过派发或至少给 1。块大小/维度组合未校验量化块大小与权重维度组合可能产出一个“0 项”的退化维度算子没在派发前断言/修正。跨后端不一致Metal 严格校验 0其它后端宽松导致“只在 macOS 崩”的迷惑现象容易让人以为是 Metal 驱动问题而非自己算错。错误定位难报错只说(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: 错误整数除法total0 时得 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极端退化(total0)直接不派发。 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_x0, total_y128)) # (0, 2, 1) - Metal 拒 print(fixed :, dispatch_fixed(total_x0, total_y128)) # (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_x0时给出gx0正是 Metal 报的(0,1,1)形态fixed用ceilmax(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::maxuint32_t(1, (total_x wg_size - 1) / wg_size); const uint32_t gy std::maxuint32_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(frozenTrue) class OrtWebGpuDispatchSizePolicy: WebGPU dispatch 尺寸计算的单一事实来源。 # 每个 dispatch 维度必须 1Metal 硬要求 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_dispatch0 维问题被统一根除Metal 不再报错。七、解决方案第三层断言 / CI 守护把“每维1、ceil、退化跳过”做成断言。下面用 pytest 守护import pytest def test_no_zero_dim(policy): g policy.compute((200, 128), wg64) 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), wg64) (4, 2, 1) def test_skip_on_zero(policy): assert policy.skip_on_zero_total is True assert policy.compute((0, 128), wg64) (0, 0, 1) # 全 0 - 跳过 def test_min_group_one(policy): # 极小 total 也应至少 1 组非退化时 assert policy.compute((1, 1), wg64) (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看是不是 Metal 专属其它 WebGPU 后端不报 → 是 Metal 严格校验 0 维。查 dispatch 计算是不是用整数截断除法total0 时得 0。查退化维度量化块大小/权重维度组合是否让某维项数为 0。改 ceil max(1)每维ceil且1全 0 退化跳过派发。统一到OrtWebGpuDispatchSizePolicyCI 断言禁止 0 维 dispatch。端到端macOS Metal 上跑多种块大小/维度组合确认不报 (0,1,1)。跨后端确认 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会直接拒。