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）

注意

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 的过程对用户基本透明，用户看到的是一组函数调用。

---

相关文档

• GPU 架构基础：CPU vs GPU 差异、SIMT 模型
• 线程层次结构：Grid/Block/Thread 组织
• 内存层次结构：寄存器/共享内存/全局内存
• 执行模型：SIMT 并行执行