前言

刚写Ascend C算子那会,我不理解为什么算子里不能随便用指针,为什么数据搬运必须调rtMemcpy,为什么多流要加Event。后来看了runtime的源码才明白——runtime是NPU上的"操作系统",你写的算子是在它的管理下运行的。不理解runtime,你写的Ascend C算子就是"裸奔"——能跑,但不知道为什么跑不动、为什么跑不快。

这篇文章是我对runtime的理解:它是什么、不是什么、为什么要听它调度。

认知纠偏:runtime ≠ 驱动 ≠ 编译器

很多人把runtime、驱动、编译器搞混。它们在CANN架构中是三层完全不同的东西:

名称 做什么 运行在哪
第5层 驱动(driver) 管硬件:寄存器读写、中断处理、电源管理 内核态
第4层 runtime 管资源:内存、流、事件、设备 用户态
第3层 GE/ATC(编译器) 管计算图编译和优化 用户态(编译时)

驱动让NPU通电,runtime让NPU干活,编译器告诉NPU怎么干。

runtime不是驱动:驱动在内核态,直接操作硬件寄存器。runtime在用户态,通过系统调用跟驱动交互。你写应用代码时只跟runtime打交道,不需要直接调驱动。

runtime不是编译器:编译器在"编译时"把计算图翻译成NPU能执行的任务序列。runtime在"运行时"把这些任务调度到NPU上执行。编译器负责"翻译",runtime负责"执行"。

runtime更像操作系统:它管理NPU上的资源(内存、流、事件),调度任务执行,提供系统调用接口。你在NPU上做的一切事情,都经过runtime。

runtime的四大核心职责

职责一:内存管理

runtime管理NPU的HBM(High Bandwidth Memory)和Host内存,提供统一的内存分配/释放/搬运接口:

#include <acl/acl.h>

// 1. 分配Device内存(HBM)
void* dev_ptr = NULL;
size_t size = 1024 * 1024;  // 1MB
aclError ret = aclrtMalloc(&dev_ptr, size, ACL_MEM_MALLOC_HUGE_FIRST);

// 2. 分配Host内存(CPU侧,用于跟Device交互)
void* host_ptr = NULL;
ret = aclrtMallocHost(&host_ptr, size);

// 3. 数据搬运:Host → Device
ret = aclrtMemcpy(dev_ptr, size, host_ptr, size, ACL_MEMCPY_HOST_TO_DEVICE);

// 4. 数据搬运:Device → Host
ret = aclrtMemcpy(host_ptr, size, dev_ptr, size, ACL_MEMCPY_DEVICE_TO_HOST);

// 5. 数据搬运:Device → Device(跨卡)
ret = aclrtMemcpy(dev_ptr_dst, size, dev_ptr_src, size, ACL_MEMCPY_DEVICE_TO_DEVICE);

// 6. 释放内存
aclrtFree(dev_ptr);
aclrtFreeHost(host_ptr);

为什么必须用runtime的内存接口?三个原因:

  1. 虚拟地址映射:runtime分配的HBM地址是虚拟地址,经过MMU映射到物理地址。直接用物理地址访问会越界。
  2. 内存池管理:runtime内部维护内存池,频繁分配/释放不会每次都调驱动(系统调用开销大)。
  3. 缓存一致性:Device和Host之间的数据搬运,runtime自动处理缓存刷新和失效。

职责二:流调度

流(Stream)是runtime的任务队列。你往流里提交任务(算子执行、数据搬运),runtime按顺序调度到NPU上执行。

// 1. 创建流
aclrtStream stream;
aclrtCreateStream(&stream);

// 2. 在流上执行算子
// 异步提交,立即返回
aclrtLaunch(kernel_func, block_dim, args, args_size, stream);

// 3. 在流上做数据搬运
aclrtMemcpyAsync(dev_ptr, size, host_ptr, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);

// 4. 同步等待流上所有任务完成
aclrtSynchronizeStream(stream);

// 5. 销毁流
aclrtDestroyStream(stream);

多流并行:创建多个流,不同流上的任务并行执行。比如Stream 1做计算,Stream 2做数据搬运,计算和搬运重叠:

