# IR接入简介

## 多级IR抽象架构

- 提供一系列高层抽象接口，屏蔽底层细节，将硬件无关表达映射到底层指令，提升算子开发易用性。
- 提供细粒度性能控制接口，能够精准控制片上内存地址、流水同步插入位置以及是否使能乒乓流水优化等。
- 基于多级接口支持自定义DSL/生态框架灵活对接，实现自定义算子在昇腾NPU上的高性能运行。

```text
  Torch-MLIR / Triton      （框架/DSL 层）
          |
          v
  Linalg / Tensor           （通用张量代数层）
          |
          v
  HFusion                   （硬件感知融合调度层）
          |
          v
  HIVM                      （NPU 指令层）
          |
          v
  LIR -> Binary             （机器码生成）
```

- **Linalg / Tensor层**：使用标准MLIR dialect表达算子语义，支持`Elemwise`、`Broadcast`、`Reduce`、`Transpose`、`Concat`等运算，HFusion自动完成算子融合、切分和调度。
- **HFusion层**：提供昇腾NPU感知的Named Op（如`hfusion.elemwise_unary`、`hfusion.cast`、`hfusion.select`、`hfusion.reduce_with_index`等），支持`tensor`语义，自动完成bufferization、tiling和调度。
- **HIVM层**：直接映射NPU硬件指令，显式控制存储层级（GM/UB/L1/L0）、计算流水线（Vector/Cube/MTE）和同步原语，支持精细粒度的性能调优。

上述多级接口支持自定义DSL和生态框架灵活对接。Triton和PyTorch等框架通过IR转换接入上述流程，实现自定义算子在昇腾NPU上的高性能运行。

## 编译选项与函数属性

### 编译选项

`bishengir-compile`提供以下公共编译选项：

| 选项 | 默认值 | 说明 |
|------|--------|------|
| `-target` | `Ascend<Name>` | 目标设备，用于获取核数、片上内存大小等硬件规格，可通过查询`npu-smi info` |
| `-block-dim` | 1 | 指定使用的block数量，编译后kernel携带`hacc.block_dim`属性 |
| `-enable-hfusion-compile` | `false` | 使能HFusion编译流程（融合、调度、切分） |
| `-enable-hivm-compile` | `true` | 使能HIVM编译流程（转换到HIVM指令并优化） |
| `-enable-torch-compile` | `false` | 使能Torch-MLIR编译流程 |
| `-enable-triton-kernel-compile` | `false` | 使能Triton kernel编译流程 |

支持的目标设备包括：

- Ascend 950PR/Ascend 950DT
- Atlas A3训练系列产品/Atlas A3推理系列产品
- Atlas A2训练系列产品/Atlas A2推理系列产品

### 函数属性

以下属性用于标注kernel入口函数，各接入路径通用：

| 属性 | 说明 |
|------|------|
| `hacc.entry` | 标记当前函数为kernel入口 |
| `hacc.function_kind = #hacc.function_kind<DEVICE>` | 表示函数运行在DEVICE设备侧 |
| `hacc.function_kind = #hacc.function_kind<HOST>` | 表示函数运行在HOST侧，HFusion会自动outline出设备kernel |

示例：

```mlir
func.func @kernel(...) attributes {hacc.entry, hacc.function_kind = #hacc.function_kind<DEVICE>} {
  ...
}
```

## Triton接入

Triton是主流的高性能算子开发编程语言。[Triton Ascend](https://gitcode.com/Ascend/triton-ascend/)将Triton算子转换为MLIR，接入AscendNPU IR生态，支持在昇腾NPU上运行Triton内核。详细的Triton接入指南（含安装、环境、算子映射及昇腾扩展）请参考[Triton接入](triton_interface.md)。

## TileLang接入

TileLang（tilelang-ascend）是面向昇腾NPU的领域特定语言，基于tile-lang的Python语法和[TVM](https://tvm.apache.org/)构建，支持GEMM、向量运算和注意力机制等算子，可将算子编译为AscendNPU IR（HIVM）在昇腾NPU上运行。详细的TileLang接入说明请参考[TileLang接入](tile_lang_interface.md)。

## 框架接入

AscendNPU IR支持框架（PyTorch/TensorFlow/MindSpore）接入，有两种方式：

- **DSL接入方式**：如Triton、TileLang
- **IR接入方式**：如Torch IR、Linalg/HFusion IR、HIVM IR

详细的框架接入说明请参考[框架接入](framework_interface.md)。