顺德 网站设计建设项目银行网站
cudaHostAlloc
cudaHostAlloc 是 CUDA 提供的用于在 主机端申请页锁定内存（pinned memory） 的函数
cudaHostAlloc() 和 malloc() 的核心区别
| 特性/方面 | malloc()（标准 C） | cudaHostAlloc()（CUDA） |
|---|---|---|
| 内存类型 | 普通主机内存（分页内存） | 页锁定内存（pinned memory） |
| 是否支持 DMA 传输 | ❌ 需要复制缓冲区 | ✅ 可直接使用 DMA 提升带宽 |
| 是否支持异步传输 | ❌ 不支持 | ✅ 支持与 CUDA 流异步 cudaMemcpyAsync |
| 是否支持映射给 GPU | ❌ 不支持 | ✅ 可映射到设备端地址空间（取决于 flag） |
| 释放函数 | free() | cudaFreeHost() |
| 是否影响主机调度 | 否 | 是（锁页后主机无法交换该页） |
cudaHostAlloc() 函数原型与参数
cudaError_t cudaHostAlloc(void** ptr, size_t size, unsigned int flags);
参数说明：
-
void** ptr：返回页锁定主机内存的地址； -
size：申请的字节数； -
flags：控制内存行为，常见值如下：
| Flag | 含义 |
|---|---|
cudaHostAllocDefault | 默认锁页内存，适用于异步传输 |
cudaHostAllocMapped | 允许内存映射到设备（即 GPU 可以直接访问） |
cudaHostAllocWriteCombined | 禁止读优化，仅用于主机写后设备读，提高性能 |
cudaHostAllocPortable | 多 GPU 上可见（跨 context 使用） |
什么是页锁定的内存
页锁定内存（pinned memory）是指不会被操作系统虚拟内存管理机制移到磁盘的主机内存，
它是物理内存中被“锁住”的一段区域，可用于GPU 的直接访问（DMA）和异步数据传输
在操作系统中，主机内存通常是分页（pageable）管理的：
-
普通内存（如
malloc()得到的）：是可分页的；-
如果内存长时间未访问，操作系统可能把它换出到磁盘（swap），释放物理内存。
-
-
页锁定内存（如
cudaHostAlloc()得到的）：是不可分页的；-
它始终驻留在主机物理内存中，不能被 swap 掉。
-
为什么这很重要？
因为 GPU 想要高效读取主机内存时，必须使用 DMA（直接内存访问），而 DMA 不能访问“随时可能被换出”的普通内存。
| 优势 | 说明 |
|---|---|
| ✅ 支持 异步内存传输 | 允许 cudaMemcpyAsync 与 kernel 并发 |
| ✅ 支持 更高带宽的内存传输 | 通过 DMA 直连传输，避免临时缓存 |
| ✅ 可被 GPU 映射访问 | 如使用 cudaHostAllocMapped |
| ✅ 节省系统临时复制开销 | 避免内部“拷贝到临时 pinned 缓冲区”的中间过程 |
| 问题 | 原因与建议 |
|---|---|
| ❗ pinned memory 会限制系统物理内存 | 不应大规模分配，避免影响主机性能 |
🔄 必须用 cudaFreeHost() 释放 | 不要用 free() |
❗ 必须配合流和 cudaMemcpyAsync 才能异步传输 | 否则性能不会提升 |
速度测试
#include <cuda_runtime.h>
#include <device_launch_parameters.h>#include <iostream>
#include<cstdio>
#define SIZE 10*1024*1024
void error_handling(cudaError_t res) {if (res !=cudaSuccess) {std::cout << "error!" << std::endl;}
}
//up是内存分配的方向
float cuda_malloc_test(int size, bool up) {cudaEvent_t start, stop;int* a, * dev_a;float elapsedTime;error_handling(cudaEventCreate(&start));error_handling(cudaEventCreate(&stop));a = new int[size];error_handling(cudaMalloc((void**)&dev_a,size*sizeof(int)));error_handling(cudaEventRecord(start,0));//执行100复制操作for (int i = 0; i < 100; i++) {if (up) {//主机到设备error_handling(cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice));}else {//设备到主机error_handling(cudaMemcpy(a, dev_a, size * sizeof(int), cudaMemcpyDeviceToHost));}}error_handling(cudaEventRecord(stop, 0));error_handling(cudaEventSynchronize(stop));error_handling(cudaEventElapsedTime(&elapsedTime, start, stop));delete[]a;error_handling(cudaFree(dev_a));error_handling(cudaEventDestroy(start));error_handling(cudaEventDestroy(stop));return elapsedTime;
}
float cuda_host_malloc_test(int size, bool up) {cudaEvent_t start, stop;int* a, * dev_a;float elapsedTime;error_handling(cudaEventCreate(&start));error_handling(cudaEventCreate(&stop));//error_handling(cudaHostAlloc((void**)&a, size * sizeof(int),cudaHostAllocDefault));//error_handling(cudaMalloc((void**)&dev_a, size * sizeof(int)));error_handling(cudaEventRecord(start, 0));//执行100复制操作for (int i = 0; i < 100; i++) {if (up) {//主机到设备error_handling(cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice));}else {//设备到主机error_handling(cudaMemcpy(a, dev_a, size * sizeof(int), cudaMemcpyDeviceToHost));}}error_handling(cudaEventRecord(stop, 0));error_handling(cudaEventSynchronize(stop));error_handling(cudaEventElapsedTime(&elapsedTime, start, stop));error_handling(cudaFreeHost(a));error_handling(cudaFree(dev_a));error_handling(cudaEventDestroy(start));error_handling(cudaEventDestroy(stop));return elapsedTime;
}
int main() {float elapsedTime;float MB = (float)100 * SIZE * sizeof(int) / 1024/1024;//可以自行修改调用的函数，进行速度测试elapsedTime = cuda_malloc_test(SIZE, true);printf("Time using cudaMalloc: %3.1f ms\n", elapsedTime);printf("\tMB/s during copy up: %3.1f\n", MB / (elapsedTime / 1000));
}
流
CUDA 中采用 流（stream） 能带来性能提升的核心原因是：
流允许数据传输（拷贝）和 GPU 计算（kernel） 之间 并行 / 重叠执行，最大化利用 CPU、PCIe 带宽、GPU 核心 等多个资源。
默认情况下（不使用流），操作是串行执行的，比如：
cudaMemcpy(...); // 主机到设备拷贝 (同步，阻塞)
my_kernel<<<...>>>(...); // kernel 启动
cudaMemcpy(...); // 设备到主机拷贝 (同步)
这些操作会：
-
一个结束后才开始下一个（传输和计算不能并发）；
-
阻塞 CPU，无法做其他事情；
-
GPU 计算资源闲置（等数据）；
-
PCIe 传输通道利用率低。
流可以带来哪些 并发重叠？
传输和计算重叠
通过 流 + pinned memory + async copy，可以做到：
-
在
stream1中异步拷贝 A 到 GPU； -
同时
stream2中 GPU 正在处理 B 数据； -
然后把 B 结果拷回主机的同时拷入 C 数据……
这称为 pipeline 式数据处理。
| 时间线 | Stream 1 | Stream 2 | Stream 3 |
|---|---|---|---|
| T1 | HtoD A | ||
| T2 | Kernel A | HtoD B | |
| T3 | DtoH A | Kernel B | HtoD C |
| T4 | DtoH B | Kernel C | |
| T5 | DtoH C |
这样：数据传输 和 GPU 计算可以重叠，整体吞吐量大幅提升。
多核函数并发调度
多个独立 kernel（如多个图像处理任务），如果：
-
都用默认流：会串行排队；
-
用不同的流：调度器会并行启动它们！
前提是设备支持并发执行（Concurrent Kernel Execution）。
什么条件下流才能提升效果？
| 条件 | 是否必须 | 说明 |
|---|---|---|
✔ 使用 cudaMemcpyAsync | ✅ 必须 | 否则不会异步执行 |
| ✔ 主机内存是 pinned memory | ✅ 必须 | 普通 malloc 内存不支持异步 |
| ✔ 操作放在非默认流上 | ✅ 必须 | 默认流（stream 0）会阻塞其他流 |
| ✔ 设备支持并发（多数支持） | ✅ 强烈建议 | 多流多 kernel 才能并行调度 |
| ❌ 没有内存访问依赖关系 | ✅ 最佳 | 否则执行顺序会被限制 |
代码示例
#include <cuda_runtime.h>
#include <device_launch_parameters.h>#include <iostream>
#include<cstdio>
const int N = 1024 * 1024;
#define FULL_DATA_SIZE N*20
void error_handling(cudaError_t res) {if (res !=cudaSuccess) {std::cout << "error!" << std::endl;}
}
//核函数不关键，可以不看
__global__ void kernel(int* a, int* b, int* c) {int idx = threadIdx.x + blockIdx.x * blockDim.x;if (idx < N) {int idx1 = (idx + 1) % 256;int idx2 = (idx + 2) % 256;float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;c[idx] = (as + bs) / 2;}
}
int main() {cudaDeviceProp prop;int whichDevice;error_handling(cudaGetDevice(&whichDevice));error_handling(cudaGetDeviceProperties_v2(&prop, whichDevice));if (!prop.deviceOverlap) {//选择一个支持设备重叠功能的设备printf("Device will not handle overlappps,so no spped up from streams\n");return 0;}cudaEvent_t start, stop;float elappsedTime;//启动计时器error_handling(cudaEventCreate(&start));error_handling(cudaEventCreate(&stop));error_handling(cudaEventRecord(start, 0));//初始化流cudaStream_t stream;error_handling(cudaStreamCreate(&stream));int* host_a, * host_b, * host_c;int* dev_a, * dev_b, * dev_c;//分配设备内存error_handling(cudaMalloc((void**)&dev_a,N*sizeof(int)));error_handling(cudaMalloc((void**)&dev_b, N * sizeof(int)));error_handling(cudaMalloc((void**)&dev_c, N * sizeof(int)));//分配由流使用的页锁定内存error_handling(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));error_handling(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));error_handling(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));for (int i = 0; i < FULL_DATA_SIZE; i++) {host_a[i] = rand();host_b[i] = rand();}//每一批做三件事，都通过同一个 stream 异步调度：for (int i = 0; i < FULL_DATA_SIZE; i += N) {error_handling(cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));error_handling(cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));kernel << <N / 256, 256, 0, stream >> > (dev_a, dev_b, dev_c);error_handling(cudaMemcpyAsync(host_c + i, dev_c, N * sizeof(int), cudaMemcpyHostToDevice, stream));}error_handling(cudaStreamSynchronize(stream));error_handling(cudaEventRecord(stop, 0));error_handling(cudaEventSynchronize(stop));error_handling(cudaEventElapsedTime(&elappsedTime, start, stop));printf("Time taken : %31f ms\n", elappsedTime);error_handling(cudaFreeHost(host_a));error_handling(cudaFreeHost(host_b));error_handling(cudaFreeHost(host_c));error_handling(cudaFree(dev_a));error_handling(cudaFree(dev_b));error_handling(cudaFree(dev_c));error_handling(cudaStreamDestroy(stream));error_handling(cudaEventDestroy(start));error_handling(cudaEventDestroy(stop));
}
多流 pipelining（真正 overlap）
目前所有操作都用一个 stream，所以 一个 batch 的三个操作仍然串行排队。
可以使用多个 stream 轮流调度每个 batch：
cudaStream_t streams[4];
for (int i = 0; i < 4; ++i)cudaStreamCreate(&streams[i]);for (int i = 0; i < FULL_DATA_SIZE; i += N) {int s = (i / N) % 4;cudaMemcpyAsync(..., streams[s]);kernel<<<..., ..., 0, streams[s]>>>(...);cudaMemcpyAsync(..., streams[s]);
}
这样可实现拷贝和 kernel 的真正 overlap。
流+事件计时
例子：两个 kernel 用两个 stream 并发执行，各自测时间
cudaStream_t stream1, stream2;
cudaEvent_t start1, stop1, start2, stop2;
float time1 = 0.0f, time2 = 0.0f;cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);cudaEventCreate(&start1); cudaEventCreate(&stop1);
cudaEventCreate(&start2); cudaEventCreate(&stop2);// kernel1 时间测量（绑定到 stream1）
cudaEventRecord(start1, stream1);
my_kernel<<<gridDim, blockDim, 0, stream1>>>(...);
cudaEventRecord(stop1, stream1);// kernel2 时间测量（绑定到 stream2）
cudaEventRecord(start2, stream2);
my_kernel2<<<gridDim, blockDim, 0, stream2>>>(...);
cudaEventRecord(stop2, stream2);// 等待各自完成
cudaEventSynchronize(stop1);
cudaEventSynchronize(stop2);// 计算时间（单位 ms）
cudaEventElapsedTime(&time1, start1, stop1);
cudaEventElapsedTime(&time2, start2, stop2);printf("Kernel 1 time: %.2f ms\n", time1);
printf("Kernel 2 time: %.2f ms\n", time2);// 清理
cudaEventDestroy(start1); cudaEventDestroy(stop1);
cudaEventDestroy(start2); cudaEventDestroy(stop2);
cudaStreamDestroy(stream1); cudaStreamDestroy(stream2);