[Bug] SIMT 指令(MSCATTER/MGATHER)通过 SU dispatcher fn-pointer 路径调用时 chip hang

[ChaoWao]

Platform

a5 (Ascend 950 hardware)

Runtime Variant

tensormap_and_ringbuffer

Description

在 a5 平台 tensormap_and_ringbuffer 运行时上,如果一个 user incore kernel 在内部直接或间接使用任何 SIMT 类指令(pto::MSCATTER、pto::MGATHER 及其内部展开的 cce::async_invoke<...>),通过当前 SU dispatcher 模型(src/a5/runtime/tensormap_and_ringbuffer/aicore/aicore_executor.cpp)调用进入这个 kernel 之后,SIMT scheduler 无法在 AICore 上正确启动,表现为 chip 静默 hang(host 上 stream sync 不返回,task-submit 超时被 kill,无 errcode 抛出)。

PR #764(Fix: a5 AICore SIMT launch — set localMemorySize + inject SIMT TLVs)已经把 launch-registration 阶段的 cfg.localMemorySize = 216 KB + ELF .ascend.meta.<func> 的 5 条 SIMT TLV(COMPILER_ALLOC_UB_SIZE, SU_STACK_SIZE, SIMT_WARP_STACK_SIZE, SIMT_DVG_WARP_STACK_SIZE, AIV_TYPE_FLAG=MIX_VF)装好,runtime 不再抛 ACL_ERROR_RT_PARAM_INVALID (107000),launch 注册成功。但进入 user kernel body 执行到 MSCATTER 时仍然 hang。

根因:bisheng 把 dispatcher aicore_executor.cpp:42 处通过 payload->function_bin_addr 加载的 fn-pointer call lowering 成 HiIPUISD::LongCALL(4 字编码,目标地址通过 linker reloc 填入,decode 期不可解析),而 SIMT scheduler 的启动需要 decode 期能解析目标 PC 才能预热。短分支 HiIPUISD::CALL(3 字 + 立即数偏移)路径上,bisheng 自动发了一组 SIMT-aware setup 指令(0x0c.. + 0x1c.. 系列),LongCALL 路径不发。bisheng 也未把这组 setup 暴露成 user-callable builtin。

详细实验记录会在评论里补。

Steps to Reproduce

1. checkout PR Add: a5 AICore SIMT (mscatter) launch support #764 (ChaoZheng109:fix-a5-aicore-simt-tlv),包含 tests/st/a5/tensormap_and_ringbuffer/simt_basic/
2. 在 a5 实机上跑 simt_basic:

   task-submit --timeout 600 --max-time 600 --device auto --device-num 2 \
     --run "python -m pytest \
       tests/st/a5/tensormap_and_ringbuffer/simt_basic/test_simt_basic.py \
       --platform a5 --device \$TASK_DEVICE -v -s"

3. 观察到:
   • Orchestration .so 编译 ✅ (583KB)
   • AIV kernel kernel_simt_scatter.o 编译 ✅ (197KB,含 SIMT TLV)
   • 12 个 CANN runtime 线程 spawn(说明 runtime 已 init,chip handle 已开)
   • 接下来在 chip 上0 输出持续 600s,task-submit 超时 kill

Expected Behavior

simt_basic 用例应在硬件上完成 8x32→256 elements 的 MSCATTER,golden 比对 PASS,跟 pto-isa 自家 tests/npu/a5/src/st/testcase/mscatter 在同硬件上的 5332ms PASS 形态一致。

具体:

• AIV 三个 block(block_dim=3)各自的 MSCATTER 应该把 8x32=256 float 元素按 identity 索引散写到 256-slot dst,使 out == src
• max diff = 0,err count = 0
• 任务在 ~10 秒内完成,不应该出现长时间静默

Actual Behavior

[npu-lock] 获取设备 0 的锁 (无超时)...
[npu-lock] 已获取设备 0 的锁 (pid=...)
[npu-lock] 获取设备 1 的锁 (无超时)...
[npu-lock] 已获取设备 1 的锁 (pid=...)
============================= test session starts ==============================
platform linux -- Python 3.10.20, pytest-9.0.3, pluggy-1.6.0
collected 1 item

