runtime:昇腾NPU的“操作系统“,为什么你的算子要听它调度?
本文深入解析了昇腾NPU的runtime系统,指出其核心定位是"NPU上的操作系统",负责资源管理和任务调度。文章首先澄清了runtime与驱动、编译器的区别,强调runtime在用户态管理NPU资源。随后详细阐述了runtime的四大职责:内存管理(虚拟地址映射、内存池优化)、流调度(多流并行)、同步机制(Event)和设备管理(多卡切换)。通过与CUDA Runtime的对比,突出了昇腾run
前言
刚写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的内存接口?三个原因:
- 虚拟地址映射:runtime分配的HBM地址是虚拟地址,经过MMU映射到物理地址。直接用物理地址访问会越界。
- 内存池管理:runtime内部维护内存池,频繁分配/释放不会每次都调驱动(系统调用开销大)。
- 缓存一致性: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(©_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
所有评论(0)