已关闭
[Bug]: mssanitizer检测triton tl.dot()时出现OOB误报 #190
Naputao创建于 8月19日关闭于 9月14日
went_code
8月19日 评论：
8月19日 评论：
👋 您好，欢迎向 MindStudio-Sanitizer 提交 Issue！
我们已收到您的反馈，感谢你对开源社区的支持。🎉
📅 处理时效 维护团队将在工作日 24 小时内查看并回复您的问题。
🔍 自助排查（推荐优先查看） 在等待回复期间，您可以先查阅以下资源，多数问题可快速解决：
- MindStudio-Sanitizer官方文档
- 历史 Issue 中相似问题的解决方案
💡 **为了更快定位问题，请您确保 Issue 包含：
- 清晰的问题描述
- 可复现的操作步骤
- 相关日志、截图或环境信息
我们会尽快跟进，感谢您的理解与配合！
8月19日 添加了label：bug
Naputao
8月24日 评论：
8月24日 评论：
新的最小化测试用例
"""Even smaller reproducer: single tl.dot, N split, minimal dims.
Strips the triangular matrix (uses ones), shrinks M/K to one fractal (16),
keeps only the essential: an [M,K] @ [K,N] matmul whose N is split across
programs (grid.x>1) writing a compact fp32 GM output.
"""
import os
import torch, torch_npu # noqa: F401
import triton
import triton.language as tl
@triton.jit
def k(a, b, o, N: tl.constexpr, M: tl.constexpr, K: tl.constexpr, BN: tl.constexpr):
i = tl.program_id(0)
om = tl.arange(0, M)
ok = tl.arange(0, K)
on = i * BN + tl.arange(0, BN)
a_t = tl.load(a + om[:, None] * K + ok[None, :]).to(tl.float32)
b_t = tl.load(b + ok[:, None] * N + on[None, :], mask=on[None, :] < N, other=0.0).to(tl.float32)
o_t = tl.dot(a_t, b_t)
tl.store(o + om[:, None] * N + on[None, :], o_t, mask=on[None, :] < N)
def main():
M = int(os.environ.get("M", "16"))
K = int(os.environ.get("K", "16"))
N = int(os.environ.get("N", "32"))
BN = int(os.environ.get("BN", "16"))
dev = "npu"
a = torch.randn(M, K, device=dev, dtype=torch.bfloat16)
b = torch.randn(K, N, device=dev, dtype=torch.bfloat16)
o = torch.empty(M, N, device=dev, dtype=torch.float)
grid = (triton.cdiv(N, BN),)
print(f"REPRO M={M} K={K} N={N} BN={BN} grid.x={grid[0]} split={'YES' if grid[0]>1 else 'NO'}")
k[grid](a, b, o, N=N, M=M, K=K, BN=BN, num_warps=4, num_stages=2)
torch.npu.synchronize()
print("REPRO done")
if __name__ == "__main__":
main()
运行结果
debian-95-31(@:):workspace# TRITON_DISABLE_LINE_INFO=false mssanitizer python3 A/xpu-ai-stack-hw-npu/mssanitizer/triton/FLA/repro_min2
.py
=================================================================
>>>>> MindStudio <<<<<
THE END-TO-END TOOLCHAIN TO UNLEASH HUAWEI ASCEND COMPUTE
=================================================================
[mssanitizer] logging to file: ./mindstudio_sanitizer_log/mssanitizer_20260824190016_2198342.log
[W824 19:00:17.371394834 FunctionLoader.cpp:48] Warning: LD_PRELOAD detected, FunctionLoader prefers RTLD_DEFAULT for symbol resolution. (function operator())
[mssanitizer] Start memcheck sanitizer on kernel StatelessNormal_9c8573a88cba4fd81bb9851413a4ebba_high_performance_100
[mssanitizer] Sanitizer finished on kernel StatelessNormal_9c8573a88cba4fd81bb9851413a4ebba_high_performance_100. No error detected.
[mssanitizer] Start memcheck sanitizer on kernel StatelessNormal_9c8573a88cba4fd81bb9851413a4ebba_high_performance_100
REPRO M=16 K=16 N=32 BN=16 grid.x=2 split=YES
[mssanitizer] Sanitizer finished on kernel StatelessNormal_9c8573a88cba4fd81bb9851413a4ebba_high_performance_100. No error detected.
[mssanitizer] Start memcheck sanitizer on kernel k
REPRO done
====== WARNING: out of bounds of size 4
====== at 0x12004d798000 on GM when writing data in k
====== in block aic(1) on device 0
====== code in pc current 0x1ac (serialNo:186)
====== #0 /usr1/npuir-a5/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:50:16
====== #1 /usr1/npuir-a5/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:158:52
====== #2 /usr1/npuir-a5/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:263:3
====== #3 /usr1/npuir-a5/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:730:1
====== #4 /workspace/A/xpu-ai-stack-hw-npu/mssanitizer/triton/FLA/repro_min2.py:21:22
====== #5 /workspace/A/xpu-ai-stack-hw-npu/mssanitizer/triton/FLA/repro_min2.py:21:22
[mssanitizer] Sanitizer finished on kernel k. See all detected errors above.
8月26日 将 liusa002 设为负责人
liusa002
9月3日 评论：
9月3日 评论：
@Naputao
【结论】非误报，是NPU-IR的实现存在竞争问题，已与NPU-IR同事确认是问题，通过issue 509跟踪。
【分析】out of bounds是指多核踩踏，即多个核访问同一块内存，可能存在数据踩踏，告警等级为WARNING（含义说明见文档），需要结合竞争检测进一步分析。如果多核之间有正确同步，则竞争检测不会报告警，WARNING可忽略；如果缺少同步，则竞争检测会报ERROR，说明确实存在多核踩踏。
使用竞争检测拉起脚本，确实会报ERROR，说明WARNING的结论准确。
竞争检测拉起方式：
export TRITON_ENABLE_SANITIZER=1
export TRITON_ALWAYS_COMPILE=1
export PYTORCH_NO_NPU_MEMORY_CACHING=1
mssanitizer -t racecheck python3 repro_min2.py
out of bounds拉起方式和打屏：
mssanitizer python3 repro_min2.py
====== WARNING: out of bounds of size 4
====== at 0x1201b555e000 on GM when writing data in k
====== in block aic(1) on device 0
====== code in pc current 0x1f0 (serialNo:277)
====== #0 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:42:14
====== #1 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:157:48
====== #2 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:326:3
====== #3 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:846:1
====== #4 /home/triton_test/internal:0:0
====== #5 /home/triton_test/repro_min2.py:21:22
竞争检测打屏（仅截取部分用于展示）：
====== ERROR: Potential RAW hazard detected at GM in k:
====== PIPE_S Write at RAW()+0x1201b559e000 in block 0 (aic) on device 0 at pc current 0x1f0 (serialNo:293)
====== #0 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:42:14
====== #1 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:157:48
====== #2 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:326:3
====== #3 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:846:1
====== #4 /home/triton_test/internal:0:0
====== #5 /home/triton_test/repro_min2.py:21:22
====== PIPE_S Read at RAW()+0x1201b559e000 in block 1 (aic) on device 0 at pc current 0x1ac (serialNo:437)
====== #0 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:42:20
====== #1 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:157:48
====== #2 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:326:3
====== #3 /usr1/npuir-a3/bishengir/lib/Template/lib/RegBase/Cube/LocalMmad.cpp:846:1
====== #4 /home/triton_test/internal:0:0
====== #5 /home/triton_test/repro_min2.py:21:22
out of bounds和竞争检测都报在NPU-IR的源码LocalMmad.cpp:42中（调用栈与你的不同，是因为NPU-IR版本不同，bishengir-compile --version可查看版本）：
__aicore__ inline uint32_t getL1MmadPingPongId() {
static uint32_t pingPongId = 1;
pingPongId = 1 - pingPongId;
return pingPongId;
}
pingPongId = 1 - pingPongId; <---报在这一行。
static 变量在 .data/.rodata 段上，runtime 会映射到GM。并且runtime 装载二进制的时候不同核的 .text/.data/.rodata 是同一份。pingPongId被声明为static，运行时地址在GM上，不同核访问同一块地址，且硬件不会自动保证跨核的Scalar流水同步。因此认为存在多核踩踏和读写/写写竞争。
9月3日 添加了label：resolved
ascend-robot
9月10日 评论：
9月10日 评论：
您好，当前Issue标记为resolved且有一段时间未进一步更新，因此我们将其标记为'stale'（闲置）状态。若您认为这是误操作，可通过添加任意评论来去除'stale'标签。标记为stale的Issue在4天内无更新活动将自动关闭。
9月10日 添加了label：stale
9月14日 关闭了 issue
9月14日 issue状态由 TODO 改变为 DONE
在提交新问题之前，请确保您已经在社区中搜索过相关问题，并使用了社区中提供的资源/工具后，仍未找到满意的解决方式。
⚠️ 安全信息提醒：请仔细检查提供的文本内容，确保其不包含敏感数据信息，包括但不限于：
在分享配置信息或代码示例时，请将敏感信息脱敏处理，或使用
<TOKEN>等占位符替代原有内容。环境信息
🐛 问题描述
这个 kernel 的核心作用是：对输入 s 在每个 chunk 内沿时间维做前缀累加。
triton源码
运行命令
运行结果
npuir
欢迎加入社区，感谢您对社区的贡献 🎉!