aclrtStream compute_stream, copy_stream;
aclrtCreateStream(&compute_stream);
aclrtCreateStream(&copy_stream);

// Stream 1: 计算当前batch
aclrtLaunch(kernel_func, ..., compute_stream);

// Stream 2: 搬运下一个batch的数据
aclrtMemcpyAsync(next_batch_dev, size, next_batch_host, size,
                 ACL_MEMCPY_HOST_TO_DEVICE, copy_stream);

// 等两条流都完成
aclrtSynchronizeStream(compute_stream);
aclrtSynchronizeStream(copy_stream);

流是硬件队列,不是软件线程。CUDA的Stream也是硬件队列,但昇腾的流调度延迟更低(~2μs vs CUDA的~5μs),因为昇腾的调度器在硬件上实现。

职责三:同步机制

Event是runtime的同步原语,用于流间同步:

// 1. 创建Event
aclrtEvent event;
aclrtCreateEvent(&event);

// 2. 在Stream A上记录Event
aclrtRecordEvent(event, stream_a);

// 3. 在Stream B上等待Event
aclrtStreamWaitEvent(stream_b, event);
// Stream B会阻塞,直到Stream A执行到Event记录点

// 4. 销毁Event
aclrtDestroyEvent(event);

为什么需要Event? 多流并行时,流之间有依赖关系。比如Stream 1搬数据到Device,Stream 2用这些数据做计算——Stream 2必须等Stream 1搬完才能开始。Event就是"搬完了"的信号。

职责四:设备管理

多卡场景下,runtime管理多张NPU设备:

// 1. 获取设备数量
uint32_t device_count;
aclrtGetDeviceCount(&device_count);

// 2. 设置当前设备(类似cudaSetDevice)
aclrtSetDevice(0);  // 使用第0张NPU

// 3. 在当前设备上分配内存、创建流、执行算子
// ...

// 4. 切换到另一张NPU
aclrtSetDevice(1);

// 5. 重置设备(释放该设备上所有资源)
aclrtResetDevice(0);

runtime vs CUDA Runtime:关键差异

维度 昇腾 runtime CUDA Runtime
内存分配 aclrtMalloc cudaMalloc
数据搬运 aclrtMemcpy(显式指定方向) cudaMemcpy(自动推断方向)
流调度延迟 ~2μs ~5μs
Event同步 aclrtStreamWaitEvent cudaStreamWaitEvent
多卡管理 aclrtSetDevice cudaSetDevice
错误处理 返回aclError枚举 返回cudaError_t枚举

最大的差异aclrtMemcpy需要显式指定方向(HOST_TO_DEVICE / DEVICE_TO_HOST / DEVICE_TO_DEVICE),cudaMemcpy会根据指针自动推断。这是runtime的设计选择——显式比隐式更安全,避免方向推断错误导致的性能陷阱。

完整代码示例:runtime的完整使用流程

import torch
import torch_npu

# 1. 初始化runtime(PyTorch集成后自动初始化)
device = torch.device("npu:0")

# 2. 分配内存(PyTorch的.npu()封装了aclrtMalloc)
a = torch.randn(1024, 1024, dtype=torch.float16, device=device)
b = torch.randn(1024, 1024, dtype=torch.float16, device=device)

# 3. 创建流
stream_compute = torch.npu.Stream(device=device)
stream_copy = torch.npu.Stream(device=device)

# 4. 多流并行:计算和搬运重叠
with torch.npu.stream(stream_compute):
    c = torch.matmul(a, b)  # Stream 1: 计算

with torch.npu.stream(stream_copy):
    d = a.npu_to_cpu()  # Stream 2: Device → Host搬运

# 5. 同步
stream_compute.synchronize()
stream_copy.synchronize()

print(f"计算结果: {c.shape}")    # [1024, 1024]
print(f"搬运结果: {d.shape}")    # [1024, 1024]

# 6. 释放资源(PyTorch自动管理,不需要手动aclrtFree)

PyTorch + torch-npu把runtime的大部分接口封装成了Python API,日常开发不需要直接调C接口。但理解runtime的工作原理,对调试性能问题和排查bug很有帮助。