tests/st/a5/tensormap_and_ringbuffer/simt_basic/test_simt_basic.py::TestSimtBasic::test_run
  [V5] [Orchestration] Compilation .../simt_basic_orch.cpp.orch_*.so successful: 583600 bytes
  [V5] [Incore] Compilation .../kernel_simt_scatter.cpp.incore_*.o successful: 197840 bytes
  ───── 之后整整 600s 零输出 ─────
错误: 等待超时 (600s)

对照参考:

Test	路径	结果
simpler simt_basic(本 issue 现象)	SU dispatcher → fn-pointer → user kernel → MSCATTER	❌ 600s hang
pto-isa tests/npu/a5/.../mscatter(我用同一台 a5 硬件验证)	host <<<1, nullptr, stream>>> → kernel(里头 MSCATTER)	✅ 5332ms PASS

两者跑的是同一段 MSCATTER 内核代码,唯一差别是 launch + dispatch 路径。

Git Commit ID

86b633c

CANN Version

9.1.T500 (V100R001C10B813)

Driver Version

25.6.rc1.b108 (ascendhal 7.35.23)

Host Platform

Linux (aarch64)

Additional Context

• Related PR: Add: a5 AICore SIMT (mscatter) launch support #764(引入 simt_basic 测试 + 给 dispatcher 注入 SIMT TLV)
• 上游参考:pto-isa tests/npu/a5/src/st/testcase/mgather/MGATHER.md §"Runtime Dispatch Requirement"(line 377-403)已经把这条限制写明:cce::async_invoke 需要 launch path 在 kernel 进入前装好 TID/warp/vec-pipe scheduler 状态,标准 rtKernelLaunch 这么做了,fn-pointer 直接调用 skips 这一步,所以第一条 async_invoke 起不来 warp 调度
• pto-isa 给出三条 fix 路径(MGATHER.md line 400-402):全部需要 simpler dispatcher 侧改造
• 我们做了一整套定位实验(byte-level disasm 对比、bisheng intrinsic 全扫、host-side builtin 注入、SDAG 节点分析),会在下面评论里补全

[ChaoWao]

已做的定位实验汇总

为了把"是不是这条路径"从模型推断变成可证伪的实测,在 pto-isa tests/npu/a5/src/st/testcase/mscatter/mscatter_kernel.cpp 里改写 runMSCATTER_elem2d_float_8x32_random_256size 这一个 case,跑同一台 a5 onboard,通过 task-submit 重复采集。pto-isa 那个 case 本来就是用 <<<1, nullptr, stream>>>(标准 host launch)调一段单一函数体里直接调 MSCATTER —— 是 simpler dispatcher path 的最小对照。

阶段 1 — 排除一些"看起来像但其实不是"的原因

测试	实测结果
baseline:standard DIRECT call → runElem2D<...>(out, src, indices)	✅ PASS 5332ms,max diff 0
Test A:100 × pipe_barrier(PIPE_ALL) 之后再 MSCATTER	✅ PASS
Test B:100 × set_flag/wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0)	✅ PASS
Test C:100 × dsb(DSB_DDR)	✅ PASS
Test D:dcci(out, SINGLE_CACHE_LINE) 真实 GM 地址,SINGLE_CACHE_LINE 模式	✅ PASS
Test E:复测 dcci(0, SINGLE_CACHE_LINE) —— 地址 0	❌ 161/256 错(63%)
Test F:dcci(out, ENTIRE_DATA_CACHE) —— simpler dispatcher 实际用的形式	✅ PASS
Test G:模拟整段 simpler polling loop 50 × (read DATA_MAIN_BASE + set_cond + dcci out ENTIRE_DATA_CACHE)	✅ PASS

Test E vs F 的对比是个关键 finding —— 之前怀疑过 dcci 这条指令污染 SIMT,但实测 dcci 在真实地址上完全无害。dcci(0) 触发的是地址 0 非法路径下的 fault 行为,跟 dcci 这条指令的语义无关。simpler dispatcher 用的 dcci(exec_payload, ENTIRE_DATA_CACHE) 操作完全 OK,不是 root cause。

