[Bug] bisheng (CANN 9.0.0-beta.1) 无法选择 intrinsic `%llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV`

### Component

PTO Dialect / ODS (include/PTO/IR)

### Description


## 问题描述

在编译包含 `pto.tinsert` (Acc→Vec ND 路径) 的 TileLang ST 用例时，bisheng 设备编译器在 instruction selection 阶段崩溃，报告无法选择 ptoas 生成的硬件 intrinsic。

### 复现命令

```bash
source tmp/cj_mk_env.sh && ptoas_build
python test/tilelang_st/script/run_st.py -r sim -v a5 -t tinsert
```

### 错误日志

```
fatal error: error in backend: Cannot select: intrinsic %llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV

PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/
and include the crash backtrace, preprocessed source, and associated run script.

Stack dump:
0. Program arguments: /usr/local/CANN/cann-9.0.0-beta.1/bin/bisheng 
   --target=hiipu64-hisilicon-cce -march=dav-c310-cube 
   --cce-aicore-arch=dav-c310-cube --cce-aicore-only -O2 -dc -mllvm 
   -cce-dyn-kernel-stack-size=true -c -x ir - -o <output>.o

1. Code generation
2. Running pass 'Function Pass Manager' on module '<stdin>'
3. Running pass 'HiIPU Non VF DAG->DAG Pattern Instruction Selection' 
   on function '@TINSERT_acc2vec_nd_f16_16x16_mix_aic'

Stack dump without symbol names:
/usr/local/CANN/cann-9.0.0-beta.1/bin/bisheng(_ZN4llvm16SelectionDAGISel15CannotYetSelect...)
```

---

## 根因分析

PTOAS 编译器与 CANN 设备编译器之间存在 **intrinsic 版本不兼容**：

| 组件 | 版本 | 状态 |
|------|------|------|
| ptoas (LLVM) | 19.1.7 (v0.46) | 生成 V2 intrinsic |
| bisheng (CANN) | clang 15.0.5 (cann-9.0.0-beta.1) | **不支持 V2 intrinsic** |

**问题本质**：
- ptoas 将 `pto.mte_l0c_ub` (Acc→UB 数据搬运) 下降为 LLVM IR 时，生成了硬件 intrinsic `@llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV`
- bisheng 后端（基于 LLVM 15.0.5）在 instruction selection 阶段无法识别该 intrinsic
- intrinsic 名称中的 `V2` 后缀表明这是较新版本的变体

---

## 触发条件

**触发路径**：必须使用 Acc→UB (L0C→UB) 数据搬运操作

**会触发的 pto op**：
- `pto.mte_l0c_ub` - 将 Acc 缓冲区数据搬到 UB 缓冲区
- 典型场景：`pto.tinsert` 中 Acc→Vec ND/DN/NZ 路径

**不会触发的操作**：
- `pto.mte_l0c_gm` - Acc 直接写回 GM（跳过 UB）
- `pto.mte_l1_ub` - L1→UB 搬运（使用不同 intrinsic）
- `pto.mte_ub_gm` - UB→GM 搬运
- `pto.textract` 的所有路径（验证回路走 `mte_l0c_gm` 或纯向量通路）

---

## 环境信息

**软件环境**：
- CANN 版本：cann-9.0.0-beta.1
- bisheng 版本：clang-5c68a1cb1231 (clang 15.0.5)
- ptoas 版本：0.46 (LLVM 19.1.7)
- OS：Linux Ubuntu (从栈帧可见 libc.so.6)

**硬件目标**：
- 架构：Hisilicon CCE (dav-c310-cube)
- 编译选项：`--cce-aicore-arch=dav-c310-cube`

---

## 影响范围

### 受影响的算子
- ✅ **tinsert**：Acc→Vec ND 路径 (`template_tinsert_acc_to_vec_nd`)
- ✅ 任何使用 `pto.mte_l0c_ub` 的自定义 TileOp

### 不受影响的算子

| 算子 | 路径 | 原因 |
|------|------|------|
| textract | Mat→Left/Right | 使用 `mte_l0c_gm` (Acc→GM) |
| textract_fp | Acc→Mat | 使用 `mte_l0c_gm` (Acc→GM) |
| textract_v2v | Vec→Vec | 纯向量通路 (`tload`→`textract`→`tstore`) |
| tmatmul | Mat×Mat | 不使用 UB 搬运输出 |
| tload/tstore | GM↔Tile | 不涉及 Acc→UB |

**验证**：textract ST 测试套件已验证可正常编译，说明问题仅限于 Acc→UB 路径。

---

## 代码位置

**触发点**：
- `lib/TileOps/tinsert_template.py:418-422` - `template_tinsert_acc_to_vec_nd` 调用 `pto.mte_l0c_ub`
- `test/tilelang_st/npu/a5/src/st/testcase/tinsert/tinsert.pto` - Case 3 & Case 6 (Acc→Vec ND kernel)

**PTO IR 示例** (会触发)：
```mlir
pto.tinsert ins(%acc_tile, %c0_idx, %c0_idx : 
                  !pto.tile_buf<acc, 16x16xf32, blayout=col_major, slayout=row_major, fractal=1024>, 
                  index, index)
            outs(%dst_vec_tile : !pto.tile_buf<vec, 16x16xf16, blayout=row_major, slayout=none_box, fractal=512>)
```

