English | 中文版

2. Hello World：第一个 NPU 程序

安装

ascend-rs 以自包含分发包的形式提供，包含预构建的编译器后端和用于宿主机与内核 API 的 Rust 源码 crate。

前置条件：

• 目标机器上已安装 CANN 工具包（8.x 或 9.x）
• Rust nightly 工具链（由分发包中的 rust-toolchain.toml 自动安装）

安装步骤：

# 1. 解压分发包
tar xzf ascend-rs-0.1.1-$(uname -m).tar.gz
cd ascend-rs-0.1.1

# 2. 加载 CANN 环境
source /usr/local/Ascend/ascend-toolkit/latest/bin/setenv.bash

# 3. 使编译器后端可被发现
export LD_LIBRARY_PATH="$(pwd)/lib:$LD_LIBRARY_PATH"

# 4. 验证（编译内核并在 NPU 上运行）
bash test.sh --run

分发包内容：

ascend-rs-0.1.1/
├── lib/librustc_codegen_mlir.so   # 编译器后端（Rust → NPU 二进制）
├── crates/
│   ├── ascend_rs/                 # 宿主机 API：设备、流、内存、内核启动
│   ├── ascend_sys/                # FFI 绑定（从 CANN 头文件自动生成）
│   ├── ascend_std/                # 内核运行时：缓冲区操作、向量指令
│   ├── ascend_std_macros/         # #[aiv_kernel] 属性宏
│   ├── ascend_rs_builder/         # 构建时内核编译器（KernelBuilder）
│   └── ascend_rs_builder_config/  # CANN 路径检测
├── examples/vec_add/              # 可运行的入门项目
├── test.sh                        # 冒烟测试
└── rust-toolchain.toml            # 固定的 nightly 版本

编译器后端（librustc_codegen_mlir.so）在内核编译时由 rustc 加载。它将 Rust 内核代码经由 MLIR 转换为 AscendC C++，然后调用 CANN 的 bisheng 编译器生成 NPU 二进制文件。用户通过 build.rs 脚本中的 KernelBuilder 间接使用它。

让我们从最简单的例子开始。这个 Hello World 示例展示了 ascend-rs 宿主机 API 的基本用法——用 Rust 安全地初始化 NPU、创建执行上下文、启动内核。

内核代码（C++）

在当前阶段，Hello World 使用 C++ 内核，这是 CANN SDK 的原生方式：

// hello_world.cpp
#include "kernel_operator.h"

extern "C" __global__ __aicore__ void hello_world() {
    AscendC::printf("Hello World!!!\n");
}

extern "C" void hello_world_do(uint32_t blockDim, void *stream) {
    hello_world<<<blockDim, nullptr, stream>>>();
}

这里的 __global__ 标记函数为可从宿主机调用的入口点，__aicore__ 表明它运行在昇腾的 AI Core 上。<<<...>>> 语法与 CUDA 类似，指定了并行度和执行流。

宿主机代码（Rust）

宿主机代码展示了 ascend-rs 最重要的设计理念——RAII 资源管理和生命周期安全：

use ascend_rs::prelude::*;
use std::error::Error;

// 声明 C++ 内核的 FFI 接口
unsafe extern "C" {
    fn hello_world_do(dim: u32, stream: *mut std::ffi::c_void);
}

fn main() -> Result<(), Box<dyn Error>> {
    // 步骤 1: 初始化 ACL 运行时
    let acl = Acl::new()?;

    // 步骤 2: 选择并初始化设备
    let device = Device::new(&acl)?;

    // 步骤 3: 创建执行上下文和流
    let context = AclContext::new(&device)?;
    let stream = AclStream::new(&context)?;

    // 步骤 4: 启动内核（8 个并行块）
    unsafe {
        hello_world_do(8, stream.to_raw());
    }

    // 步骤 5: 同步等待内核完成
    stream.synchronize()?;

    // 步骤 6: 所有资源自动释放（RAII）
    // Drop 顺序: stream → context → device → acl
    Ok(())
}

关键设计：生命周期链

注意这段代码的类型签名：

Acl                    → 生命周期根
  Device<'acl>         → 必须在 Acl 之前析构
    AclContext<'d>     → 必须在 Device 之前析构
      AclStream<'c>   → 必须在 Context 之前析构

如果你试图以错误的顺序使用这些资源，代码将无法通过编译。 这是 Rust 类型系统的力量——在编译期保证了资源管理的正确性，而 C++ 只能依赖程序员的纪律。

对比：C++ 版本的隐患

等价的 C++ 代码需要手动管理每个资源的生命周期：

// C++ 版本：每个资源都需要手动释放
aclInit(nullptr);
aclrtSetDevice(0);
aclrtContext ctx;
aclrtCreateContext(&ctx, 0);
aclrtStream stream;
aclrtCreateStream(&stream);

hello_world_do(8, stream);
aclrtSynchronizeStream(stream);

// 必须按正确顺序手动释放，否则导致未定义行为
aclrtDestroyStream(stream);
aclrtDestroyContext(ctx);
aclrtResetDevice(0);
aclFinalize();

如果任何一步抛出异常或提前返回，后续的清理代码将被跳过。而 Rust 版本中，Drop trait 保证了无论控制流如何变化，资源都会被正确释放。