→ 排除掉:dcci、pipe barrier、pipe sync、polling loop 的 SU 操作、整套 SPMD prelude。这些都不破坏后续 SIMT。

阶段 2 — 定位到 fn-pointer 这一步

测试	实测结果
direct call(bisheng 静态可解析 callee)	✅ PASS 5332ms
direct call,callee 加 __attribute__((noinline))	✅ PASS 6123ms
static const FnT k_fn = ⌖ k_fn(args)(const fn-pointer,bisheng 编译期 devirtualize)	✅ PASS 5075ms
volatile FnT fn = ⌖ fn(args)(强制走 memory roundtrip → 必发 indirect call)	❌ 200s task-submit timeout(chip hang)
同上,__attribute__((aiv)) 加在 fn-pointer 类型上	❌ 同 hang
同上,__attribute__((musttail)) 强制 tail-call(BR 而非 BLR)	❌ 同 hang
volatile fn-pointer 到非 SIMT body(out[0] = 42 之类)	✅ 调用本身完成(用例因 golden 不匹配 fail 但 chip 没 hang)

最后一项是决定性证据:同样的 volatile fn-pointer 间接调用,被调函数体里不带 SIMT 指令时调用完成、chip 干净;带 MSCATTER 时立刻 hang。证明 hang 不是 fn-pointer 这条 BLR 指令本身导致的,是 BLR + 被调函数体里的 SIMT execution 这个组合致命。

阶段 3 — 字节级反查 bisheng codegen 差异

extract .aicore_binary section 从 DIRECT 和 INDIRECT(volatile fn-pointer)两版的 .so,逐字对比 runMSCATTER_elem2d_float_8x32_random_256size 函数体:

• DIRECT 版本:232 字节
• INDIRECT 版本:256 字节(多 24 字节 = 6 条指令)
• 前 60 字节(15 条指令)的 prologue 完全相同
• 末尾 ~50 字节的 epilogue 完全相同
• 中间差 = call lowering 差异

DIRECT 的 call 序列(7 条 + 2 条 pre-call setup):
```
1cc01000 ← 0x1c LD-class,bisheng 自动发,INDIRECT 不发
1c821008 ← 0x1c LD-class,bisheng 自动发,INDIRECT 不发
0706ffb0 ┐
0745ffff │ 3-word short BL,immediate offset 已编译期填好
0785ffff ┘
40423000 ← 1 条 post-call control/sync
```

INDIRECT 的 call 序列(13 条):
```
0706fec4 ┐
0745ffff │
0785ffff │ 4-word long form indirect call,目标地址通过 linker reloc 填入
07c7ffff ┘ (.rela.text 里有一条 entry 指向 target_body 符号,type 0x15c)
02080880 ┐
00063001 │
1cc01000 │ post-call cleanup(7 条,补 BLR 蹭脏的 spill 状态)
0cc21110 │
03c5e878 │
1cc5e878 │
02440282 ┘
40443000 ← post-call control/sync
```

.rela.text 在 INDIRECT 版本里有 1 条 entry,offset = call site,symbol = target,reloc type = 0x15c(unknown to standard tooling,推断为 HiIPU 内部"4 字 long-form 调用目标填入");DIRECT 版本无 .rela.text(call site 全部是立即数偏移)。

bisheng 内部 SDAG 节点(从 bisheng binary strings 扒出):

• HiIPUISD::CALL — 短分支
• HiIPUISD::LongCALL — 长分支
• HiIPUISD::BUFFER_SYMBOL_MOV16 / MOV32 — 长分支用来 materialize 目标地址的辅助节点
• HiIPUISD::UBUF_MOV_SIMD vs HiIPUISD::UBUF_MOV_SIMT — 同一个 UB move 操作分两个独立 SDAG 节点,证明 bisheng 在 SDAG 层就分了 SIMD/SIMT codegen path

