[Bug] [Compiler] SIMTVF 循环内 if 分支导致 END 指令带谓词

[KurrinQu]

问题概述

在 VPTO SIMTVF kernel 中，如果 scf.for 循环体内包含 scf.if 分支，编译后的 AICore 反汇编中可能出现带谓词的 END 指令，例如：

END @!P0

SIMTVF 的限制是：END 指令不应出现在分支路径中，也不应通过分支跳转后条件执行。否则未执行 END 的线程可能继续向后执行，硬件行为不可预期。

该现象发生在编译器处理 SIMTVF 控制流之后，生成代码不符合上述约束。

1. 最小复现用例

失败场景是：SIMTVF 内循环 16 次，但只有前 10 次写 UB。

完整 PTO 复现如下：

module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
func.func @simt_vf_0(%arg0: !pto.ptr<i8, ub>, %arg1: f32, %arg2: i32)
    attributes {pto.simt_entry, pto.simt_max_regs = 64 : i32} {
  %tid = pto.get_tid_x : i32

  %c0 = arith.constant 0 : index
  %c16 = arith.constant 16 : index
  %c1 = arith.constant 1 : index
  %c0_i32 = arith.constant 0 : i32
  %c1_i32 = arith.constant 1 : i32
  %c2_i32 = arith.constant 2 : i32
  %c10_i32 = arith.constant 10 : i32
  %c512_i32 = arith.constant 512 : i32
  %c8192_i32 = arith.constant 8192 : i32
  %c21520_i32 = arith.constant 21520 : i32
  %c4_i64 = arith.constant 4 : i64

  %fpconst = arith.constant 1.250000e+00 : f32
  %undef = llvm.mlir.undef : vector<2xf32>
  %ins0 = "llvm.insertelement"(%undef, %fpconst, %c0_i32)
      : (vector<2xf32>, f32, i32) -> vector<2xf32>
  %val = "llvm.insertelement"(%ins0, %fpconst, %c1_i32)
      : (vector<2xf32>, f32, i32) -> vector<2xf32>

  %ub_int = "pto.castptr"(%arg0) : (!pto.ptr<i8, ub>) -> i64
  %ub_llvm = "llvm.inttoptr"(%ub_int) : (i64) -> !llvm.ptr<6>

  scf.for %i = %c0 to %c16 step %c1 {
    %i32 = arith.index_cast %i : index to i32
    %active = arith.cmpi slt, %i32, %c10_i32 : i32

    scf.if %active {
      %eid = arith.andi %arg2, %c1_i32 : i32
      %ping = arith.muli %eid, %c8192_i32 : i32
      %chunk = arith.muli %i32, %c512_i32 : i32
      %base0 = arith.addi %ping, %chunk : i32
      %tid2 = arith.muli %tid, %c2_i32 : i32
      %base1 = arith.addi %base0, %tid2 : i32
      %idx = arith.addi %base1, %c21520_i32 : i32
      %idx_index = arith.index_cast %idx : i32 to index
      %idx_i64 = arith.index_cast %idx_index : index to i64
      %byte_off = arith.muli %idx_i64, %c4_i64 : i64
      %ptr = "llvm.getelementptr"(%ub_llvm, %byte_off)
          {elem_type = i8, rawConstantIndices = array<i32: -2147483648>}
          : (!llvm.ptr<6>, i64) -> !llvm.ptr<6>
      "llvm.store"(%val, %ptr) : (vector<2xf32>, !llvm.ptr<6>) -> ()
    }
  }

  pto.syncthreads
  return
}

