wenyi-li 调试与排查清单
完整调试清单
内存访问问题
- 所有 load/store 是否都有 mask 或 boundary_check？
- stride 参数设置是否正确？
- 数组索引是否越界？
控制流检查
- 是否误用了 return/break/continue？
- 复杂条件是否用 mask 组合实现？
-
tl.constexpr是否只在内核参数中使用？
Grid 与 Block 配置检查
- Grid 总大小是否不超过 65535？
- 对于大 shape 算子，是否采用了交错循环
for i in range(pid, total, core_num)？ - Grid 维度是否为 tuple 类型且不超过 3 维？
并发与原子操作检查
- 并发写入是否使用了原子操作（
tl.atomic_add等）？ - 原子操作是否必要（能否避免）？
切片与索引检查
- 是否避免了Python风格的直接切片（如
b[0]、b[i:j]）？ - 是否对
tl.arange生成的张量误用了tl.get_element？ - 切片操作是否使用了正确的API（
tl.get_element、tl.extract_slice等）？
性能优化检查
- 内存访问是否连续（避免跨步访问）？
- 是否充分利用了块内并行？
- 复杂算子是否考虑拆分为多个简单kernel？
禁止使用的语法（Ascend 后端）
| 禁止写法 | 替代方案 |
|---|---|
return / break / continue |
使用 mask 控制流程 |
| lambda 表达式 | 内联函数或 tl.where |
链式布尔运算 a and b |
分步计算 mask：m1 = ...; m2 = ...; m = m1 & m2 |
张量直接索引 tensor[i] |
tl.load(ptr + offset) / tl.store(ptr + offset, val) |
Python 切片 b[0] / b[i:j] |
tl.get_element / tl.extract_slice / tl.insert_slice |
对 tl.arange 结果用 get_element |
直接计算索引值 |
while 循环 |
for i in range(MAX): if i < n: |
range() 混用运行时变量和 constexpr |
全 constexpr 的 range(0, N, BLOCK_K) + 循环体内运行时 if |
tl.float16(scalar) |
scalar.to(tl.float16) |
tl.constexpr 在 host 侧使用 |
仅在 kernel 参数中使用 |
| if-else 中负偏移 | tl.maximum(offset, 0) |
复杂 tl.where 用于内存偏移 |
拆分为 if-else 静态分支 |
常见错误速查表
编译错误
| 错误类型 | 典型症状 | 常见原因 | 解决方案 |
|---|---|---|---|
| UB/CBUF 溢出 | ub overflow, requires X bits while 1572864 bits available |
BLOCK 尺寸过大或中间变量过多 | 缩小 BLOCK 尺寸；减少同时活跃的 tensor 数 |
| HiVM vsel 错误 | hivm.hir.vsel: Unsupported op for finding the root alloc |
嵌套 mask + tl.where 组合过于复杂 | 用乘法替代 tl.where：a * mask.to(dtype) |
| 内存越界访问 | 运行时错误、结果异常、随机崩溃 | load/store缺少mask或boundary_check | 添加正确的mask或boundary_check保护 |
| Grid超限 | 编译失败或运行时错误 | grid总大小超过65535 | 使用交错循环for i in range(pid, total, core_num)或连续分块处理 |
| 控制流错误 | unsupported AST node type: Continue |
使用了return/break/continue | 改用 if-else 包裹逻辑 |
| while循环错误 | 编译失败（Ascend后端） | 使用了while循环 | 改用for + if替代：for i in range(MAX): if i < n: |
| constexpr 索引 | ValueError('unsupported tensor index: constexpr[0]') |
对 tl.sum 等返回的标量做 [0] 索引 |
直接使用标量结果，不要索引 |
| 切片语法错误 | 编译失败 | 使用了b[0]或b[i:j]直接切片 |
使用tl.get_element或tl.extract_slice |
| tl.arange索引错误 | 编译失败 | 对tl.arange结果使用get_element |
直接计算索引值而非提取 |
| 类型转换错误 | cast incompatible |
隐式 cast 或使用tl.float16(scalar) |
用 .to(tl.float16) 或显式指定 acc dtype |
| constexpr误用 | 编译失败 | 在host侧使用tl.constexpr | 仅在kernel参数中使用tl.constexpr |
| Stride设置错误 | 计算结果错误、数据错位 | stride参数计算或传递错误 | 验证stride设置，检查tensor.stride() |
| 数值不稳定 | 结果为NaN或Inf | softmax/sqrt等操作溢出 | 减去最大值、检查非负、使用float32 |
| 数据竞争 | 结果不确定、每次运行不同 | 多program并发写入同一位置 | 使用tl.atomic_add等原子操作 |
| BLOCK_SIZE过大 | 编译失败或运行时错误 | BLOCK_SIZE超过65536或硬件限制 | 减小BLOCK_SIZE，使用循环处理 |
| tl.where偏移计算 | 编译失败（Ascend后端） | 在内存偏移中使用tl.where | 改用if-else静态分支处理 |
| 性能低下 | 运行缓慢 | 内存访问不连续、切分不合理 | 优化内存布局、调整BLOCK_SIZE、使用block_ptr |
| 运行时range边界崩溃 | bishengIR crash | range()的start/stop混用运行时变量和constexpr | 改用全constexpr的range(0, N, BLOCK_K)，循环体内用运行时if跳过 |