bisheng 选 CALL 还是 LongCALL 的依据只看"TU compile time 是否能解析目标 PC offset",跟函数属性无关:可 devirtualize 的 const fn-pointer 也走 CALL,无法解析的 volatile fn-pointer 必走 LongCALL。-mllvm --cce-aicore-long-call=true、-mlong-calls、-mllvm --cce-long-call=true 这几个 flag 在 dav-c310-vec 上实测均无效(产物字节 MD5 完全一致),所以 ISel 这条 rule 是 hardcoded 在 bisheng 二进制里。

→ simpler aicore_executor.cpp:42 的 kernel(args) 是从 cached GM payload 读出来的 64-bit 运行时值,bisheng 永远会选 LongCALL,无论怎么改 flag/attribute 都改不掉。

阶段 4 — Path Y 探索(不静态 link,device 端注入 setup)

DIRECT 多发的那 2 条 0x1c LD-class 指令是经验上的 hang prevention key。试图用 device-side builtin 复现:

已扫完的 builtin 清单(bisheng cce_aicore_intrinsics.h 全表)

Builtin	a5 (dav-c310-vec) 可用?	注入到 INDIRECT call 之前,实测
__builtin_cce_store_vfsimt_info(uint64)	✅	hang(无改善)
__builtin_cce_store_vfsimt_info(packed dim3{32,32,1}) 按 MSCATTER 内部 grid	✅	hang(无改善)
__builtin_cce_dc_preload(__gm__ ulong*, ulong)	✅	hang(无改善)
__builtin_cce_preload(const void*, ulong)	✅(stack 指针)	没单独跑
__builtin_cce_kernel_type_set(int)	❌ frontend 拒(必须用 __vector__/__cube__/__mix__ attribute)	—
__builtin_cce_set_kernel(...)	❌ "does not support the given target feature"	—
__builtin_cce___ib_set_stub(int, int, bool)	❌ target-feature gated	—
__builtin_cce_mvf_dci()	❌ target-feature gated	—
__builtin_cce_mov_vspr(VSPR_t, ...)	✅ 但 VSPR enum 全是 RPN/DATA_EXP_*,无 SIMT 相关	没跑(无相关枚举)
__builtin_cce_init_vector_align_data()	✅ 但返回 vector_align 类型构造数据,非 setup 操作	没跑(语义不对路)
手工 inline asm .4byte 0x0cc02080 / 0x1cc42010(从 DIRECT 编译产物里抠的 bytes)	✅(汇编层)	✓ 不 hang,但 161/256 错
Combo:store_vfsimt + dc_preload + raw bytes	✅	不 hang,仍 161/256 错(跟 raw bytes 单独效果一致)

关键 empirical 事实

1. 唯一能消 hang 的是手工注入 DIRECT 编译产物里 bisheng 自己发的那 2 条 bytes;其他所有 user-callable builtin 单独/组合均无任何错误率改善。
2. 注入 raw bytes 之后错误率精确固定在 161/256(63% 写丢失,act zero 比例完全一致),不论隔多远、组合啥都一样 → 这 2 条指令写的是 sticky HW state(像 SPR 写),不是 cacheline preload 这种 ephemeral 操作。
3. 161/256 这个 floor 用 user-level 任何注入组合都打不下来。剩下的差距在 BL imm vs BLR reg 这个指令本身的 decode-period 语义上(操作数为 DIRECT 上下文编码,套到 INDIRECT 的寄存器布局上无意义)。

bisheng 内部能看见但用户无 builtin 可调的相关 IR intrinsic

从 bisheng binary strings 扒出:

• llvm.hivm.cce.tinit —— 名字强烈像 "thread/task init"
• llvm.hivm.SET.KERNEL.SPRb8 —— 写 KERNEL SPR bit 8
• llvm.hivm.DC.PRELOAD.IMM.STACK.C310 —— a5 (C310) 专属栈 preload(跟 __builtin_cce_dc_preload lower 出来的不是同一个变体)

这些是 bisheng 在某个 codegen pass 里自动发出来的,没有暴露给用户的 __builtin_cce_* 接口。

阶段 5 — 其他相关发现