func.func @main_kernel_kernel(%arg0: !pto.ptr<f32, gm>, %arg1: !pto.ptr<f32, gm>,
                              %arg2: !pto.ptr<f32, gm>, %arg3: !pto.ptr<f32, gm>,
                              %arg4: f32)
    attributes {dyn_shared_memory_buf = 152640 : i64, pto.aicore} {
  %c0_i64 = arith.constant 0 : i64
  %ub_base = pto.castptr %c0_i64 : i64 -> !pto.ptr<i8, ub>
  %bx = "pto.get_block_idx"() : () -> i64
  %core = arith.trunci %bx : i64 to i32

  %c0_i32 = arith.constant 0 : i32
  %c1_i32 = arith.constant 1 : i32
  %c64 = arith.constant 64 : index
  %c256_i32 = arith.constant 256 : i32
  %c5120_i32 = arith.constant 5120 : i32
  %c8192_i32 = arith.constant 8192 : i32
  %c21520_i32 = arith.constant 21520 : i32
  %c327680_i32 = arith.constant 327680 : i32
  %c4 = arith.constant 4 : index
  %c0 = arith.constant 0 : index
  %c1 = arith.constant 1 : index
  %c1_i64 = arith.constant 1 : i64
  %c20480_i64 = arith.constant 20480 : i64

  "pto.set_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                    dst_pipe = #pto.pipe<PIPE_V>,
                    event_id = #pto.event<EVENT_ID0>} : () -> ()
  "pto.set_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                    dst_pipe = #pto.pipe<PIPE_V>,
                    event_id = #pto.event<EVENT_ID1>} : () -> ()

  scf.for %t = %c0 to %c64 step %c1 {
    %t_i32 = arith.index_cast %t : index to i32
    %eid = arith.andi %t_i32, %c1_i32 : i32
    %eid_index = arith.index_cast %eid : i32 to index

    "pto.wait_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_MTE3>,
                                     dst_pipe = #pto.pipe<PIPE_V>} : (index) -> ()
    pto.simt_launch @simt_vf_0<<<%c256_i32, %c1_i32, %c1_i32>>>
        (%ub_base, %arg4, %t_i32) : (!pto.ptr<i8, ub>, f32, i32) -> ()
    "pto.set_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_V>,
                                    dst_pipe = #pto.pipe<PIPE_MTE3>} : (index) -> ()
    "pto.wait_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_V>,
                                     dst_pipe = #pto.pipe<PIPE_MTE3>} : (index) -> ()

    %t_stride = arith.muli %t_i32, %c327680_i32 : i32
    %core_stride = arith.muli %core, %c5120_i32 : i32
    %gm_idx_i32 = arith.addi %t_stride, %core_stride : i32
    %gm_idx = arith.index_cast %gm_idx_i32 : i32 to index
    %gm_y = "pto.addptr"(%arg3, %gm_idx)
        : (!pto.ptr<f32, gm>, index) -> !pto.ptr<f32, gm>

    %ping = arith.muli %eid, %c8192_i32 : i32
    %ub_idx_i32 = arith.addi %ping, %c21520_i32 : i32
    %ub_idx = arith.index_cast %ub_idx_i32 : i32 to index
    %byte_off = arith.muli %ub_idx, %c4 : index
    %ub_elem = "pto.addptr"(%ub_base, %byte_off)
        : (!pto.ptr<i8, ub>, index) -> !pto.ptr<i8, ub>
    %ub_y = "pto.castptr"(%ub_elem) : (!pto.ptr<i8, ub>) -> !pto.ptr<f32, ub>

    pto.mte_ub_gm %ub_y, %gm_y, %c20480_i64
      nburst(%c1_i64, %c20480_i64, %c20480_i64)
      : !pto.ptr<f32, ub>, !pto.ptr<f32, gm>, i64, i64, i64, i64
    "pto.set_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_MTE3>,
                                    dst_pipe = #pto.pipe<PIPE_V>} : (index) -> ()
  }

  "pto.wait_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                     dst_pipe = #pto.pipe<PIPE_V>,
                     event_id = #pto.event<EVENT_ID0>} : () -> ()
  "pto.wait_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                     dst_pipe = #pto.pipe<PIPE_V>,
                     event_id = #pto.event<EVENT_ID1>} : () -> ()
  return
}
}

该用例的语义很简单：每个 SIMTVF lane 按 tid 计算地址，前 10 个 chunk 写 UB，后 6 个 chunk 不写。

2. 问题反汇编示例

问题版本的反汇编中，SIMTVF 入口里存在类似如下控制流：

0000000000000000 <simt_vf_0_simt_entry>:
  ...
  70:  BRANCH  88
  78:  ISETP.NE.AND.s32  P1, PT, S8, #0, PT
  88:  BRANCH  416, !P1
  ...
  c0:  END  @!P0
  ...
  1b8: BRANCH  0x78, P1
  ...
  220: BRANCH  0x78

关键点是：

END @!P0

END 被生成成了带谓词的指令，并且位于分支控制流相关路径中。

3. 问题说明

SIMTVF 中 END 应当是所有相关线程都能确定执行的终止点。

反汇编中 END 携带了 !P0：

END @!P0

这意味着：

• !P0 为真的线程会执行 END；
• P0 为真的线程不会执行 END；
• 未执行 END 的线程可能继续跑到后续指令；
• 这违反 SIMTVF 对 END 的约束，硬件行为不可预期。

PTO 源码中没有条件终止 SIMTVF 的语义。问题出现在带分支的 SIMTVF 控制流降低之后，END 没有保持在无条件收敛路径上。

4. 可以通过的场景

如果把逻辑改成循环 10 次，并去掉 if 分支，则可以通过。

等价通过版本如下：

module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
func.func @simt_vf_0(%arg0: !pto.ptr<i8, ub>, %arg1: f32, %arg2: i32)
    attributes {pto.simt_entry, pto.simt_max_regs = 64 : i32} {
  %tid = pto.get_tid_x : i32

  %c0 = arith.constant 0 : index
  %c10 = arith.constant 10 : index
  %c1 = arith.constant 1 : index
  %c0_i32 = arith.constant 0 : i32
  %c1_i32 = arith.constant 1 : i32
  %c2_i32 = arith.constant 2 : i32
  %c512_i32 = arith.constant 512 : i32
  %c8192_i32 = arith.constant 8192 : i32
  %c21520_i32 = arith.constant 21520 : i32
  %c4_i64 = arith.constant 4 : i64

  %fpconst = arith.constant 1.250000e+00 : f32
  %undef = llvm.mlir.undef : vector<2xf32>
  %ins0 = "llvm.insertelement"(%undef, %fpconst, %c0_i32)
      : (vector<2xf32>, f32, i32) -> vector<2xf32>
  %val = "llvm.insertelement"(%ins0, %fpconst, %c1_i32)
      : (vector<2xf32>, f32, i32) -> vector<2xf32>

  %ub_int = "pto.castptr"(%arg0) : (!pto.ptr<i8, ub>) -> i64
  %ub_llvm = "llvm.inttoptr"(%ub_int) : (i64) -> !llvm.ptr<6>

  scf.for %i = %c0 to %c10 step %c1 {
    %i32 = arith.index_cast %i : index to i32
    %eid = arith.andi %arg2, %c1_i32 : i32
    %ping = arith.muli %eid, %c8192_i32 : i32
    %chunk = arith.muli %i32, %c512_i32 : i32
    %base0 = arith.addi %ping, %chunk : i32
    %tid2 = arith.muli %tid, %c2_i32 : i32
    %base1 = arith.addi %base0, %tid2 : i32
    %idx = arith.addi %base1, %c21520_i32 : i32
    %idx_index = arith.index_cast %idx : i32 to index
    %idx_i64 = arith.index_cast %idx_index : index to i64
    %byte_off = arith.muli %idx_i64, %c4_i64 : i64
    %ptr = "llvm.getelementptr"(%ub_llvm, %byte_off)
        {elem_type = i8, rawConstantIndices = array<i32: -2147483648>}
        : (!llvm.ptr<6>, i64) -> !llvm.ptr<6>
    "llvm.store"(%val, %ptr) : (vector<2xf32>, !llvm.ptr<6>) -> ()
  }

  return
}