踩坑实录

坑1:aclrtMalloc分配的是虚拟地址

问题:拿到dev_ptr后,想用mmap映射到用户空间,报段错误。

原因aclrtMalloc返回的地址是NPU侧的虚拟地址,不在CPU的虚拟地址空间里。CPU不能直接通过指针访问NPU HBM。

解决方案:用aclrtMemcpy搬运数据,不要直接解引用:

// ❌ 错误写法(CPU不能直接访问NPU虚拟地址)
float* dev_ptr = aclrtMalloc(...);
float value = dev_ptr[0];  // 段错误!

// ✅ 正确写法(通过aclrtMemcpy搬运)
float host_value;
aclrtMemcpy(&host_value, sizeof(float), dev_ptr, sizeof(float),
            ACL_MEMCPY_DEVICE_TO_HOST);

坑2:多流共享内存必须加Event同步

问题:Stream 1写数据,Stream 2读同一块数据,读到半新半旧的值。

原因:两条流并行执行,没有同步保证。Stream 1写到一半时,Stream 2可能已经在读了。

解决方案:Stream 1写完后Record Event,Stream 2 Wait Event再读:

aclrtEvent event;
aclrtCreateEvent(&event);

// Stream 1: 写数据
aclrtLaunch(write_kernel, ..., stream1);
aclrtRecordEvent(event, stream1);  // 写完记录Event

// Stream 2: 读数据(必须等Stream 1写完)
aclrtStreamWaitEvent(stream2, event);  // 等Event
aclrtLaunch(read_kernel, ..., stream2);  // 然后读

坑3:aclrtSynchronizeStream是阻塞调用

问题:主线程调aclrtSynchronizeStream,整个进程卡住,UI没响应。

原因aclrtSynchronizeStream是阻塞调用,会等到流上所有任务完成。如果流上有长任务(比如训一个大batch),主线程可能阻塞几秒甚至几分钟。

解决方案:用非阻塞方式检查流状态,或者把同步放到子线程:

# 方案A:非阻塞检查
import torch_npu

stream = torch.npu.Stream()
# ... 提交任务到stream ...

# 非阻塞查询
if stream.query():  # 返回True表示所有任务完成
    result = get_result()
else:
    # 还没完成,先做别的事
    do_other_work()

# 方案B:子线程同步
import threading

def wait_and_get_result():
    stream.synchronize()
    return get_result()

thread = threading.Thread(target=wait_and_get_result)
thread.start()
# 主线程继续响应UI

runtime在CANN架构中的位置

runtime位于CANN五层架构的第4层(昇腾计算执行层),是连接上层框架和底层驱动的桥梁:

第1层:AscendCL(应用接口)
  ↓ 调用
第4层:Runtime(资源管理和任务调度)← 你在这里
  ↓ 调用
第5层:Driver(硬件管理)

几乎所有CANN组件最终都经过runtime:

  • AscendCL → runtime(应用开发接口封装runtime)
  • GE → runtime(图编译后通过runtime执行)
  • HCCL → runtime(集合通信通过runtime调度)
  • 算子 → runtime(算子通过runtime启动)

结尾

理解runtime是写好Ascend C算子的前提。它管内存、管流、管同步、管设备——你写的算子是在它的管理下运行的。不理解它,你不知道为什么数据搬运要调rtMemcpy(因为虚拟地址映射)、为什么多流要加Event(因为没有同步保证)、为什么内存不能直接用指针访问(因为CPU和NPU的地址空间不同)。

runtime不是你的障碍,是你的安全网。它帮你管理NPU资源,让你专注于算子逻辑本身。下次你的算子跑不动或者跑不快,先想想runtime在背后做了什么——答案往往就在那里。

https://atomgit.com/cann/runtime

Logo

openEuler 是由开放原子开源基金会孵化的全场景开源操作系统项目,面向数字基础设施四大核心场景(服务器、云计算、边缘计算、嵌入式),全面支持 ARM、x86、RISC-V、loongArch、PowerPC、SW-64 等多样性计算架构