Host/Device 编程模型
假设你已了解 CPU vs GPU 架构差异和 SIMT 执行模型，本文将帮助你深入理解 Host 与 Device 在编程中的分工、通信和协作——内存管理、数据传输、流并发、事件同步。
Host 与 Device 的职责
| Host（CPU）职责 | Device（GPU）职责 |
|---|---|
| 程序主逻辑流（顺序代码） | 并行 Kernel 执行 |
| 内存分配与释放 | 大规模数据并行处理 |
| 数据拷贝（Host ↔ Device） | 高吞吐计算 |
| Kernel 配置与启动 | 接收并执行 Kernel |
| 同步与等待 | 返回执行结果 |
关键原则：
- Host 代码是顺序的，Device 代码是并行的
- Host 控制"做什么"，Device 负责"怎么做"
- Host 和 Device 拥有独立的内存地址空间
内存管理与数据传输
内存空间
| 内存空间 | 位置 | 访问者 | 特点 |
|---|---|---|---|
| 主机内存 | CPU 系统内存 | CPU | 低延迟，容量大（8-64GB） |
| 设备内存 | GPU 显存 | GPU | 高带宽，容量适中（8-32GB） |
warning
CPU 和 GPU 无法直接访问对方的内存，必须通过 musaMemcpy 显式传输。
内存管理 API
#include <musa_runtime.h>
// 分配设备内存
float* d_data;
musaMalloc(&d_data, size * sizeof(float));
// 释放设备内存
musaFree(d_data);
数据传输 API
// Host → Device
musaMemcpy(d_data, h_data, size, musaMemcpyHostToDevice);
// Device → Host
musaMemcpy(h_result, d_result, size, musaMemcpyDeviceToHost);
// Device → Device（多 GPU 场景）
musaMemcpy(d_dst, d_src, size, musaMemcpyDeviceToDevice);
// Host ←→ Host（内存初始化）
musaMemcpy(h_data, h_src, size, musaMemcpyHostToHost);
方向枚举
| 枚举值 | 方向 |
|---|---|
musaMemcpyHostToDevice | CPU → GPU |
musaMemcpyDeviceToHost | GPU → CPU |
musaMemcpyDeviceToDevice | GPU → GPU |
musaMemcpyHostToHost | CPU → CPU |
统一内存：musaMallocManaged() 分配统一内存，CPU 和 GPU 共享同一地址空间，无需显式拷贝。但需要了解其性能特性（首次访问可能触发页面迁移）。
内核执行
内核定义
// __global__ 修饰器表示这是一个内核函数
__global__ void myKernel(float* data, int n) {
// 计算全局线程索引
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
data[idx] *= 2.0f;
}
}
内核基础语法见 GPU 并行计算基础。
内核启动
// 配置执行参数
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
// 启动内核：<<<gridDim, blockDim>>>
myKernel<<<gridSize, blockSize>>>(d_data, n);
执行配置参数
| 参数 | 含义 |
|---|---|
gridDim | 网格维度（线程块数量） |
blockDim | 线程块维度（每线程块线程数） |
blockIdx | 当前线程块索引 |
threadIdx | 当前线程索引 |
流（Stream）—— 并发执行
为什么需要流
默认情况下，Kernel 在默认流中顺序执行：
// 顺序执行：Kernel2 必须等待 Kernel1 完成
kernel1<<<gridSize, blockSize>>>(data1); // 假设耗时 10ms
kernel2<<<gridSize, blockSize>>>(data2); // 等待 Kernel1 完成后才开始
kernel3<<<gridSize, blockSize>>>(data3); // 等待 Kernel2 完成后才开始
// 总耗时：约 30ms
使用流可以并行执行独立的操作：
// 创建 3 个流
musaStream_t stream1, stream2, stream3;
musaStreamCreate(&stream1);
musaStreamCreate(&stream2);
musaStreamCreate(&stream3);
// 并行执行：3 个 Kernel 可能在不同 MP 上并发执行
kernel1<<<gridSize, blockSize, 0, stream1>>>(data1);
kernel2<<<gridSize, blockSize, 0, stream2>>>(data2);
kernel3<<<gridSize, blockSize, 0, stream3>>>(data3);
// 销毁流
musaStreamDestroy(stream1);
musaStreamDestroy(stream2);
musaStreamDestroy(stream3);
// 总耗时：约 10ms（取决于硬件资源）
流 API
#include <musa_runtime.h>
// 创建流
musaStream_t stream;
musaStreamCreate(&stream);
// 在流中启动 Kernel
myKernel<<<gridSize, blockSize, 0, stream>>>(data);
// 在流中执行数据传输（异步）
musaMemcpyAsync(d_data, h_data, size, musaMemcpyHostToDevice, stream);
// 等待流完成（阻塞 Host）
musaStreamSynchronize(stream);
// 查询流状态（非阻塞）
musaError_t status = musaStreamQuery(stream);
if (status == musaSuccess) {
// 流中所有操作已完成
} else if (status == musaErrorNotReady) {
// 流中操作仍在执行
}
// 销毁流
musaStreamDestroy(stream);
流的类型
| 流类型 | 说明 | 创建方式 |
|---|---|---|
| 默认流 | 所有操作的默认队列 | 无需创建，stream = 0 |
| 普通流 | 独立的操作队列 | musaStreamCreate() |
| 高性能流 | 优先调度，低延迟 | musaStreamCreateWithPriority() |
使用场景
| 场景 | 说明 |
|---|---|
| 多 Kernel 并发 | 独立的数据并行处理 |
| 数据流水线 | 当前 Kernel 执行时，预取下一批数据 |
| 多设备协同 | 在不同 GPU 上并发执行任务 |
事件（Event）—— 同步与计时
事件同步
事件是流中的标记点，可用于同步操作：
#include <musa_runtime.h>
// 创建事件
musaEvent_t event;
musaEventCreate(&event);
// 在流中记录事件
myKernel<<<gridSize, blockSize, 0, stream>>>(data);
musaEventRecord(event, stream);
// 等待事件完成（阻塞 Host）
musaEventSynchronize(event);
// 或者让另一个流等待该事件
musaStreamWaitEvent(stream2, event, 0);
// 销毁事件
musaEventDestroy(event);
事件计时
事件可用于精确测量 Kernel 执行时间：
musaEvent_t start, stop;
musaEventCreate(&start);
musaEventCreate(&stop);
// 记录开始事件
musaEventRecord(start, stream);
// 执行 Kernel
myKernel<<<gridSize, blockSize, 0, stream>>>(data);
// 记录结束事件
musaEventRecord(stop, stream);
// 等待完成
musaEventSynchronize(stop);
// 获取耗时（毫秒）
float elapsed_ms;
musaEventElapsedTime(&elapsed_ms, start, stop);
printf("Kernel 耗时：%f ms\n", elapsed_ms);
// 销毁事件
musaEventDestroy(start);
musaEventDestroy(stop);
事件 API 汇总
| API | 说明 |
|---|---|
musaEventCreate(&event) | 创建事件 |
musaEventRecord(event, stream) | 在流中记录事件 |
musaEventSynchronize(event) | 等待事件完成（阻塞 Host） |
musaEventElapsedTime(&ms, start, stop) | 计算两个事件之间的时间（ms） |
musaEventDestroy(event) | 销毁事件 |
任务流水线（驱动层）
从应用调用到 GPU 执行的完整流程（对用户透明）：
对用户的透明度：整个 CPU ↔ 驱动 ↔ GPU 的过程对用户基本透明，用户看到的是一组函数调用。