extract_text_section 的潜在脆弱性(不是本 issue 的 root cause,但相关)

simpler_setup/elf_parser.py::extract_text_section 在 onboard 模式下(scene_test.py:665)对 user incore .o 做了 strip,只保留 .text 段,丢弃所有 .ascend.meta.* TLV 段。

实测当前 kernel_simt_scatter.cpp 编出来的 user kernel .o:

• .text 0x278 = 632 字节,包含 kernel_entry 和 simt_mscatter_kernel(WEAK GLOBAL,MSCATTER 模板实例化的 SIMT body)
• 无 .rela.text(自包含)
• 无 .ascend.meta.* 段(bisheng 因 kernel_entry 未被标 __global__ 没发)

所以 strip 对当前这个测试无害。但这是脆弱设计:

• 一旦 user kernel 引用任何跨段符号(WEAK extern、全局变量、constpool),.rela.text 一丢就静默坏
• 一旦 bisheng 给 __simt_vf__ 发 .ascend.meta.<func> 让 firmware/HW 读,丢掉就静默坏

跟 root cause 无关,但建议留意。

CANN runtime 源码确认 PR #764 launch math 正确

clone gitcode.com/cann/runtime.git 走通了从 rtKernelLaunchWithHandleV2 → ApiImplDavid::KernelLaunchWithHandle → StreamLaunchKernelWithHandle → CheckAndGetTotalShareMemorySize 的完整链路:

```cpp
// src/runtime/core/src/launch/aix_starsv2.cc
const uint32_t kernelVfType = kernel->KernelVfType_(); // 从 ELF TLV 读
bool simtFlag = (kernelVfType == AIV_TYPE_SIMT_VF_ONLY) ||
(kernelVfType == AIV_TYPE_SIMD_SIMT_MIX_VF);
totalSmSize = kernel->ShareMemSize_() + dynamicShareMemSize; // 8KB + 216KB = 224KB
maxSmSize = simtFlag ? RT_SIMT_REMAIN_UB_SIZE : RT_SIMT_UB_SIZE; // 224KB
// strict > check → 224 > 224 false → pass
```

PR #764 的 launch 配置在 runtime 注册期完全正确,不再触发 107000。这一层不是 bug 所在层。

pto-isa 上游已经知道这个限制

tests/npu/a5/src/st/testcase/mgather/MGATHER.md §"Runtime Dispatch Requirement" 已经描述:

Dispatch kernels as a direct C function-pointer call:
```cpp
UnifiedKernelFunc kernel = (UnifiedKernelFunc)payload->function_bin_addr;
kernel(reinterpret_cast<gm int64_t *>(payload->args));
```
This is fine for SPMD ops (TLOAD, TSTORE, TADD …) but skips the SIMT-context init step, so the first cce::async_invoke inside MGATHER has no warp scheduler to dispatch into and hangs.

并给出三条 fix paths(line 400-402),全部需要 simpler dispatcher 侧改造:

1. PTO2DispatchPayload 加 is_simt flag,SIMT payload 走 host rtKernelLaunch
2. 编译期检测 SIMT-bearing kernel(ELF attribute),自动路由
3. 在 dispatcher 直接 fn-pointer call 之前device-side 重装 SIMT context

实验已经验证 (3) 这条在当前 bisheng 暴露的 builtin 集合下不可达。

已经 ruled out 的 fix 路径

• ❌ 改 dcci 模式 / 改 polling loop SU 操作 —— 实测均无害,不是 root cause
• ❌ 改 simt_meta.h 里的 TLV 值(尝试过 SIMT_VF_ONLY=3 替代 MIX_VF=4,SU dispatcher 第一条 SPMD 指令立刻 trap)
• ❌ -mllvm flag 强制 short / long call —— 对 dav-c310-vec 实测全部无效
• ❌ __attribute__((long_call)) / ((short_call)) / ((aiv)) 在 fn-pointer 类型 —— 实测全部无效(产物字节完全一致)
• ❌ device-side __builtin_cce_* 注入修复 —— 上面表里的所有可用 builtin 都试过,只能消 hang 不能改 161/256 错率

