Bug 描述
在 Ascend A3 上运行 examples/normalization/rms_norm.py 时，
_rms_norm_grad(1024, 1024, 128, 128, dtype="float") 能完成编译，但生成的
kernel 存在 UB 越界访问，执行时触发 AICore 异常。
rms_norm.py 在进入 grad 测试前已经输出 Test passed! 和
Kernel Output Match!。examples/bench_test.sh 将输出中任意位置出现的这些文字视为
成功，且不要求子进程退出码为 0，因此后续 grad 触发 AICore 异常后仍被报告为
[PASSED]。
这会导致全量 examples 显示 142/142 通过，但其中实际包含未完成全部测试的用例。
环境
- NPU：Ascend A3
- CANN：9.1.0-beta.1
- torch-npu：2.10.0
- Python：3.12.13
- TileLang-Ascend：
d2bd9ec7c97a08db31454485c978ae177a06c56f
复现代码
直接运行 example，可以观察到 grad 阶段的运行时异常：
source set_env.sh
cd examples
python normalization/rms_norm.py
通过旧 bench 运行相同目录，则会被误报为通过：
source set_env.sh
cd examples
bash bench_test.sh --dirs normalization --skip-pytest
失败配置为：
func_grad = _rms_norm_grad(
1024,
1024,
128,
128,
dtype="float",
)
dx = func_grad(dy, x)错误信息
NPU 执行时报告：
VEC instruction error: the ub address out of bounds
AICore error 507015
通过 kernel.get_kernel_source() 检查生成的 Ascend C，可以看到：
auto math_mat = ascend_ub.GetWithOffset<float>(8192, 164160);
其结束地址为：
164160 + 8192 × sizeof(float) = 196928 bytes
A3 UB 上限为 196352 字节，因此越界 576 字节。编译阶段没有报告容量错误。
旧 bench_test.sh 的成功判定近似为：
if output_matches_success_marker || custom_task_exits_zero; then
report_passed
fi普通 example 的退出码没有参与成功判断。由于 forward 完成后已经输出
Kernel Output Match!，后续 grad 阶段的 AICore 异常不会改变 bench 的判定结果。
期望行为
- 普通 example 只有在子进程退出码为 0，且所有测试阶段完成后输出最终成功标记时，才能被
判为通过。
_rms_norm_grad 不应生成有效访问范围超过 UB 容量的 kernel。
希望对 rms_norm.py 和 bench_test.sh 进行一次原子修复，使全量 examples 的成功结果能够
真实表示所有测试阶段均已完成。
实际行为
_rms_norm_grad 在执行时因 AICore 异常非零退出，但 bench_test.sh 仍报告
[PASSED]，并在 --skip-pytest 模式下返回 0。
补充信息
该 example 使用 Ascend C 后端和显式 UB 分配，相关 pass 配置为：
pass_configs = {
tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: True,
tilelang.PassConfigKey.TL_ASCEND_MEMORY_PLANNING: False,
}成功标记的可靠性
当前 bench 只判断输出中是否曾出现 Test passed! 或 Kernel Output Match!，无法区分
逐 case/阶段进度和整个脚本的最终完成。增加退出码检查可以修复本次 AICore 异常造成的
误报，但如果后续阶段被错误跳过并正常退出，早期 marker 仍可能导致假阳性。由于现有
examples 的 marker 数量并不统一，修复不应依赖固定匹配次数，而应明确区分中间进度和最终
完成信号。
Ping-pong 修复约束
旧实现为输入和转换 buffer 保留了两个交替 slot，生成代码也保留了相应寻址：
x_cast[by * 8192]
dy_cast[by * 8192]
两个 slot 在理论上可以支持相邻迭代的数据搬运与 Vector 计算重叠。当前
AUTO_SYNC 插入的 V_MTE2 event 消除了实际 overlap，因此两个 slot 仍占用 UB，
却没有形成预期流水。
将双 slot 折叠为单 slot 可以规避当前越界，但可能丢失未来编译器生成真正 ping-pong 流水的
优化空间。因此，本 Issue 不预设具体的 buffer 修复方案；最终修复应同时保证 UB 访问合法，
并避免无依据地放弃双缓冲的优化语义。
Bug 描述
在 Ascend A3 上运行
examples/normalization/rms_norm.py时，_rms_norm_grad(1024, 1024, 128, 128, dtype="float")能完成编译，但生成的kernel 存在 UB 越界访问，执行时触发 AICore 异常。
rms_norm.py在进入 grad 测试前已经输出Test passed!和Kernel Output Match!。examples/bench_test.sh将输出中任意位置出现的这些文字视为成功，且不要求子进程退出码为 0，因此后续 grad 触发 AICore 异常后仍被报告为
[PASSED]。这会导致全量 examples 显示 142/142 通过，但其中实际包含未完成全部测试的用例。
环境
d2bd9ec7c97a08db31454485c978ae177a06c56f复现代码
直接运行 example，可以观察到 grad 阶段的运行时异常：
通过旧 bench 运行相同目录，则会被误报为通过：
失败配置为：
错误信息
NPU 执行时报告：
通过
kernel.get_kernel_source()检查生成的 Ascend C，可以看到：其结束地址为：
A3 UB 上限为 196352 字节，因此越界 576 字节。编译阶段没有报告容量错误。
旧
bench_test.sh的成功判定近似为：普通 example 的退出码没有参与成功判断。由于 forward 完成后已经输出
Kernel Output Match!，后续 grad 阶段的 AICore 异常不会改变 bench 的判定结果。期望行为
判为通过。
_rms_norm_grad不应生成有效访问范围超过 UB 容量的 kernel。希望对
rms_norm.py和bench_test.sh进行一次原子修复，使全量 examples 的成功结果能够真实表示所有测试阶段均已完成。
实际行为
_rms_norm_grad在执行时因 AICore 异常非零退出，但bench_test.sh仍报告[PASSED]，并在--skip-pytest模式下返回 0。补充信息
该 example 使用 Ascend C 后端和显式 UB 分配，相关 pass 配置为：
成功标记的可靠性
当前 bench 只判断输出中是否曾出现
Test passed!或Kernel Output Match!，无法区分逐 case/阶段进度和整个脚本的最终完成。增加退出码检查可以修复本次 AICore 异常造成的
误报，但如果后续阶段被错误跳过并正常退出，早期 marker 仍可能导致假阳性。由于现有
examples 的 marker 数量并不统一，修复不应依赖固定匹配次数，而应明确区分中间进度和最终
完成信号。
Ping-pong 修复约束
旧实现为输入和转换 buffer 保留了两个交替 slot，生成代码也保留了相应寻址：
两个 slot 在理论上可以支持相邻迭代的数据搬运与 Vector 计算重叠。当前
AUTO_SYNC插入的V_MTE2event 消除了实际 overlap，因此两个 slot 仍占用 UB，却没有形成预期流水。
将双 slot 折叠为单 slot 可以规避当前越界，但可能丢失未来编译器生成真正 ping-pong 流水的
优化空间。因此，本 Issue 不预设具体的 buffer 修复方案；最终修复应同时保证 UB 访问合法，
并避免无依据地放弃双缓冲的优化语义。