---

## 建议的修复方向

### 方案 1：升级 CANN 工具链（推荐）
升级 bisheng 至支持 `MOV.UB.TO.OUT.ALIGN.V2.DV` intrinsic 的版本。

**验证方法**：检查新版 bisheng 的 release notes 或 changelog，确认是否包含该 intrinsic 支持。

### 方案 2：PTOAS 降级 intrinsic 版本
修改 ptoas 的 lowering 逻辑，生成 V1 版本的 intrinsic（如 `MOV.UB.TO.OUT.ALIGN.DV`，无 V2 后缀）。

**风险**：V1 intrinsic 可能不支持某些特性（如对齐优化），需要验证功能完整性。

### 方案 3：规避受影响路径（临时）
从 `tinsert.pto` 中注释掉使用 `mte_l0c_ub` 的 kernel，仅测试其他路径：

```bash
# 临时禁用 Case 3 和 Case 6 (Acc→Vec ND)
# 保留 Acc→Mat 和 Vec→Vec 路径验证
```

---

## 相关测试

**可通过的测试**（不涉及 Acc→UB）：
```bash
python test/tilelang_st/script/run_st.py -r sim -v a5 -t textract
python test/tilelang_st/script/run_st.py -r sim -v a5 -t textract_fp
python test/tilelang_st/script/run_st.py -r sim -v a5 -t textract_v2v
```

**失败的测试**（涉及 Acc→UB）：
```bash
python test/tilelang_st/script/run_st.py -r sim -v a5 -t tinsert
```


### Reproduction (minimal)

```markdown
见PR：https://github.com/hw-native-sys/PTOAS/pull/765
```

### Expected behavior

without error

### Actual behavior / error logs

```text
fatal error: error in backend: Cannot select: intrinsic %llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV

PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/
and include the crash backtrace, preprocessed source, and associated run script.

Stack dump:
0. Program arguments: /usr/local/CANN/cann-9.0.0-beta.1/bin/bisheng 
   --target=hiipu64-hisilicon-cce -march=dav-c310-cube 
   --cce-aicore-arch=dav-c310-cube --cce-aicore-only -O2 -dc -mllvm 
   -cce-dyn-kernel-stack-size=true -c -x ir - -o <output>.o

1. Code generation
2. Running pass 'Function Pass Manager' on module '<stdin>'
3. Running pass 'HiIPU Non VF DAG->DAG Pattern Instruction Selection' 
   on function '@TINSERT_acc2vec_nd_f16_16x16_mix_aic'

Stack dump without symbol names:
/usr/local/CANN/cann-9.0.0-beta.1/bin/bisheng(_ZN4llvm16SelectionDAGISel15CannotYetSelect...)
```

### Git commit

5efd2764452789a4d452fd6c95faf894dc313172

### Host platform

None

### Target Ascend arch (if relevant)

None

### PTOAS build level (if relevant)

None

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Bug] bisheng (CANN 9.0.0-beta.1) 无法选择 intrinsic `%llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV` #811

Component

Description

问题描述

复现命令

错误日志

根因分析

触发条件

环境信息

影响范围

受影响的算子

不受影响的算子

代码位置

建议的修复方向

方案 1：升级 CANN 工具链（推荐）

方案 2：PTOAS 降级 intrinsic 版本

方案 3：规避受影响路径（临时）

相关测试

Reproduction (minimal)

Expected behavior

Actual behavior / error logs

Git commit

Host platform

Target Ascend arch (if relevant)

PTOAS build level (if relevant)

Metadata

Assignees

Labels

Type

Fields

Projects

Milestone

Relationships

Development

组件	版本	状态
ptoas (LLVM)	19.1.7 (v0.46)	生成 V2 intrinsic
bisheng (CANN)	clang 15.0.5 (cann-9.0.0-beta.1)	不支持 V2 intrinsic

算子	路径	原因
textract	Mat→Left/Right	使用 `mte_l0c_gm` (Acc→GM)
textract_fp	Acc→Mat	使用 `mte_l0c_gm` (Acc→GM)
textract_v2v	Vec→Vec	纯向量通路 (`tload`→`textract`→`tstore`)
tmatmul	Mat×Mat	不使用 UB 搬运输出
tload/tstore	GM↔Tile	不涉及 Acc→UB

[Bug] bisheng (CANN 9.0.0-beta.1) 无法选择 intrinsic %llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV #811

Description

Component

Description

问题描述

复现命令

错误日志

根因分析

触发条件

环境信息

影响范围

受影响的算子

不受影响的算子

代码位置

建议的修复方向

方案 1：升级 CANN 工具链（推荐）

方案 2：PTOAS 降级 intrinsic 版本

方案 3：规避受影响路径（临时）

相关测试

Reproduction (minimal)

Expected behavior

Actual behavior / error logs

Git commit

Host platform

Target Ascend arch (if relevant)

PTOAS build level (if relevant)

Metadata

Metadata

Assignees

Labels

Type

Fields

Projects

Milestone

Relationships

Development

Issue actions

[Bug] bisheng (CANN 9.0.0-beta.1) 无法选择 intrinsic `%llvm.hivm.MOV.UB.TO.OUT.ALIGN.V2.DV` #811