剩下的可行路径

路径	描述	已知阻塞点
Path A	SIMT-bearing task 走独立 rtKernelLaunch,绕过 SU dispatcher	性能代价:每 SIMT task 一次 host launch(~50μs vs ringbuffer dispatch)
Path X	dispatcher 跟 user incore co-link,通过 extern __aicore__ void user_task_kernel(args) weak symbol + switch table 走 bisheng 短分支	build model 改造,每个 test build 重链 aicore_kernel.o(其他 .o 复用 cache);pypto 已经在用类似套路(framework/src/machine/compile/aicore_compiler.cpp::GenSubFuncCall 生成 switch-by-funcIdx,然后 ld.lld -m aicorelinux -static 静态链接)
Path Y	在 dispatcher 内 device-side patch SIMT scheduler 状态	需要 bisheng 暴露 cce.tinit 或者类似 builtin,或开放 -mllvm -cce-emit-simt-setup-on-indirect-call=true flag

推荐处理

1. 短期:simt_basic 测试加 _skip_reason="blocked on indirect-call SIMT scheduler init (see #970)",不阻塞 CI
2. 中期(架构选型):团队决定走 Path A 还是 Path X
3. 长期:跟 bisheng 团队提需求,把 cce.tinit / SIMT call prelude 暴露成 user-callable builtin(惠及所有自定义 dispatcher 架构,不止 simpler)

PR #764 的 cfg.localMemorySize + 5 条 SIMT TLV 注入本身是正确的,应该保留(任何路径都需要这层);simt_basic 的测试设计需要根据架构选型微调,但不该撤回。

