零拷贝内存
在流中，我们介绍了cudaHostAlloc这个函数，它有一些标志，其中cudaHostAllocMapped允许内存映射到设备，也即GPU可以直接访问主机上的内存，不用额外再给设备指针分配内存
通过下面的操作，即可让设备指针也可访问主机内存
cudaHostAlloc((void**)&a, N * sizeof(float), cudaHostAllocWriteCombined | cudaHostAllocMapped);
cudaHostGetDevicePointer(&dev_a, a, 0); // 将主机指针映射为设备可用指针
由于GPU虚拟内存空间和CPU不同，不能直接使用指针a，必须调用cudaHostGetDevicePointer函数；这样 dev_a 就是设备端可以直接访问的 host 内存。
原理简介
-
在调用
cudaHostAllocMapped时，CUDA 会在主机申请一块 页锁定内存（pinned memory）； -
再通过
cudaHostGetDevicePointer把这块主机内存映射为设备端地址空间中的指针； -
当 GPU 访问
dev_a[i]时，会通过 PCIe 总线从主机 RAM 中取数据，实现 零拷贝访问。
所以它虽然“看起来像显存指针”，但其实访问的是主机内存。
下面用该机制重写cuda编程笔记（2.5）--简易的应用代码-CSDN博客里的矢量点乘
#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <cuda_runtime.h>
#include <device_launch_parameters.h>#include <iostream>
#include<cstdio>#define threadsPerBlock 256
const int Blocks = 32;
const int N = Blocks * threadsPerBlock;void error_handling(cudaError_t res) {if (res !=cudaSuccess) {std::cout << "error!" << std::endl;}
}
__global__ void dot(float* a, float* b, float* c) {__shared__ float cache[threadsPerBlock];int tid = threadIdx.x + blockIdx.x * blockDim.x;int cacheIndex = threadIdx.x;float temp = 0;if (tid < N) temp = a[tid] * b[tid];cache[cacheIndex] = temp;__syncthreads();for (int stride = blockDim.x / 2; stride > 0; stride>>= 1) {if (cacheIndex < stride)cache[cacheIndex] += cache[cacheIndex + stride];__syncthreads();}// 将每个 block 的结果写入全局内存if (cacheIndex == 0) {c[blockIdx.x] = cache[0];}}
int main() {cudaEvent_t start, stop;float* a, * b, c, * partial_c;float* dev_a, * dev_b, * dev_partial_c;float elapsedTime;error_handling(cudaEventCreate(&start));error_handling(cudaEventCreate(&stop));//在cpu上分配内存error_handling(cudaHostAlloc((void**)&a, N * sizeof(float),cudaHostAllocWriteCombined|cudaHostAllocMapped));error_handling(cudaHostAlloc((void**)&b, N * sizeof(float), cudaHostAllocWriteCombined | cudaHostAllocMapped));error_handling(cudaHostAlloc((void**)&partial_c, Blocks * sizeof(float), cudaHostAllocWriteCombined | cudaHostAllocMapped));for (int i = 0; i < N; i++) {a[i] = i;b[i] = i * 2;}error_handling(cudaHostGetDevicePointer(&dev_a, a, 0));error_handling(cudaHostGetDevicePointer(&dev_b, b, 0));error_handling(cudaHostGetDevicePointer(&dev_partial_c, partial_c, 0));error_handling(cudaEventRecord(start, 0));dot << < Blocks, threadsPerBlock >> > (dev_a, dev_b, dev_partial_c);error_handling(cudaDeviceSynchronize());error_handling(cudaEventRecord(stop, 0));error_handling(cudaEventSynchronize(stop));error_handling(cudaEventElapsedTime(&elapsedTime, start, stop));c = 0;for (int i = 0; i < Blocks; i++)c += partial_c[i];error_handling(cudaFreeHost(a));error_handling(cudaFreeHost(b));error_handling(cudaFreeHost(partial_c));error_handling(cudaEventDestroy(start));error_handling(cudaEventDestroy(stop));printf("Value calculated: %f\n", c);printf("Time consumed:%f\n", elapsedTime);
}
| 优点 | 说明 |
|---|---|
减少显式 cudaMemcpy 调用 | 主机 → 设备零拷贝 |
| 避免重复申请/释放显存 | 数据只分配一次 |
| 简化代码结构 | 多个内核之间共享同一 host 指针 |
| 适合小规模、实时更新场景 | 如 GUI 控件、摄像头图像 |
| 缺点 | 说明 |
|---|---|
| 访问速度远慢于 global memory | 因为要通过 PCIe |
| 仅适用于某些 GPU（如支持 UVA） | 非所有设备支持 |
| 最佳性能只在小数据量/零延迟访问场景 | 比如小型图像处理、调试等 |
| 受限于 CPU 内存页 | 页大小影响效率，不能高并发 |
使用条件
| 要点 | 说明 |
|---|
| GPU 必须支持 UVA（统一虚拟地址空间） | 可用 cudaGetDeviceProperties() 查询 unifiedAddressing 是否为 1 |
最好配合 WriteCombined | 适合只写不读场景（如从主机写入，GPU 读取） |
| 不适合大规模数据训练/推理 | 会严重拖慢 GPU 性能，PCIe 带宽远小于显存带宽 |
启动多GPU
使用多个线程，就可以同时启动多个 GPU 来并行计算，这是现代 CUDA 编程中非常推荐且常用的做法。
CUDA 的执行模型是：
-
每个 CPU 线程 通过
cudaSetDevice(id)绑定到某个 GPU -
每个线程可以在绑定的 GPU 上：
-
分配显存
-
启动 kernel
-
执行 memcpy
-
做同步
-
CUDA runtime 为每个 CPU 线程维护独立的 GPU 上下文（context），所以 不同线程绑定不同 GPU，就可以各自独立调度、执行自己的 kernel。
#include <thread>
#include <iostream>__global__ void kernel(int id) {printf("Hello from GPU %d, thread %d\n", id, threadIdx.x);
}void gpu_task(int device_id) {cudaSetDevice(device_id);kernel<<<1, 4>>>(device_id);cudaDeviceSynchronize(); // 等待 GPU 完成
}int main() {int num_devices = 0;cudaGetDeviceCount(&num_devices);std::vector<std::thread> threads;for (int i = 0; i < num_devices; ++i) {threads.emplace_back(gpu_task, i); // 每个线程负责一个 GPU}for (auto& t : threads) t.join(); // 等待所有线程完成return 0;
}
多 GPU 场景下共享主机内存
cudaHostAlloc中当flags传入cudaHostAllocPortable时
就意味着：
✅ 分配出的主机内存是跨 GPU 可见（portable）的，不属于某个特定的 GPU 上下文。
为什么多 GPU 编程中需要 cudaHostAllocPortable？
在默认情况下（无 cudaHostAllocPortable）：
-
使用
cudaHostAlloc()分配的内存只绑定到当前 GPU 上下文； -
如果你在另一个 GPU 上使用该内存（比如调用
cudaMemcpyAsync），就会报错或性能下降。
加上 cudaHostAllocPortable 后：
-
这块页锁定内存在所有 GPU 上都能直接访问（只要硬件支持 UVA）。
典型用法：多 GPU + Portable 内存
float *host_ptr;
cudaHostAlloc((void**)&host_ptr, N * sizeof(float), cudaHostAllocPortable);
然后每个线程可以这样操作：
void run_on_device(int device_id, float* shared_host) {cudaSetDevice(device_id);float *dev_ptr;cudaMalloc(&dev_ptr, N * sizeof(float));// 每个 GPU 从共享主机内存拷贝数据cudaMemcpy(dev_ptr, shared_host, N * sizeof(float), cudaMemcpyHostToDevice);kernel<<<blocks, threads>>>(dev_ptr);cudaDeviceSynchronize();cudaFree(dev_ptr);
}
这样，每个 GPU 都能用同一块主机内存 shared_host 来做数据初始化、写回、交换数据等操作。
常见组合：
cudaHostAllocPortable | cudaHostAllocWriteCombined
GPU A 写结果，GPU B 读取验证
GPU A 写入 shared host memory，GPU B 读取验证是完全可能出现同步问题的
线程之间需要加同步
#include <cuda_runtime.h>
#include <iostream>
#include <thread>
#include <vector>
#include <cassert>#define N 16__global__ void write_kernel(int *data, int val) {int idx = threadIdx.x;if (idx < N) {data[idx] = val * 100 + idx;}
}__global__ void read_kernel(int *data) {int idx = threadIdx.x;if (idx < N) {printf("GPU 1 reads: data[%d] = %d\n", idx, data[idx]);}
}// GPU 0 线程函数：写入共享主机内存
void gpu0_writer(int *host_data, cudaEvent_t write_done_event) {cudaSetDevice(0);cudaStream_t stream;cudaStreamCreate(&stream);int *dev_data;cudaMalloc(&dev_data, N * sizeof(int));write_kernel<<<1, N, 0, stream>>>(dev_data, 1);// 将数据从设备拷贝到共享主机内存cudaMemcpyAsync(host_data, dev_data, N * sizeof(int), cudaMemcpyDeviceToHost, stream);// 记录写入完成事件cudaEventRecord(write_done_event, stream);cudaStreamSynchronize(stream);cudaFree(dev_data);cudaStreamDestroy(stream);std::cout << "[GPU 0] 写入完成\n";
}// GPU 1 线程函数：等待事件后读取共享主机内存
void gpu1_reader(int *host_data, cudaEvent_t write_done_event) {cudaSetDevice(1);cudaStream_t stream;cudaStreamCreate(&stream);// 等待 GPU 0 写入完成cudaStreamWaitEvent(stream, write_done_event, 0);int *dev_data;cudaMalloc(&dev_data, N * sizeof(int));// 从共享主机内存拷贝到 GPU 1 上的显存cudaMemcpyAsync(dev_data, host_data, N * sizeof(int), cudaMemcpyHostToDevice, stream);read_kernel<<<1, N, 0, stream>>>(dev_data);cudaStreamSynchronize(stream);cudaFree(dev_data);cudaStreamDestroy(stream);std::cout << "[GPU 1] 读取完成\n";
}int main() {int gpu_count = 0;cudaGetDeviceCount(&gpu_count);if (gpu_count < 2) {std::cerr << "需要至少两个 GPU！\n";return -1;}// 分配共享主机内存（portable）int *shared_host_data;cudaHostAlloc((void**)&shared_host_data, N * sizeof(int), cudaHostAllocPortable);// 创建用于跨 GPU 通信的事件cudaEvent_t write_done_event;cudaEventCreateWithFlags(&write_done_event, cudaEventDisableTiming); // faster event// 启动两个线程std::thread t0(gpu0_writer, shared_host_data, write_done_event);std::thread t1(gpu1_reader, shared_host_data, write_done_event);t0.join();t1.join();cudaEventDestroy(write_done_event);cudaFreeHost(shared_host_data);return 0;
}
cudaEventCreateWithFlags
事件创建：cudaEventCreateWithFlags
cudaEvent_t evt;
cudaEventCreateWithFlags(&evt, cudaEventDisableTiming); // 推荐带标志创建更轻量
| 标志 | 含义 | 说明 |
|---|---|---|
cudaEventDefault | 默认行为 | 会记录耗时，可用于性能计时 |
cudaEventDisableTiming | 禁用计时功能 | 更轻量，推荐用于同步控制 |
cudaEventInterprocess | 可用于多进程共享事件 | 不常用于多 GPU 同步（属于高级功能） |
cudaEventRecord
表示 之前所有 stream中的操作都完成时，该事件被标记完成。
cudaStreamWaitEvent
cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags);
| 参数 | 类型 | 含义 |
|---|---|---|
stream | cudaStream_t | 要等待事件的 CUDA 流。这个 stream 将在 event 被触发后才开始执行其后续任务。 |
event | cudaEvent_t | 要等待的事件。这个事件应该在其他设备或流上通过 cudaEventRecord 创建。 |
flags | unsigned int | 当前必须设为 0。（CUDA 12.4 以前不支持其他选项） |
![《[系统底层攻坚] 张冬〈大话存储终极版〉精读计划启动——存储架构原理深度拆解之旅》-系统性学习笔记（适合小白与IT工作人员）](http://pic.xiahunao.cn/《[系统底层攻坚] 张冬〈大话存储终极版〉精读计划启动——存储架构原理深度拆解之旅》-系统性学习笔记（适合小白与IT工作人员）)
)