func.func @main_kernel_kernel(%arg0: !pto.ptr<f32, gm>, %arg1: !pto.ptr<f32, gm>,
                              %arg2: !pto.ptr<f32, gm>, %arg3: !pto.ptr<f32, gm>,
                              %arg4: f32)
    attributes {dyn_shared_memory_buf = 152640 : i64, pto.aicore} {
  %c0_i64 = arith.constant 0 : i64
  %ub_base = pto.castptr %c0_i64 : i64 -> !pto.ptr<i8, ub>
  %bx = "pto.get_block_idx"() : () -> i64
  %core = arith.trunci %bx : i64 to i32

  %c0_i32 = arith.constant 0 : i32
  %c1_i32 = arith.constant 1 : i32
  %c64 = arith.constant 64 : index
  %c256_i32 = arith.constant 256 : i32
  %c5120_i32 = arith.constant 5120 : i32
  %c8192_i32 = arith.constant 8192 : i32
  %c21520_i32 = arith.constant 21520 : i32
  %c327680_i32 = arith.constant 327680 : i32
  %c4 = arith.constant 4 : index
  %c0 = arith.constant 0 : index
  %c1 = arith.constant 1 : index
  %c1_i64 = arith.constant 1 : i64
  %c20480_i64 = arith.constant 20480 : i64

  "pto.set_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                    dst_pipe = #pto.pipe<PIPE_V>,
                    event_id = #pto.event<EVENT_ID0>} : () -> ()
  "pto.set_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                    dst_pipe = #pto.pipe<PIPE_V>,
                    event_id = #pto.event<EVENT_ID1>} : () -> ()

  scf.for %t = %c0 to %c64 step %c1 {
    %t_i32 = arith.index_cast %t : index to i32
    %eid = arith.andi %t_i32, %c1_i32 : i32
    %eid_index = arith.index_cast %eid : i32 to index

    "pto.wait_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_MTE3>,
                                     dst_pipe = #pto.pipe<PIPE_V>} : (index) -> ()
    pto.simt_launch @simt_vf_0<<<%c256_i32, %c1_i32, %c1_i32>>>
        (%ub_base, %arg4, %t_i32) : (!pto.ptr<i8, ub>, f32, i32) -> ()
    "pto.set_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_V>,
                                    dst_pipe = #pto.pipe<PIPE_MTE3>} : (index) -> ()
    "pto.wait_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_V>,
                                     dst_pipe = #pto.pipe<PIPE_MTE3>} : (index) -> ()

    %t_stride = arith.muli %t_i32, %c327680_i32 : i32
    %core_stride = arith.muli %core, %c5120_i32 : i32
    %gm_idx_i32 = arith.addi %t_stride, %core_stride : i32
    %gm_idx = arith.index_cast %gm_idx_i32 : i32 to index
    %gm_y = "pto.addptr"(%arg3, %gm_idx)
        : (!pto.ptr<f32, gm>, index) -> !pto.ptr<f32, gm>

    %ping = arith.muli %eid, %c8192_i32 : i32
    %ub_idx_i32 = arith.addi %ping, %c21520_i32 : i32
    %ub_idx = arith.index_cast %ub_idx_i32 : i32 to index
    %byte_off = arith.muli %ub_idx, %c4 : index
    %ub_elem = "pto.addptr"(%ub_base, %byte_off)
        : (!pto.ptr<i8, ub>, index) -> !pto.ptr<i8, ub>
    %ub_y = "pto.castptr"(%ub_elem) : (!pto.ptr<i8, ub>) -> !pto.ptr<f32, ub>

    pto.mte_ub_gm %ub_y, %gm_y, %c20480_i64
      nburst(%c1_i64, %c20480_i64, %c20480_i64)
      : !pto.ptr<f32, ub>, !pto.ptr<f32, gm>, i64, i64, i64, i64
    "pto.set_flag_dyn"(%eid_index) {src_pipe = #pto.pipe<PIPE_MTE3>,
                                    dst_pipe = #pto.pipe<PIPE_V>} : (index) -> ()
  }

  "pto.wait_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                     dst_pipe = #pto.pipe<PIPE_V>,
                     event_id = #pto.event<EVENT_ID0>} : () -> ()
  "pto.wait_flag"() {src_pipe = #pto.pipe<PIPE_MTE3>,
                     dst_pipe = #pto.pipe<PIPE_V>,
                     event_id = #pto.event<EVENT_ID1>} : () -> ()
  return
}
}

这个版本没有循环内 if，只做 10 次 UB 写入，不会生成上述带谓词的 END 问题。

5. 规避方法

可以在 PTOAS 的 VPTO pipeline 中引入 MLIR 公共优化，在进入最终 VPTO/LLVM emission 之前消除这类小常量循环和常量条件分支。

规避思路：

1. 只针对 pto.simt_entry 函数；
2. 对小 trip count 的常量 scf.for 做 full unroll；
3. 运行 canonicalize；
4. 运行 sccp；
5. 再运行 canonicalize 和 cse。

对于如下模式：

for i in 0..16:
  if i < 10:
    store

展开后每个 i 都变成常量：

• i = 0..9 时，i < 10 折叠为 true；
• i = 10..15 时，i < 10 折叠为 false；
• canonicalize/sccp 删除恒 false 分支；
• 最终只剩 10 段直线 store。

本地验证结果：

• 最小复现的同步版本上板通过；
• 真实 RMSNorm 5120 VPTO 用例上板通过；
• 对应 SIMTVF 中原始 for 0..16 + if i < 10 已经被消除。

6. 需要编译器侧分析

PTOAS 可以通过循环展开和 if 消除规避该问题，但生成 END @!P0 的根因仍需要在编译器侧处理。

编译器仍需要分析：

1. SIMTVF 中存在循环内分支时，END @!P0 的生成路径；
2. END 被放入分支控制流相关路径的原因；
3. SIMTVF 控制流收敛分析对该场景的覆盖情况；
4. SIMTVF 的 END 应保持无谓词、无条件执行；

[KurrinQu]

RMSNorm d=5120 PTOAS 与 ASC 差异定位更新

1. ASC d=5120 代码形态：带不带 unroll 都通过

d=5120 的 ASC 生成代码。核心 SIMTVF 逻辑如下：

#include <tl_templates/ascend/common.h>

#include <limits>
__simt_vf__ __launch_bounds__(256) inline void simt_vf_2(__ubuf__ uint8_t* buf_dyn_shmem, int32_t t, float eps) {
  int32_t simtvf_tx = threadIdx.x;
  int32_t simtvf_ty = threadIdx.y;
  int32_t simtvf_tz = threadIdx.z;
  float x_frag[32];
  float sum_sq[1];
  #pragma unroll
  for (int32_t i = 0; i < 16; ++i) {
    *(float2*)(x_frag + (i * 2)) = *(__ubuf__ float2*)(((__ubuf__ float*)buf_dyn_shmem) + (((((t & 1) * 8192) + (i * 512)) + (simtvf_tx * 2)) + 5120));
  }
  sum_sq[0] = float(0x0p+0f/*0.000000e+00*/);
  #pragma unroll
  for (int32_t i_1 = 0; i_1 < 32; ++i_1) {
    if (i_1 < 20) {
      sum_sq[0] = (sum_sq[0] + (x_frag[i_1] * x_frag[i_1]));
    }
  }
  sum_sq[0] = tl::AscendAllReduce<tl::SumOp, 256, 1, 0>::run(sum_sq[0], (&(((__ubuf__ float*)buf_dyn_shmem)[37904])));
  float var = ((sum_sq[0] / float(0x1.4p+12f/*5.120000e+03*/)) + eps);
  float rstd_val = rsqrtf(var);
  ((__ubuf__ float*)buf_dyn_shmem)[(((t & 1) * 8) + 21504)] = rstd_val;
  #pragma unroll
  for (int32_t i_2 = 0; i_2 < 16; ++i_2) {
    if (i_2 < 10) {
      *(__ubuf__ float2*)(((__ubuf__ float*)buf_dyn_shmem) + (((((t & 1) * 8192) + (i_2 * 512)) + (simtvf_tx * 2)) + 21520)) = ((*(float2*)(x_frag + (i_2 * 2)) * make_float2(rstd_val, rstd_val)) * *(__ubuf__ float2*)(((__ubuf__ float*)buf_dyn_shmem) + ((i_2 * 512) + (simtvf_tx * 2))));
    }
  }
}

extern "C"
__global__ __vector__ void main_kernel(__gm__ float* RSTD, __gm__ float* W, __gm__ float* X, __gm__ float* Y, float eps) {
  int32_t bx = blockIdx.x;
  __ubuf__ uint8_t *buf_dyn_shmem = (__ubuf__ uint8_t *)0;
  copy_gm_to_ubuf_align_v2((__ubuf__ uint8_t*)((&(((__ubuf__ float*)buf_dyn_shmem)[0]))), (__gm__ uint8_t*)((&(W[0]))), 0, 1, 20480, 0, 0, 0, 0, 20480, 20480);
  AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(3);
  AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(3);
  AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(0);
  AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(0);
  AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(1);
  AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(1);
  for (int32_t t = 0; t < 64; ++t) {
    AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>((t & 1));
    copy_gm_to_ubuf_align_v2((__ubuf__ uint8_t*)((&(((__ubuf__ float*)buf_dyn_shmem)[(((t & 1) * 8192) + 5120)]))), (__gm__ uint8_t*)((&(X[((t * 327680) + (bx * 5120))]))), 0, 1, 20480, 0, 0, 0, 0, 20480, 20480);
    AscendC::SetFlag<AscendC::HardEvent::MTE2_V>((t & 1));
    AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>((t & 1));
    AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>((t & 1));
    asc_vf_call<simt_vf_2>(cce::dim3(256), buf_dyn_shmem, t, eps);
    AscendC::SetFlag<AscendC::HardEvent::V_MTE3>((t & 1));
    AscendC::SetFlag<AscendC::HardEvent::V_MTE2>((t & 1));
    AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>((t & 1));
    copy_ubuf_to_gm_align_v2((__gm__ void*)((&(Y[((t * 327680) + (bx * 5120))]))), (__ubuf__ void*)((&(((__ubuf__ float*)buf_dyn_shmem)[(((t & 1) * 8192) + 21520)]))), 0, 1, 20480, 4, 20480, 20480);
    copy_ubuf_to_gm_align_v2((__gm__ void*)((&(RSTD[((t * 64) + bx)]))), (__ubuf__ void*)((&(((__ubuf__ float*)buf_dyn_shmem)[(((t & 1) * 8) + 21504)]))), 0, 1, 4, 4, 4, 4);
    AscendC::SetFlag<AscendC::HardEvent::MTE3_V>((t & 1));
  }
  AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(0);
  AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(1);
  AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(0);
  AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(1);
}

这个 ASC 版本保留了 tail 分支：

for (int32_t i_2 = 0; i_2 < 16; ++i_2) {
  if (i_2 < 10) {
    store;
  }
}

当前验证结果：

ASC d=5120 非 unroll 版本：通过
ASC d=5120 加 pragma unroll 版本：通过

这说明 for i in 0..16 { if (i < 10) store } 这类 tail 分支语义本身可以被 ASC/BiSheng 正确处理。

2. ASC unroll 版本与 PTOAS 手工展开版本等价

ASC 加 pragma unroll 后，本质上等价于在 PTOAS 侧把 tail 循环：

for i in 0..16:
  if i < 10:
    store

展开成：

store i=0
store i=1
store i=2
store i=3
store i=4
store i=5
store i=6
store i=7
store i=8
store i=9

PTOAS 侧将该 tail 循环展开成 10 次无条件 vector store 后，d=5120 上板通过。

3. ASC 非 unroll 版本通过，但 PTOAS 对应逻辑失败

失败的 PTO 源文件是：

https://github.com/KurrinQu/PTOAS/blob/9388e8d36eb6af81cf573951deb9f50ca24a482b/test/vpto/cases/kernels/rmsnorm-fp32-5120-tilelang/kernel.pto

PTOAS 中 SIMTVF 内也有对应的 tail 分支逻辑：

func.func @simt_vf_0(%arg0: !pto.ptr<i8, ub>, %arg1: f32, %arg2: i32)
    attributes {pto.simt_entry, pto.simt_max_regs = 64 : i32} {
  ...

  scf.for %i_268 = %cast_263 to %ub_267 step %c1_266 {
    %cast_269 = arith.index_cast %i_268 : index to i32
    %c10_i32_270 = arith.constant 10 : i32
    %cmp_271 = arith.cmpi slt, %cast_269, %c10_i32_270 : i32

    scf.if %cmp_271 {
      ...
      "llvm.store"(%mul_297, %gep_ub_ptr_316)
        : (vector<2xf32>, !llvm.ptr<6>) -> ()
    }
  }

  return
}

因此，问题更可能出在 ASC 非 unroll 路径和 PTOAS 路径在后续 LLVM IR 或最终 SIMTVF 汇编生成上的差异。

下一步建议重点对比：

1. ASC 非 unroll 版本进入后端前的 LLVM IR
2. PTOAS 传入至bisheng的LLVM IR
3. 两者最终生成的 SIMTVF 汇编
4. tail if (i < 10) 对应的分支结构
5. END 指令布局及其是否落在分支可达路径中

[katakuri123]

1. ASC (EmitC) 非 unroll 版本进入后端前的 LLVM IR

EmitC 路径不经过 SCFToCF。scf.for/scf.if 被 SCFToEmitC 转 C++，
Bisheng 以 -O0 编译。等价的 LLVM IR（alloca-based）：

entry:
  %i = alloca i32
  store i32 0, i32* %i
  br label %for.cond

for.cond:                    ; load → cmp → cond_br
  %i.val = load i32, i32* %i
  %loop_cond = icmp slt i32 %i.val, 16
  br i1 %loop_cond, label %for.body, label %for.end

for.body:
  %active = icmp slt i32 %i.val, 10
  br i1 %active, label %if.then, label %if.end   ; inner cond_br

if.then:  store volatile float %val, ptr %ptr
          br label %if.end
if.end:   br label %for.inc

for.inc:                     ; latch: add → store → br
  %next = add nsw i32 %i.val, 1
  store i32 %next, i32* %i
  br label %for.cond

for.end:
  call void @...sync()
  ret void                   ; ← cond_br 可达

循环变量在栈上（alloca），跨 block 数据通过 load/store 传递。

---

2. PTOAS 传入 Bisheng 的 LLVM IR

define linkonce_odr simt_entry void @simt_vf_0(ptr addrspace(6) %0, float %1, i32 %2) {
  %4 = call i32 @llvm.hivm.get.TID.X()
  ...
  br label %7

7:                         ; preds = %24, %3
  %8 = phi i64 [ %25, %24 ], [ 0, %3 ]
  %9 = icmp slt i64 %8, 16
  br i1 %9, label %10, label %26              ; ★ 循环出口 cond_br

10:                        ; preds = %7
  %11 = trunc i64 %8 to i32
  %12 = icmp slt i32 %11, 10
  br i1 %12, label %13, label %24              ; ★ if 内层 cond_br

13:                        ; preds = %10        store ...
  store <2 x float> ..., ptr addrspace(6) %23
  br label %24

24:                        ; preds = %13, %10   ← join point
  %25 = add i64 %8, 1
  br label %7                                   ; 循环回边

26:                        ; preds = %7          ← 仅 cond_br 可达
  call void @llvm.hivm.sync.workitems()
  ret void                                       ; ★ END 落在这里
}

循环变量通过 phi 在 block 间直接传递（SSA form），无 alloca。

---

3. 两者最终生成的 SIMTVF 汇编

3a. 2×2 对比

核心模式来自 ASC rmsnorm kernel (simt_vf_2) 的循环：

#pragma unroll                         // ← ASC 侧：有/无 #pragma unroll
for (int32_t i = 0; i < 16; ++i) {
    if (i < 10) {                      // scf.if 等价物
        *(__ubuf__ float2*)(...) = ...;  // store
    }
}

维度	ASC (EmitC)	PTOAS (VPTO)
IR 形式	alloca + load/store	phi node (SSA)
bisheng 优化	-O0 (C++ 编译默认)	-O2 (ptoas 硬编码)
unroll 控制	#pragma unroll 有/无	PTOAS MLIR pass 有/无

	ASC 不加 unroll	ASC 加 unroll	VPTO 不加 unroll	VPTO 加 unroll
优化级别	-O0	-O0	-O2	-O2
BRANCH 数	3	0	7	0
循环结构	标量迭代	直线	部分展开 4×	直线
IFU pred_loc	—	—	pred_loc:2 (45 条)	—
触发 bug	❌	❌	✅	❌

3b. 关键对比：同样有分支，ASC 不触发，VPTO 触发

ASC 不加 unroll @ -O0 — 3 条 BRANCH，END 干净：

simt_vf_0_simt_entry:
    MOVI   R6, #0                     ; i = 0
.LBB0_1:                               ; ← loop header
    ISETP.GT P1, PT, R6, R8, PT       ; i > 15 → exit
    BRANCH .LBB0_5, P1                ; ★ cond_br → END block
    ISETP.GT P0, PT, R6, #9, PT       ; i > 9 ? (if)
    BRANCH .LBB0_4, P0                ; ★ 内层 cond_br
    STS.b64  R8, [R5]                 ; store (conditional)
.LBB0_4:
    IADD.X  R6, R6, #1                ; i++
.LBB0_5:
    BRANCH  .LBB0_1, !P0              ; loop back-edge
    BAR.THREAD_BLOCK
    END                                ; ← clean (bisheng -O0 不做 predicate 传播)

VPTO 不加 unroll @ -O2 — 7 条 BRANCH，IFU trace 显示谓词传播：

simt_vf_0_simt_entry:
    MOVI   R4, #16
    BRANCH .LBB0_1
.LBB0_9:                               ; ← increment by 4
    ISETP.NE P0, PT, R8, R4, PT        ; i != 16 ?
    BRANCH .LBB0_10, !P0               ; ★ cond_br → exit (END)
.LBB0_1:
    ISETP.GT P0, PT, R8, #9, PT        ; i > 9 ?
    BRANCH .LBB0_3, P0                 ; ★ 内层 cond_br
    STS.b64                              ; store iter 0
.LBB0_3: ... (iter 1 check + store)
.LBB0_5: ... (iter 2 check + store)
.LBB0_7: ... (iter 3 check + store)
    BRANCH .LBB0_9                     ; ★ 7 条 BRANCH 总计
.LBB0_10:
    BAR.THREAD_BLOCK
    END                                 ; ← 静态无谓词，但 IFU trace 有 pred_loc

bisheng -O2 自动 unroll 4×，同时做 predicate 传播。

两个加 unroll 的版本都是 10 个连续 STS.b64 + BAR + END，0 条 BRANCH，不再展开。

3c. IFU 模拟器 trace（仅 VPTO 不加 unroll）

@2049 update_pred_info, start_pc:0x10d0da60, pred_loc:2, end_pos:2
@2049 save: MOV_SRC_TO_DST_ALIGNv2  pc=0x10d0da60
@2049 save: JUMPC                    pc=0x10d0da64    ; SIMT body 返回

pred_loc:2（bit 1 置位）= 返回跳转被谓词化。共 45 条谓词化指令，其他三个版本无此现象。

---

4. 分支结构与 END 布局

SCFToCF 产出的 CFG（用 §2 的 block 号标注）：

                [%3: entry]
                    |
                    v
               [%7: header]
               /           \
     cond_br (i<16)     cond_br (!(i<16))
     true →              false →
          /                    \
         v                      v
   [%10: if_body]          [%26: exit]
   /            \              |
cond_br(i<10) cond_br(!(i<10))|
true →         false →         |
   v               v           |
[%13:store]  [%24:latch]      |
   |               |           |
   └── br ────────┘           |
          |                    |
          └── br(i+1) ──→ [%7] ──┘ (i>=16)

%26（exit，含 ret void）的唯一前驱是 %7 的 false 分支（cond_br），
到达 ret 的路径上存在两层条件分支。

逐迭代路径：

i=0..9:   %7(true) → %10(true) → %13(store) → %24 → %7
i=10..15: %7(true) → %10(false) → %24 → %7
i=16:     %7(false) → %26 → sync → ret void (★ 仅 cond_br 可达)

---

[KurrinQu]

1. ASC (EmitC) 非 unroll 版本进入后端前的 LLVM IR

EmitC 路径不经过 SCFToCF。scf.for/scf.if 被 SCFToEmitC 转 C++， Bisheng 以 -O0 编译。等价的 LLVM IR（alloca-based）：

entry:
%i = alloca i32
store i32 0, i32* %i
br label %for.cond

for.cond: ; load → cmp → cond_br
%i.val = load i32, i32* %i
%loop_cond = icmp slt i32 %i.val, 16
br i1 %loop_cond, label %for.body, label %for.end

for.body:
%active = icmp slt i32 %i.val, 10
br i1 %active, label %if.then, label %if.end ; inner cond_br

if.then: store volatile float %val, ptr %ptr
br label %if.end
if.end: br label %for.inc

for.inc: ; latch: add → store → br
%next = add nsw i32 %i.val, 1
store i32 %next, i32* %i
br label %for.cond

for.end:
call void @...sync()
ret void ; ← cond_br 可达
循环变量在栈上（alloca），跨 block 数据通过 load/store 传递。

2. PTOAS 传入 Bisheng 的 LLVM IR

define linkonce_odr simt_entry void @simt_vf_0(ptr addrspace(6) %0, float %1, i32 %2) {
%4 = call i32 @llvm.hivm.get.TID.X()
...
br label %7

7: ; preds = %24, %3
%8 = phi i64 [ %25, %24 ], [ 0, %3 ]
%9 = icmp slt i64 %8, 16
br i1 %9, label %10, label %26 ; ★ 循环出口 cond_br

10: ; preds = %7
%11 = trunc i64 %8 to i32
%12 = icmp slt i32 %11, 10
br i1 %12, label %13, label %24 ; ★ if 内层 cond_br

13: ; preds = %10 store ...
store <2 x float> ..., ptr addrspace(6) %23
br label %24

24: ; preds = %13, %10 ← join point
%25 = add i64 %8, 1
br label %7 ; 循环回边

26: ; preds = %7 ← 仅 cond_br 可达
call void @llvm.hivm.sync.workitems()
ret void ; ★ END 落在这里
}
循环变量通过 phi 在 block 间直接传递（SSA form），无 alloca。

3. 两者最终生成的 SIMTVF 汇编

3a. 2×2 对比

核心模式来自 ASC rmsnorm kernel (simt_vf_2) 的循环：

#pragma unroll // ← ASC 侧：有/无 #pragma unroll
for (int32_t i = 0; i < 16; ++i) {
if (i < 10) { // scf.if 等价物
(ubuf float2)(...) = ...; // store
}
}
维度 ASC (EmitC) PTOAS (VPTO)
IR 形式 alloca + load/store phi node (SSA)
bisheng 优化 -O0 (C++ 编译默认) -O2 (ptoas 硬编码)
unroll 控制 #pragma unroll 有/无 PTOAS MLIR pass 有/无
ASC 不加 unroll ASC 加 unroll VPTO 不加 unroll VPTO 加 unroll
优化级别 -O0 -O0 -O2 -O2
BRANCH 数 3 0 7 0
循环结构 标量迭代 直线 部分展开 4× 直线
IFU pred_loc — — pred_loc:2 (45 条) —
触发 bug ❌ ❌ ✅ ❌

3b. 关键对比：同样有分支，ASC 不触发，VPTO 触发

ASC 不加 unroll @ -O0 — 3 条 BRANCH，END 干净：

simt_vf_0_simt_entry:
MOVI R6, #0 ; i = 0
.LBB0_1: ; ← loop header
ISETP.GT P1, PT, R6, R8, PT ; i > 15 → exit
BRANCH .LBB0_5, P1 ; ★ cond_br → END block
ISETP.GT P0, PT, R6, #9, PT ; i > 9 ? (if)
BRANCH .LBB0_4, P0 ; ★ 内层 cond_br
STS.b64 R8, [R5] ; store (conditional)
.LBB0_4:
IADD.X R6, R6, #1 ; i++
.LBB0_5:
BRANCH .LBB0_1, !P0 ; loop back-edge
BAR.THREAD_BLOCK
END ; ← clean (bisheng -O0 不做 predicate 传播)
VPTO 不加 unroll @ -O2 — 7 条 BRANCH，IFU trace 显示谓词传播：

simt_vf_0_simt_entry:
MOVI R4, #16
BRANCH .LBB0_1
.LBB0_9: ; ← increment by 4
ISETP.NE P0, PT, R8, R4, PT ; i != 16 ?
BRANCH .LBB0_10, !P0 ; ★ cond_br → exit (END)
.LBB0_1:
ISETP.GT P0, PT, R8, #9, PT ; i > 9 ?
BRANCH .LBB0_3, P0 ; ★ 内层 cond_br
STS.b64 ; store iter 0
.LBB0_3: ... (iter 1 check + store)
.LBB0_5: ... (iter 2 check + store)
.LBB0_7: ... (iter 3 check + store)
BRANCH .LBB0_9 ; ★ 7 条 BRANCH 总计
.LBB0_10:
BAR.THREAD_BLOCK
END ; ← 静态无谓词，但 IFU trace 有 pred_loc

bisheng -O2 自动 unroll 4×，同时做 predicate 传播。

两个加 unroll 的版本都是 10 个连续 STS.b64 + BAR + END，0 条 BRANCH，不再展开。

3c. IFU 模拟器 trace（仅 VPTO 不加 unroll）

@2049 update_pred_info, start_pc:0x10d0da60, pred_loc:2, end_pos:2
@2049 save: MOV_SRC_TO_DST_ALIGNv2  pc=0x10d0da60
@2049 save: JUMPC                    pc=0x10d0da64    ; SIMT body 返回

pred_loc:2（bit 1 置位）= 返回跳转被谓词化。共 45 条谓词化指令，其他三个版本无此现象。

4. 分支结构与 END 布局

SCFToCF 产出的 CFG（用 §2 的 block 号标注）：

                [%3: entry]
                    |
                    v
               [%7: header]
               /           \
     cond_br (i<16)     cond_br (!(i<16))
     true →              false →
          /                    \
         v                      v
   [%10: if_body]          [%26: exit]
   /            \              |
cond_br(i<10) cond_br(!(i<10))|
true →         false →         |
   v               v           |
[%13:store]  [%24:latch]      |
   |               |           |
   └── br ────────┘           |
          |                    |
          └── br(i+1) ──→ [%7] ──┘ (i>=16)

%26（exit，含 ret void）的唯一前驱是 %7 的 false 分支（cond_br）， 到达 ret 的路径上存在两层条件分支。

逐迭代路径：

i=0..9:   %7(true) → %10(true) → %13(store) → %24 → %7
i=10..15: %7(true) → %10(false) → %24 → %7
i=16:     %7(false) → %26 → sync → ret void (★ 仅 cond_br 可达)

有2点疑问：

1. ASC的代码和EmitC有什么联系？我认为这两者没有关系，ASC是指这里的代码，是由TileLang直接产生；EmitC指的是通过PTOAS产生C++ pto-isa指令调用的代码。
2. 这里分析-O0和-O2选项的原因是因为ASC使用-O0或者-O2选项的结果有差异吗？如果没有差异，建议都对齐使用-O2选项对比。

[KurrinQu]

根因定位

该问题的根因是：PTOAS 生成的 LLVM IR 进入 Bisheng 编译时，使用的 driver action 路径与 CCE/ASC 源码编译路径不同，导致 Bisheng 没有为 AICore CCE 编译自动添加 -disable-ifcvt- simple=true / -disable-ifcvt-simple-false=true 等后端选项。

最终表现为：

• release 版本 Bisheng 在 if-convert 相关 pass 中走到错误路径，最终生成非法 SIMTVF 代码；
• debug 版本 Bisheng 会在对应路径触发断言。

具体原因

问题出在 Bisheng driver 中 CCEActionBuilder::addDeviceDepences 的判断逻辑。

CCE/ASC 源码路径

对于 CCE/ASC 源码输入，例如 .cpp：

• types::isCCE(InputType) 返回 true
• CCEActionBuilder 处于 active 状态
• driver 会创建带 Action::OFK_CCE 的 OffloadAction
• 后续 Clang::ConstructJob 中：

IsCCE = JA.isOffloading(Action::OFK_CCE)

返回 true

因此会进入 IsCCE 分支，自动添加 CCE/AICore 编译相关选项，包括：

-mllvm -disable-ifcvt-simple=true
-mllvm -disable-ifcvt-simple-false=true

LLVM IR 路径

对于 PTOAS 生成的 LLVM IR，例如 .ll：

• types::isCCE(InputType) 返回 false
• LLVM IR 本身不会被识别为 CCE 源码
• isCCEAiCoreIR 只有在传入以下选项时才会返回 true：

--cce-aicore-jit

或

--cce-bitcode-is-aicore

如果 PTOAS 编译 .ll 时没有传这些标志：

• isCCEAiCoreIR 返回 false
• CCEActionBuilder 变为 inactive
• 不会创建 Action::OFK_CCE 的 OffloadAction
• Clang::ConstructJob 中 IsCCE 为 false
• 不会进入 CCE 分支
• 也就不会添加 if-convert 相关的禁用选项

相关代码注释

Bisheng driver 中已有注释指出了这个问题：

static bool isCCEAiCoreIR(const ArgList &Args, types::ID Ty) {
  // NOTE: We really should be checking for --cce-aicore-only, however there are
  // many legacy test cases in CI that compiles host-only bitcode with
  // aicore-only. We need to resolve those at some point.
}

理论上 --cce-aicore-only 也应该参与判断，但由于历史 CI 用例中存在 host-only bitcode 携带 --cce-aicore-only 的情况，目前没有直接使用该选项判断输入是否为 AICore IR。

影响

PTOAS 当前编译 LLVM IR 时虽然传了：

--cce-aicore-only

但该选项不足以让 CCEActionBuilder 进入 active 状态。

因此 PTOAS IR 编译路径和 CCE/ASC 源码路径实际使用的 Bisheng 后端选项不一致，最终导致 if-convert 相关 pass 行为不同，生成了错误代码。

解决方案

1. PTOAS 编译 .ll 时显式传入 AICore IR 标志，使 Bisheng driver 走 CCE offload/action 路径：

--cce-bitcode-is-aicore

2. 删除PTOAS--target=hiipu64-hisilicon-cce编译选项
   添加--cce-bitcode-is-aicore选项后，可以使得device编译真正进入到CCEActionBuilder ，其中会判断Host Target Triple不能为hiipu64-hisilicon-cce，而当前PTOAS添加了--target=hiipu64-hisilicon-cce选项，会触发编译如下报错，需要删除：

 bisheng: error: unsupported use of hiipu64-hisilicon-cce for host compilation.