cc @ChaoZheng109 (PR #764 作者)

[ChaoWao]

补:最 clean 的 standalone repro —— BLR before SIMT,跟 simpler dispatcher 完全无关

为了把"是 BLR 这条指令本身 + SIMT body 的组合致命"这个结论钉死,不让"是不是 simpler dispatcher 的别的东西在作怪"留任何尾巴,刚刚跑了一个最小化对照:

Test setup

在 pto-isa tests/npu/a5/src/st/testcase/mscatter/mscatter_kernel.cpp 里,保留 pto-isa 原版 <<<1, nullptr, stream>>> 标准 launch 路径(= rtKernelLaunch),只改 runMSCATTER_elem2d_float_8x32_random_256size 的 body 加一层 BLR 中转:

// pto-isa 原版的 launch wrapper,host 侧 <<<1, nullptr, stream>>> 不动
AICORE __attribute__((noinline)) void simt_body_for_disasm(
    __gm__ float *out, __gm__ float *src, __gm__ int32_t *indices)
{
    // SIMT body —— 调 MSCATTER (内部展开 cce::async_invoke + simt_kernel)
    runElem2D<pto::ScatterAtomicOp::None, pto::ScatterOOB::Undefined,
              pto::ScatterConflict::Last, float, int32_t, 8, 32, 256>(out, src, indices);
}

extern "C" __global__ AICORE void runMSCATTER_elem2d_float_8x32_random_256size(
    __gm__ float *out, __gm__ float *src, __gm__ int32_t *indices)
{
    using FnT = void (*)(__gm__ float *, __gm__ float *, __gm__ int32_t *);
    volatile FnT fn = &simt_body_for_disasm;  // volatile 强制 bisheng 发 HiIPUISD::LongCALL
    fn(out, src, indices);                      // 这一行就是 BLR;然后 SIMT body 开始
}

没注入任何 raw bytes,没用任何额外 builtin,只是把 SIMT body 套了一层 noinline 函数 + 一个 volatile fn-pointer 跳进去。pto-isa host 侧 launcher 一行没动。

Result

[npu-lock] 获取设备 2 的锁 (无超时)...
[npu-lock] 已获取设备 2 的锁 (pid=3208750)
Note: Google Test filter = MSCATTERTest.case_elem2d_float_8x32_random_256size
[==========] Running 1 test from 1 test suite.
[ RUN      ] MSCATTERTest.case_elem2d_float_8x32_random_256size
错误: 等待超时 (200s)        ← chip hang

跟对照组对比:

Test	launch path	call into SIMT body 形态	SIMT 在 body 里	结果
pto-isa 原 baseline	<<<1, nullptr, stream>>>	inline 展开,无 BLR	✓ MSCATTER	✅ PASS 5332ms
本 repro	<<<1, nullptr, stream>>>(完全同一条)	加一层 volatile FnT; fn(...) → BLR	✓ MSCATTER	❌ 200s hang
(之前 #970 评论里) BLR 到非 SIMT body	<<<1, nullptr, stream>>>	volatile fn-pointer BLR	✗ out[0]=42 标量 store	✅ 调用完成 (golden 不匹配 fail 但 chip 不 hang)

这次实测把什么钉死

1. 跟 simpler dispatcher 完全无关 —— 没有 SU polling loop、没有 dcci、没有 AICPU↔AICore ringbuffer dispatch,纯 pto-isa 原版 launch 流程
2. 跟 launch 路径无关 —— 用的就是 pto-isa 自己验证 OK 的 <<<...>>> (rtKernelLaunchWithHandleV2),firmware 装 SIMT context 这一步跟 baseline 一样
3. 跟 dcci / SPMD prelude / polling loop SU 操作无关 —— body 里啥都没,就一个 volatile fn-pointer assignment + indirect call
4. 唯一变量是这一行 indirect call —— 把它换成 direct call(simt_body_for_disasm(out, src, indices))立即 PASS,加 volatile 让它走 BLR 立即 hang
5. 跟 simpler 的 LongCALL 在 hardware level 是完全一回事:bisheng 都是发 HiIPUISD::LongCALL 长形式 4-word indirect call,目标地址 linker reloc 填入

→ 这是 bisheng codegen + HiIPU 硬件 SIMT scheduler 启动机制的固有限制,不是 simpler dispatcher 设计上的 bug,任何代码只要"运行时 fn-pointer 间接调用进入一个会跑 SIMT 指令的函数"都会撞上。

这个 repro 的价值

• 可以直接给 bisheng/CANN 团队看,不需要任何 simpler 上下文就能 reproduce
• 把"问题层在 SIMT API + indirect call ABI"这层钉死,后续要 push 上游的话有可附的 standalone case

cc @ChaoZheng109 (PR #764) — 这次实验进一步确认 simt_basic hang 跟 simpler dispatcher 设计无关,是 BLR+SIMT 的硬性组合;PR #764 的 TLV/localMemorySize 配置都对,不需要回退,只是 simt_basic 测试本身需要等 indirect call SIMT 路径被解决(或换 Path A/X 之一)才能跑得通。

[ChaoZheng109]

Closing — this is resolved by PR #764 (merged 2026-06-11).

The hang turned out not to be the HiIPUISD::LongCALL decode-time-PC limitation hypothesized in the comments above (those were written on 06-02, before the fix landed on 06-11). The actual root cause was child-kernel binary alignment: the SIMT kernel binary was landing at an 8-mod-16 address, and a5 SIMT kernels (mscatter et al.) require a 16-byte-aligned code address, otherwise the chip hangs silently.

PR #764 fixes it with two load-bearing pieces, both required:

1. SIMT TLVs — cfg.localMemorySize = 216 KB + the SIMT TLVs in .ascend.meta (UB-size, AIV_TYPE_FLAG = MIX_VF, …), so launch registration succeeds and no longer throws ACL_ERROR_RT_PARAM_INVALID (107000).
2. 16-byte child-kernel alignment — CALLABLE_CHILD_ALIGN = 16 (src/common/task_interface/callable.h), so storage_ and the kernel binary inside it land at a 16-byte-aligned address. This is what actually clears the MSCATTER hang.

With both in place the full st-onboard-a5 suite passes on hardware, including tests/st/a5/tensormap_and_ringbuffer/simt_basic/ (golden match, no hang). The SU dispatcher fn-pointer path itself is fine — no dispatcher redesign / Path A/X needed.

Closing as fixed. (#156 keeps the SIMT-via-indirect-call regression coverage.)