常见编译/运行错误速查【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills触发条件当你在以下场景中遇到问题时查阅本文档kernel 编译报MLIRCompilationError编译报ub overflow或local memory limit exceeded编译报coreDim UINT16_MAX运行时报illegal memory access或misaligned address运行时报deadlock detected或barrier timeoutNPU 结果与参考结果差异过大kernel 执行超时或性能远低于预期需要快速定位错误类型和解决方案核心知识错误分类速查表错误类别典型错误信息发生阶段严重程度UB 溢出ub overflow, requires xxxx bits while 1572864 bits available!编译/运行阻断coreDim 超限coreDimxxxx cant be greater than UINT16_MAX编译阻断数据类型不支持unsupported type: uint8/fp64编译阻断操作不支持failed to legalize operation tt.dot编译阻断对齐错误misaligned memory access编译/运行阻断内存访问越界illegal memory access/segfault运行阻断死锁deadlock detected运行阻断同步超时barrier timeout/event timeout运行阻断精度异常结果与参考值差异过大运行功能Scalar 退化Vector 流水利用率 10%运行性能离散访存MTE2 搬运时间过长运行性能编译错误定位方法第一步确定错误发生阶段export TRITON_DEBUG1 export TRITON_DISABLE_CACHE1 python your_program.py检查~/.triton/dump/目录文件状态含义下一步无kernel.ttir.mlirPython Kernel - TTIR 阶段失败检查 kernel 语法有ttir.mlir但无ttadapter.mlirTTIR - TTAdapter 转换失败检查操作是否受 NPU 支持有ttadapter.mlirTTAdapter - 二进制阶段失败使用MLIR_ENABLE_DUMP1定位失败 Pass第二步查看 Pass 详情export MLIR_ENABLE_DUMP1 export TRITON_DEBUG1 export TRITON_DISABLE_CACHE1 python your_program.py 2mlir_dump.log第三步手动编译 TTAdapter IRbishengir-compile kernel.ttadapter.mlir \ --targetAscend910B3 \ --enable-auto-multi-bufferTrue \ --enable-hfusion-compiletrue \ --enable-hivm-compiletrue \ --enable-triton-kernel-compiletrue \ --hivm-compile-argsbishengir-print-ir-afterhivm-inject-sync运行时错误定位方法方法环境变量/命令适用场景解释器模式TRITON_INTERPRET1精度问题、逻辑错误设备打印TRITON_DEVICE_PRINT1tl.device_printNPU 上的中间结果UB 占用分析ENABLE_PRINT_UB_BITS1UB 溢出排查内存分析TRITON_MEMORY_DISPLAY1内存使用详情msProf 上板msprof op --kernel-namexxx python3 xxx.py性能瓶颈msProf 仿真msprof op simulator --soc-versionAscend910B3 ...指令级分析环境变量速查场景推荐环境变量组合编译错误排查TRITON_DEBUG1MLIR_ENABLE_DUMP1TRITON_DISABLE_CACHE1精度调试TRITON_INTERPRET1TRITON_DISABLE_CACHE1运行时打印TRITON_DEVICE_PRINT1TRITON_DISABLE_CACHE1UB 溢出排查ENABLE_PRINT_UB_BITS1TRITON_DEBUG1TRITON_DISABLE_CACHE1coreDim 超限TRITON_ALL_BLOCKS_PARALLEL1深度编译调试TRITON_DEBUG1MLIR_ENABLE_DUMP1TRITON_ENABLE_LLVM_DEBUG1LLVM_DEBUG_ONLYisel性能分析msprof op --kernel-namexxx python3 xxx.py编译加速TRITON_ASCEND_COMPILE_SPEED_OPT1代码模式错误 1UB 溢出错误信息ub overflow, requires 3072256 bits while 1572864 bits available! (possible reason: large or block number is more than what user expect due to multi-buffer feature is enabled and some ops need extra local buffer.)原因单个 AI Core 一次处理的数据量超过 UB 容量。UB 容量参考硬件型号单 bufferdouble buffer说明A2/A3 系列192 KB96 KB x 2最常见910_95 系列256 KB128 KB x 2较新型号解决方案# 方案 1减小 BLOCK_SIZE BLOCK_SIZE 1024 # 从 32768 减小 # 方案 2增加核内 TilingBLOCK_SIZE_SUB triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr, BLOCK_SIZE_SUB: tl.constexpr): pid tl.program_id(axis0) base_offset pid * BLOCK_SIZE num_sub_blocks tl.cdiv(BLOCK_SIZE, BLOCK_SIZE_SUB) for sub_idx in range(num_sub_blocks): offsets base_offset sub_idx * BLOCK_SIZE_SUB tl.arange(0, BLOCK_SIZE_SUB) mask offsets N x tl.load(x_ptr offsets, maskmask) tl.store(out_ptr offsets, x * 2.0, maskmask) # 方案 3关闭 multibuffer triton.Config({XS: 128, multibuffer: False}) # 方案 4查看 UB 占用量 # export ENABLE_PRINT_UB_BITS1注意int8 类型会占用更大的片上空间约为其他类型的 2 倍考虑使用 int32 替代。错误 2coreDim 超限错误信息coreDim524288 cant be greater than UINT16_MAX原因Grid 维度超过 NPU 硬件限制65535。coreDim ceil(N / BLOCK_SIZE)。解决方案# 方案 1环境变量kernel 逻辑必须对执行顺序不敏感 # export TRITON_ALL_BLOCKS_PARALLEL1 # 方案 2动态计算 BLOCK_SIZE N x.numel() min_block_size triton.next_power_of_2(triton.cdiv(N, 65535)) BLOCK_SIZE max(32768, min_block_size)警告TRITON_ALL_BLOCKS_PARALLEL1要求 kernel 逻辑对执行顺序不敏感有原子操作或跨核同步时不能开启否则会死锁。错误 3数据类型不支持错误信息error: unsupported type: uint8 / uint64 / fp64原因NPU 不支持 uint8/uint16/uint32/uint64/fp64 等数据类型。解决方案# uint8 - int32 x_int32 x.to(tl.int32) # uint64 - int64 x_int64 x.to(tl.int64) # fp64 - fp32 x_fp32 x.to(tl.float32)NPU 支持的数据类型int8/16/32/64, fp16/32, bf16, bool910_95 额外支持 fp8NPU 不支持的数据类型uint8/16/32/64, fp64错误 4对齐错误错误信息error: misaligned memory access原因Vector 算子要求 32 字节访存对齐Cube-Vector 融合算子要求 512 字节对齐。解决方案# 方案 1确保 BLOCK_SIZE * element_size 满足对齐要求 # fp32: BLOCK_SIZE 应为 8 的倍数8 * 4B 32B # fp16/bf16: BLOCK_SIZE 应为 16 的倍数16 * 2B 32B # Cube-Vector 融合: BLOCK_SIZE 应为 512B 对齐 # 方案 2使用 tl.multiple_of 提示编译器对齐信息 offsets pid * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) offsets tl.multiple_of(offsets, 128) # 提示 offsets 是 128 的倍数错误 5dot 操作降级失败错误信息error: failed to legalize operation tt.dot原因dot 操作的输入数据类型不被支持或 BLOCK_SIZE 不满足对齐要求。解决方案import triton.language.extra.cann.extension as al # 方案 1确保输入为 fp16/fp32/bf16/int8 a tl.load(a_ptr offsets).to(tl.float16) b tl.load(b_ptr offsets).to(tl.float16) # 方案 2添加 compile_hint a al.compile_hint(a, dot_pad_only_k) b al.compile_hint(b, dot_pad_only_k) result tl.dot(a, b) # 方案 3检查 BLOCK_M 和 BLOCK_N 是否为合法值 # 推荐 512B 对齐的 BLOCK_SIZEBLOCK_M128, BLOCK_N256, BLOCK_K256错误 6Block Pointer 与复杂控制流搭配编译失败错误信息error: failed to legalize operation tt.advance原因tl.advance与复杂循环/分支语句搭配时可能出现编译问题。解决方案# 错误写法tl.advance 与复杂循环搭配 block_ptr tl.make_block_ptr(...) for i in range(N): x tl.load(block_ptr) block_ptr tl.advance(block_ptr, [BLOCK_SIZE]) # 可能编译失败 # 正确写法重新创建 block_ptr 替代 advance for i in range(N): offsets i * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) block_ptr tl.make_block_ptr(basebase_ptr, shapeshape, stridesstrides, offsets[i * BLOCK_SIZE], block_shape[BLOCK_SIZE], order[0]) x tl.load(block_ptr)错误 7stride 顺序不支持错误信息error: unsupported stride pattern for transpose原因NPU 不支持通过调整 stride 参数顺序实现转置语义。解决方案# 错误写法通过 stride 实现转置 block_ptr tl.make_block_ptr(baseptr, shape[M, N], strides[1, M], # 转置 stride offsets[0, 0], block_shape[BLOCK_M, BLOCK_N], order[0, 1]) # 正确写法使用 order 参数表达转置stride 反映真实内存布局 block_ptr tl.make_block_ptr(baseptr, shape[N, M], strides[M, 1], # 真实内存布局 offsets[0, 0], block_shape[BLOCK_N, BLOCK_M], order[1, 0]) # order 表达转置错误 8精度异常典型表现NPU 运行结果与 PyTorch/CPU/GPU 参考结果差异过大。排查步骤# 第一步解释器模式验证逻辑 # export TRITON_INTERPRET1 # export TRITON_DISABLE_CACHE1 # 第二步在 NPU 上打印中间结果 # export TRITON_DEVICE_PRINT1 triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr): x tl.load(x_ptr offsets, maskmask) tl.device_print(x after load: , x) result tl.sqrt(x) tl.device_print(result after sqrt: , result) tl.store(out_ptr offsets, result, maskmask)常见精度问题原因原因表现解决方案浮点计算顺序差异微小差异1e-6 量级正常现象使用torch.allclose数据类型退化较大差异检查 int64/int32 是否导致 Scalar 退化bf16/fp16 精度损失中等差异关键计算在 fp32 下进行mask 处理差异特定位置差异检查 boundary_check 和 mask 逻辑tl.exp 近似Softmax 差异NPU 上 exp 是近似计算差异 1e-7 量级属正常fp16 精度不够的通用解决方案triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr): x tl.load(x_ptr offsets, maskmask) # kernel 内部使用 fp32 精度计算 x_fp32 x.to(tl.float32) result tl.sqrt(x_fp32) # 写回时转换 tl.store(out_ptr offsets, result.to(tl.float16), maskmask)错误 9Scalar 退化典型表现Vector 流水利用率极低10%Scalar 流水成为瓶颈。常见 Scalar 退化原因操作不支持的数据类型退化行为Vector ADDint64退化为 Scalar 加法Vector CMPint64, int32退化为 Scalar 比较解决方案# 优化前cols N 导致 Scalar CMPcols 为 int64 cols tl.arange(0, BLOCK_N) xbar tl.where(cols N, x - mean, 0.0) # 优化后转换为 fp32 使用 Vector CMP cols_cmp cols.to(tl.float32) xbar tl.where(cols_cmp N, x - mean, 0.0)错误 10device_print 缺少 prefix错误信息error: device_print requires a prefix string解决方案# 错误写法 tl.device_print(x) # 正确写法 tl.device_print(x: , x)错误 11Tensor 形状为空错误信息error: all shapes must have size 1原因NPU 不支持 shape 中某个维度 size 小于 1 的 tensor。解决方案# 在调用 kernel 前检查数据规模 if x.numel() 0: return # 跳过空 tensor kernel[grid](https://link.gitcode.com/i/ab3cbea374e4dbb4b8e38ee55cfc7da6)错误 12bool 类型不支持错误信息error: unsupported type: bool原因部分操作不支持 bool 类型。解决方案Triton 内部会将 bool 转为 int8 进行运算通常无需手动处理。如果特定操作报错显式转换mask_int8 mask.to(tl.int8)错误 13atomic_add 多核累加不支持错误信息error: atomic_add not supported for multi-core accumulation解决方案# 方案 1使用 atomic_cas 自旋锁保护共享缓冲区 while True: old tl.atomic_cas(lock_ptr, 0, 1) if old 0: break # 临界区普通 add tl.store(out_ptr offsets, current_val new_val) # 释放锁 tl.atomic_xchg(lock_ptr, 0) # 方案 2两阶段归约策略 # 阶段 1每个 Block 写入自己的部分结果 # 阶段 2单独的 kernel 做最终归约错误 14MLIR Pass 失败错误信息error: Failed to run BishengHIR pipeline解决方案# 使用 MLIR_ENABLE_DUMP1 查看具体哪个 Pass 失败 export MLIR_ENABLE_DUMP1 export TRITON_DEBUG1 export TRITON_DISABLE_CACHE1 python your_program.py 2mlir_dump.log # 手动编译 TTAdapter IR 进行调试 bishengir-compile kernel.ttadapter.mlir \ --targetAscend910B3 \ --enable-auto-multi-bufferTrue \ --enable-hfusion-compiletrue \ --enable-hivm-compiletrue \ --enable-triton-kernel-compiletrue错误 15运行时死锁错误信息deadlock detected常见原因与解决方案原因解决方案atomic_cas自旋锁未释放确保锁的获取和释放在同一执行路径上避免条件分支导致锁未释放TRITON_ALL_BLOCKS_PARALLEL1与同步操作冲突有同步操作时不能开启此选项sync_block_set/wait不配对确保每个 set 有对应 wait且参数一致条件分支中获取锁避免在 if/else 分支中获取锁确保所有路径都能释放锁错误 16NPU 执行超时可能原因Grid 分核数过多分批调度开销大死锁kernel 逻辑对执行顺序敏感数据规模异常大解决方案# 方案 1Grid 对齐物理核数 num_aicore triton.runtime.get_npu_properties()[num_aicore] grid (min(num_aicore, triton.cdiv(N, BLOCK_SIZE)),) # 方案 2检查 kernel 是否存在跨核依赖 # 方案 3减小数据规模进行测试错误 17结果全零或全 NaN排查方法# 使用 device_print 逐步排查 triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr): x tl.load(x_ptr offsets, maskmask) tl.device_print(x: , x) # 检查数据是否正确加载 result tl.sqrt(x) tl.device_print(result: , result) # 检查是否有 NaN tl.store(out_ptr offsets, result, maskmask)常见原因数据未正确加载mask 错误数据类型转换错误除零操作或 0 的负数次幂使用tl.where保护危险操作# 安全的 sqrt result tl.where(x 0, tl.sqrt(x), 0.0) # 安全的除法 result tl.where(denom ! 0, numer / denom, 0.0)910_95 特别注意UB 容量差异硬件型号单 bufferdouble bufferA2/A3192 KB96 KB x 2910_95256 KB128 KB x 2910_95 的 UB 容量更大UB 溢出的阈值更高。但 int8 类型在 910_95 上同样会占用更大的片上空间。multibuffer 默认关闭910_95 上 multibuffer 默认关闭A2/A3 默认开启。如果遇到 UB 溢出检查是否手动开启了 multibuffer# 910_95 上默认不需要 multibuffer除非有特殊需求 triton.Config({XS: 128, multibuffer: False}) # 默认行为L0C - UB 直通FixPipe910_95 支持 L0C - UB 直通通路通过 FixPipeA2/A3 不支持。这意味着 Cube 计算结果可以直接搬入 UB减少了中间搬运步骤。但也改变了同步的 Pipe 组合可能影响同步行为。fp8 数据类型910_95 额外支持 fp8 数据类型A2/A3 不支持。使用 fp8 时注意fp8 的动态范围更小更容易溢出确保计算在 fp32 下进行只在存储时转换为 fp8SIMT 模式910_95 支持 SIMT 模式A2/A3 不支持。SIMT 模式下每个 Scalar Core 可以独立执行不同的指令路径减少了分支发散的影响。同步机制差异910_95 采用 Reg-based 同步架构SetFlag/WaitFlag而非 A2/A3 的 FFTS 机制。在 910_95 上不需要设置ffts_base_addr跨核同步更轻量延迟更低同步语义与 FFTS 一致Triton 层 API 无需修改编译目标差异手动编译 TTAdapter IR 时注意--target参数# A2/A3 bishengir-compile kernel.ttadapter.mlir --targetAscend910B3 ... # 910_95 bishengir-compile kernel.ttadapter.mlir --targetAscend910_95 ...相关文档01-debug-overview.md - 调试方法总览02-interpreter-mode.md - 解释器模式调试03-compile-errors.md - 编译错误排查04-runtime-errors.md - 运行时错误排查05-environment-variables.md - 调试相关环境变量03-error-codes.md - 错误码参考04-env-variables.md - 环境变量完整参考05-faq.md - 常见问题速查【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
常见编译/运行错误速查
发布时间：2026/7/11 11:20:32
常见编译/运行错误速查【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills触发条件当你在以下场景中遇到问题时查阅本文档kernel 编译报MLIRCompilationError编译报ub overflow或local memory limit exceeded编译报coreDim UINT16_MAX运行时报illegal memory access或misaligned address运行时报deadlock detected或barrier timeoutNPU 结果与参考结果差异过大kernel 执行超时或性能远低于预期需要快速定位错误类型和解决方案核心知识错误分类速查表错误类别典型错误信息发生阶段严重程度UB 溢出ub overflow, requires xxxx bits while 1572864 bits available!编译/运行阻断coreDim 超限coreDimxxxx cant be greater than UINT16_MAX编译阻断数据类型不支持unsupported type: uint8/fp64编译阻断操作不支持failed to legalize operation tt.dot编译阻断对齐错误misaligned memory access编译/运行阻断内存访问越界illegal memory access/segfault运行阻断死锁deadlock detected运行阻断同步超时barrier timeout/event timeout运行阻断精度异常结果与参考值差异过大运行功能Scalar 退化Vector 流水利用率 10%运行性能离散访存MTE2 搬运时间过长运行性能编译错误定位方法第一步确定错误发生阶段export TRITON_DEBUG1 export TRITON_DISABLE_CACHE1 python your_program.py检查~/.triton/dump/目录文件状态含义下一步无kernel.ttir.mlirPython Kernel - TTIR 阶段失败检查 kernel 语法有ttir.mlir但无ttadapter.mlirTTIR - TTAdapter 转换失败检查操作是否受 NPU 支持有ttadapter.mlirTTAdapter - 二进制阶段失败使用MLIR_ENABLE_DUMP1定位失败 Pass第二步查看 Pass 详情export MLIR_ENABLE_DUMP1 export TRITON_DEBUG1 export TRITON_DISABLE_CACHE1 python your_program.py 2mlir_dump.log第三步手动编译 TTAdapter IRbishengir-compile kernel.ttadapter.mlir \ --targetAscend910B3 \ --enable-auto-multi-bufferTrue \ --enable-hfusion-compiletrue \ --enable-hivm-compiletrue \ --enable-triton-kernel-compiletrue \ --hivm-compile-argsbishengir-print-ir-afterhivm-inject-sync运行时错误定位方法方法环境变量/命令适用场景解释器模式TRITON_INTERPRET1精度问题、逻辑错误设备打印TRITON_DEVICE_PRINT1tl.device_printNPU 上的中间结果UB 占用分析ENABLE_PRINT_UB_BITS1UB 溢出排查内存分析TRITON_MEMORY_DISPLAY1内存使用详情msProf 上板msprof op --kernel-namexxx python3 xxx.py性能瓶颈msProf 仿真msprof op simulator --soc-versionAscend910B3 ...指令级分析环境变量速查场景推荐环境变量组合编译错误排查TRITON_DEBUG1MLIR_ENABLE_DUMP1TRITON_DISABLE_CACHE1精度调试TRITON_INTERPRET1TRITON_DISABLE_CACHE1运行时打印TRITON_DEVICE_PRINT1TRITON_DISABLE_CACHE1UB 溢出排查ENABLE_PRINT_UB_BITS1TRITON_DEBUG1TRITON_DISABLE_CACHE1coreDim 超限TRITON_ALL_BLOCKS_PARALLEL1深度编译调试TRITON_DEBUG1MLIR_ENABLE_DUMP1TRITON_ENABLE_LLVM_DEBUG1LLVM_DEBUG_ONLYisel性能分析msprof op --kernel-namexxx python3 xxx.py编译加速TRITON_ASCEND_COMPILE_SPEED_OPT1代码模式错误 1UB 溢出错误信息ub overflow, requires 3072256 bits while 1572864 bits available! (possible reason: large or block number is more than what user expect due to multi-buffer feature is enabled and some ops need extra local buffer.)原因单个 AI Core 一次处理的数据量超过 UB 容量。UB 容量参考硬件型号单 bufferdouble buffer说明A2/A3 系列192 KB96 KB x 2最常见910_95 系列256 KB128 KB x 2较新型号解决方案# 方案 1减小 BLOCK_SIZE BLOCK_SIZE 1024 # 从 32768 减小 # 方案 2增加核内 TilingBLOCK_SIZE_SUB triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr, BLOCK_SIZE_SUB: tl.constexpr): pid tl.program_id(axis0) base_offset pid * BLOCK_SIZE num_sub_blocks tl.cdiv(BLOCK_SIZE, BLOCK_SIZE_SUB) for sub_idx in range(num_sub_blocks): offsets base_offset sub_idx * BLOCK_SIZE_SUB tl.arange(0, BLOCK_SIZE_SUB) mask offsets N x tl.load(x_ptr offsets, maskmask) tl.store(out_ptr offsets, x * 2.0, maskmask) # 方案 3关闭 multibuffer triton.Config({XS: 128, multibuffer: False}) # 方案 4查看 UB 占用量 # export ENABLE_PRINT_UB_BITS1注意int8 类型会占用更大的片上空间约为其他类型的 2 倍考虑使用 int32 替代。错误 2coreDim 超限错误信息coreDim524288 cant be greater than UINT16_MAX原因Grid 维度超过 NPU 硬件限制65535。coreDim ceil(N / BLOCK_SIZE)。解决方案# 方案 1环境变量kernel 逻辑必须对执行顺序不敏感 # export TRITON_ALL_BLOCKS_PARALLEL1 # 方案 2动态计算 BLOCK_SIZE N x.numel() min_block_size triton.next_power_of_2(triton.cdiv(N, 65535)) BLOCK_SIZE max(32768, min_block_size)警告TRITON_ALL_BLOCKS_PARALLEL1要求 kernel 逻辑对执行顺序不敏感有原子操作或跨核同步时不能开启否则会死锁。错误 3数据类型不支持错误信息error: unsupported type: uint8 / uint64 / fp64原因NPU 不支持 uint8/uint16/uint32/uint64/fp64 等数据类型。解决方案# uint8 - int32 x_int32 x.to(tl.int32) # uint64 - int64 x_int64 x.to(tl.int64) # fp64 - fp32 x_fp32 x.to(tl.float32)NPU 支持的数据类型int8/16/32/64, fp16/32, bf16, bool910_95 额外支持 fp8NPU 不支持的数据类型uint8/16/32/64, fp64错误 4对齐错误错误信息error: misaligned memory access原因Vector 算子要求 32 字节访存对齐Cube-Vector 融合算子要求 512 字节对齐。解决方案# 方案 1确保 BLOCK_SIZE * element_size 满足对齐要求 # fp32: BLOCK_SIZE 应为 8 的倍数8 * 4B 32B # fp16/bf16: BLOCK_SIZE 应为 16 的倍数16 * 2B 32B # Cube-Vector 融合: BLOCK_SIZE 应为 512B 对齐 # 方案 2使用 tl.multiple_of 提示编译器对齐信息 offsets pid * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) offsets tl.multiple_of(offsets, 128) # 提示 offsets 是 128 的倍数错误 5dot 操作降级失败错误信息error: failed to legalize operation tt.dot原因dot 操作的输入数据类型不被支持或 BLOCK_SIZE 不满足对齐要求。解决方案import triton.language.extra.cann.extension as al # 方案 1确保输入为 fp16/fp32/bf16/int8 a tl.load(a_ptr offsets).to(tl.float16) b tl.load(b_ptr offsets).to(tl.float16) # 方案 2添加 compile_hint a al.compile_hint(a, dot_pad_only_k) b al.compile_hint(b, dot_pad_only_k) result tl.dot(a, b) # 方案 3检查 BLOCK_M 和 BLOCK_N 是否为合法值 # 推荐 512B 对齐的 BLOCK_SIZEBLOCK_M128, BLOCK_N256, BLOCK_K256错误 6Block Pointer 与复杂控制流搭配编译失败错误信息error: failed to legalize operation tt.advance原因tl.advance与复杂循环/分支语句搭配时可能出现编译问题。解决方案# 错误写法tl.advance 与复杂循环搭配 block_ptr tl.make_block_ptr(...) for i in range(N): x tl.load(block_ptr) block_ptr tl.advance(block_ptr, [BLOCK_SIZE]) # 可能编译失败 # 正确写法重新创建 block_ptr 替代 advance for i in range(N): offsets i * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) block_ptr tl.make_block_ptr(basebase_ptr, shapeshape, stridesstrides, offsets[i * BLOCK_SIZE], block_shape[BLOCK_SIZE], order[0]) x tl.load(block_ptr)错误 7stride 顺序不支持错误信息error: unsupported stride pattern for transpose原因NPU 不支持通过调整 stride 参数顺序实现转置语义。解决方案# 错误写法通过 stride 实现转置 block_ptr tl.make_block_ptr(baseptr, shape[M, N], strides[1, M], # 转置 stride offsets[0, 0], block_shape[BLOCK_M, BLOCK_N], order[0, 1]) # 正确写法使用 order 参数表达转置stride 反映真实内存布局 block_ptr tl.make_block_ptr(baseptr, shape[N, M], strides[M, 1], # 真实内存布局 offsets[0, 0], block_shape[BLOCK_N, BLOCK_M], order[1, 0]) # order 表达转置错误 8精度异常典型表现NPU 运行结果与 PyTorch/CPU/GPU 参考结果差异过大。排查步骤# 第一步解释器模式验证逻辑 # export TRITON_INTERPRET1 # export TRITON_DISABLE_CACHE1 # 第二步在 NPU 上打印中间结果 # export TRITON_DEVICE_PRINT1 triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr): x tl.load(x_ptr offsets, maskmask) tl.device_print(x after load: , x) result tl.sqrt(x) tl.device_print(result after sqrt: , result) tl.store(out_ptr offsets, result, maskmask)常见精度问题原因原因表现解决方案浮点计算顺序差异微小差异1e-6 量级正常现象使用torch.allclose数据类型退化较大差异检查 int64/int32 是否导致 Scalar 退化bf16/fp16 精度损失中等差异关键计算在 fp32 下进行mask 处理差异特定位置差异检查 boundary_check 和 mask 逻辑tl.exp 近似Softmax 差异NPU 上 exp 是近似计算差异 1e-7 量级属正常fp16 精度不够的通用解决方案triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr): x tl.load(x_ptr offsets, maskmask) # kernel 内部使用 fp32 精度计算 x_fp32 x.to(tl.float32) result tl.sqrt(x_fp32) # 写回时转换 tl.store(out_ptr offsets, result.to(tl.float16), maskmask)错误 9Scalar 退化典型表现Vector 流水利用率极低10%Scalar 流水成为瓶颈。常见 Scalar 退化原因操作不支持的数据类型退化行为Vector ADDint64退化为 Scalar 加法Vector CMPint64, int32退化为 Scalar 比较解决方案# 优化前cols N 导致 Scalar CMPcols 为 int64 cols tl.arange(0, BLOCK_N) xbar tl.where(cols N, x - mean, 0.0) # 优化后转换为 fp32 使用 Vector CMP cols_cmp cols.to(tl.float32) xbar tl.where(cols_cmp N, x - mean, 0.0)错误 10device_print 缺少 prefix错误信息error: device_print requires a prefix string解决方案# 错误写法 tl.device_print(x) # 正确写法 tl.device_print(x: , x)错误 11Tensor 形状为空错误信息error: all shapes must have size 1原因NPU 不支持 shape 中某个维度 size 小于 1 的 tensor。解决方案# 在调用 kernel 前检查数据规模 if x.numel() 0: return # 跳过空 tensor kernel[grid](https://link.gitcode.com/i/ab3cbea374e4dbb4b8e38ee55cfc7da6)错误 12bool 类型不支持错误信息error: unsupported type: bool原因部分操作不支持 bool 类型。解决方案Triton 内部会将 bool 转为 int8 进行运算通常无需手动处理。如果特定操作报错显式转换mask_int8 mask.to(tl.int8)错误 13atomic_add 多核累加不支持错误信息error: atomic_add not supported for multi-core accumulation解决方案# 方案 1使用 atomic_cas 自旋锁保护共享缓冲区 while True: old tl.atomic_cas(lock_ptr, 0, 1) if old 0: break # 临界区普通 add tl.store(out_ptr offsets, current_val new_val) # 释放锁 tl.atomic_xchg(lock_ptr, 0) # 方案 2两阶段归约策略 # 阶段 1每个 Block 写入自己的部分结果 # 阶段 2单独的 kernel 做最终归约错误 14MLIR Pass 失败错误信息error: Failed to run BishengHIR pipeline解决方案# 使用 MLIR_ENABLE_DUMP1 查看具体哪个 Pass 失败 export MLIR_ENABLE_DUMP1 export TRITON_DEBUG1 export TRITON_DISABLE_CACHE1 python your_program.py 2mlir_dump.log # 手动编译 TTAdapter IR 进行调试 bishengir-compile kernel.ttadapter.mlir \ --targetAscend910B3 \ --enable-auto-multi-bufferTrue \ --enable-hfusion-compiletrue \ --enable-hivm-compiletrue \ --enable-triton-kernel-compiletrue错误 15运行时死锁错误信息deadlock detected常见原因与解决方案原因解决方案atomic_cas自旋锁未释放确保锁的获取和释放在同一执行路径上避免条件分支导致锁未释放TRITON_ALL_BLOCKS_PARALLEL1与同步操作冲突有同步操作时不能开启此选项sync_block_set/wait不配对确保每个 set 有对应 wait且参数一致条件分支中获取锁避免在 if/else 分支中获取锁确保所有路径都能释放锁错误 16NPU 执行超时可能原因Grid 分核数过多分批调度开销大死锁kernel 逻辑对执行顺序敏感数据规模异常大解决方案# 方案 1Grid 对齐物理核数 num_aicore triton.runtime.get_npu_properties()[num_aicore] grid (min(num_aicore, triton.cdiv(N, BLOCK_SIZE)),) # 方案 2检查 kernel 是否存在跨核依赖 # 方案 3减小数据规模进行测试错误 17结果全零或全 NaN排查方法# 使用 device_print 逐步排查 triton.jit def kernel(x_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr): x tl.load(x_ptr offsets, maskmask) tl.device_print(x: , x) # 检查数据是否正确加载 result tl.sqrt(x) tl.device_print(result: , result) # 检查是否有 NaN tl.store(out_ptr offsets, result, maskmask)常见原因数据未正确加载mask 错误数据类型转换错误除零操作或 0 的负数次幂使用tl.where保护危险操作# 安全的 sqrt result tl.where(x 0, tl.sqrt(x), 0.0) # 安全的除法 result tl.where(denom ! 0, numer / denom, 0.0)910_95 特别注意UB 容量差异硬件型号单 bufferdouble bufferA2/A3192 KB96 KB x 2910_95256 KB128 KB x 2910_95 的 UB 容量更大UB 溢出的阈值更高。但 int8 类型在 910_95 上同样会占用更大的片上空间。multibuffer 默认关闭910_95 上 multibuffer 默认关闭A2/A3 默认开启。如果遇到 UB 溢出检查是否手动开启了 multibuffer# 910_95 上默认不需要 multibuffer除非有特殊需求 triton.Config({XS: 128, multibuffer: False}) # 默认行为L0C - UB 直通FixPipe910_95 支持 L0C - UB 直通通路通过 FixPipeA2/A3 不支持。这意味着 Cube 计算结果可以直接搬入 UB减少了中间搬运步骤。但也改变了同步的 Pipe 组合可能影响同步行为。fp8 数据类型910_95 额外支持 fp8 数据类型A2/A3 不支持。使用 fp8 时注意fp8 的动态范围更小更容易溢出确保计算在 fp32 下进行只在存储时转换为 fp8SIMT 模式910_95 支持 SIMT 模式A2/A3 不支持。SIMT 模式下每个 Scalar Core 可以独立执行不同的指令路径减少了分支发散的影响。同步机制差异910_95 采用 Reg-based 同步架构SetFlag/WaitFlag而非 A2/A3 的 FFTS 机制。在 910_95 上不需要设置ffts_base_addr跨核同步更轻量延迟更低同步语义与 FFTS 一致Triton 层 API 无需修改编译目标差异手动编译 TTAdapter IR 时注意--target参数# A2/A3 bishengir-compile kernel.ttadapter.mlir --targetAscend910B3 ... # 910_95 bishengir-compile kernel.ttadapter.mlir --targetAscend910_95 ...相关文档01-debug-overview.md - 调试方法总览02-interpreter-mode.md - 解释器模式调试03-compile-errors.md - 编译错误排查04-runtime-errors.md - 运行时错误排查05-environment-variables.md - 调试相关环境变量03-error-codes.md - 错误码参考04-env-variables.md - 环境变量完整参考05-faq.md - 常见问题速查【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
