CUDA编程：cudaMalloc、cudaHostAlloc和cudaMallocManaged三种方式所创建内存在主机和GPU间传输速度的差异
在上一篇文章《CUDA编程： GPU与CPU之间使用统一内存的完整示例代码》里，为了在主机开辟巨大的内存，只读开放给GPU访问，测试了cudaMallocManaged的用法，并提出了一个疑问，到底CUDA里最传统的内存拷贝cudaMalloc，和显式地在主机创建内存cudaHostAlloc再开放给GPU访问，以及全局内存寻址cudaMallocManaged，这三种方式，哪种在GPU与CPU之间传输数据时，有更高的效率？
为了回答这个问题，于是有了以下对上述三种方式所创建的内存，在主机与设备之间传输速度的比较代码。
先贴测试的结果，显然cudaHostAlloc方式创始的主机内存，共享给设备读取时，效率最高：
$ ./test_mem
[+] 比较cudaMalloc、cudaHostAlloc和cudaMallocManaged三种方式创建的内存数据,在主机和GPU间传输的性能差异....
[1] 使用 cudaMalloc 创建[128.00]MB内存空间并从主机传输到GPU 100次耗时: 802.2 ms
传输速度: 15956.9 (MB/s)
[2] 使用 cudaMalloc 创建[128.00]MB内存空间并从GPU传输到主机100次耗时: 1486.6 ms
传输速度: 8610.4 (MB/s)
[3] 使用 cudaHostAlloc 创建[128.00]MB内存空间并从主机传输到GPU 100次耗时: 515.6 ms
传输速度: 24827.8 (MB/s)
[4] 使用 cudaHostAlloc 创建[128.00]MB内存空间并从GPU传输到主机100次耗时: 509.2 ms
传输速度: 25136.0 (MB/s)
[5] 使用 cudaMallocManaged 创建[128.00]MB内存空间并从主机传输到GPU 100次耗时: 1161.2 ms
传输速度: 11022.7 (MB/s)
[6] 使用 cudaMallocManaged 创建[128.00]MB内存空间并从GPU传输到主机100次耗时: 1064.9 ms
传输速度: 12019.5 (MB/s)
结论是显著的，cudaHostAlloc有着最高的性能，速度是另外两方式的二至三倍。
以下是测试程序 cuda_mem_cpy.cu 完整的源代码：
/*******************************************************************************************************************
* 文件名 cuda_mem_cpy.cu
* 编译： nvcc -o test_uvm uvm_test.cu
* 功能：比较cudaMalloc、cudaHostAlloc和cudaMallocManaged三种方式创建的内存数据,在主机和GPU间传输的性能差异
* author: Ryan
********************************************************************************************************************/
#include <iostream>
#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <device_launch_parameters.h>
#define SIZE (32*1024*1024)
using namespace std;
//------------------------------------------------------------------------------------------------------------------
void HANDLE_ERROR(cudaError_t cuda_error_code){
if(cuda_error_code != cudaSuccess)
printf("[E] CUDA返回错误: %s\n", cudaGetErrorString(cuda_error_code));
}
//------------------------------------------------------------------------------------------------------------------
// 测试 cudaMalloc 方法
float cuda_malloc_test(int size, bool up) {
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
HANDLE_ERROR(cudaEventCreate(&start));
HANDLE_ERROR(cudaEventCreate(&stop));
a = (int*)malloc(size * sizeof(*a));
HANDLE_ERROR(cudaMalloc((void**)&dev_a,
size * sizeof(*dev_a)));
HANDLE_ERROR(cudaEventRecord(start, 0));
for (int i = 0; i<100; i++) {
if (up)
HANDLE_ERROR(cudaMemcpy(dev_a, a,
size * sizeof(*dev_a),
cudaMemcpyHostToDevice));
else
HANDLE_ERROR(cudaMemcpy(a, dev_a,
size * sizeof(*dev_a),
cudaMemcpyDeviceToHost));
}
HANDLE_ERROR(cudaEventRecord(stop, 0));
HANDLE_ERROR(cudaEventSynchronize(stop));
HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
start, stop));
free(a);
HANDLE_ERROR(cudaFree(dev_a));
HANDLE_ERROR(cudaEventDestroy(start));
HANDLE_ERROR(cudaEventDestroy(stop));
return elapsedTime;
}
//------------------------------------------------------------------------------------------------------------------
// 测试 cudaHostAlloc 方法
float cuda_host_alloc_test(int size, bool up) {
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
HANDLE_ERROR(cudaEventCreate(&start));
HANDLE_ERROR(cudaEventCreate(&stop));
HANDLE_ERROR(cudaHostAlloc((void**)&a,
size * sizeof(*a),
cudaHostAllocDefault));
HANDLE_ERROR(cudaMalloc((void**)&dev_a,
size * sizeof(*dev_a)));
HANDLE_ERROR(cudaEventRecord(start, 0));
for (int i = 0; i<100; i++) {
if (up)
HANDLE_ERROR(cudaMemcpy(dev_a, a,
size * sizeof(*a),
cudaMemcpyHostToDevice));
else
HANDLE_ERROR(cudaMemcpy(a, dev_a,
size * sizeof(*a),
cudaMemcpyDeviceToHost));
}
HANDLE_ERROR(cudaEventRecord(stop, 0));
HANDLE_ERROR(cudaEventSynchronize(stop));
HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
start, stop));
HANDLE_ERROR(cudaFreeHost(a));
HANDLE_ERROR(cudaFree(dev_a));
HANDLE_ERROR(cudaEventDestroy(start));
HANDLE_ERROR(cudaEventDestroy(stop));
return elapsedTime;
}
//------------------------------------------------------------------------------------------------------------------
// 测试 cudaMallocManaged 方法
float cuda_alloc_managed_test(int size, bool up) {
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
int deviceId;
HANDLE_ERROR(cudaGetDevice(&deviceId));
HANDLE_ERROR(cudaEventCreate(&start));
HANDLE_ERROR(cudaEventCreate(&stop));
//用cudaMallocManaged开辟内存
HANDLE_ERROR(cudaMallocManaged((void**)&a, size * sizeof(*a)));
if(up)
cudaMemAdvise(a, size * sizeof(*a), cudaMemAdviseSetReadMostly, deviceId);
//用cudaMalloc在设备端开辟内存
HANDLE_ERROR(cudaMallocManaged((void**)&dev_a, size * sizeof(*dev_a)));
//把a同步给设备可见
cudaMemPrefetchAsync(a, size, deviceId);
cudaMemPrefetchAsync(dev_a, size, deviceId);
HANDLE_ERROR(cudaEventRecord(start, 0));
for (int i = 0; i<100; i++) {
if (up){
//HANDLE_ERROR(cudaMemcpy(dev_a, a,size * sizeof(*a),cudaMemcpyHostToDevice));
memcpy(dev_a,a,size * sizeof(*a));
}
else{
//HANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*a),cudaMemcpyDeviceToHost));
memcpy(a,dev_a,size * sizeof(*a));
}
}
HANDLE_ERROR(cudaEventRecord(stop, 0));
HANDLE_ERROR(cudaEventSynchronize(stop));
HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
start, stop));
HANDLE_ERROR(cudaFree(a));
HANDLE_ERROR(cudaFree(dev_a));
HANDLE_ERROR(cudaEventDestroy(start));
HANDLE_ERROR(cudaEventDestroy(stop));
return elapsedTime;
}
//------------------------------------------------------------------------------------------------------------------
int main(void) {
float elapsedTime;
float MB = (float)100 * SIZE*sizeof(int) / 1024 / 1024;
int test_index = 1;
cout << "[+] 比较cudaMalloc、cudaHostAlloc和cudaMallocManaged三种方式创建的内";
cout << "存数据,在主机和GPU间传输的性能差异...." << endl;
// 测试 cudaMalloc
elapsedTime = cuda_malloc_test(SIZE, true);
printf("[%d] 使用 cudaMalloc 创建[%.2f]MB内存空间并从主机传输到GPU 100次耗时: %3.1f ms\n",
test_index,MB/100,elapsedTime);
printf("\t传输速度: %3.1f (MB/s)\n",
MB / (elapsedTime / 1000));
test_index++;
elapsedTime = cuda_malloc_test(SIZE, false);
printf("[%d] 使用 cudaMalloc 创建[%.2f]MB内存空间并从GPU传输到主机100次耗时: %3.1f ms\n",
test_index,MB/100,elapsedTime);
printf("\t传输速度: %3.1f (MB/s)\n",
MB / (elapsedTime / 1000));
test_index++;
// 测试 cudaHostAlloc
elapsedTime = cuda_host_alloc_test(SIZE, true);
printf("[%d] 使用 cudaHostAlloc 创建[%.2f]MB内存空间并从主机传输到GPU 100次耗时: %3.1f ms\n",
test_index,MB/100,elapsedTime);
printf("\t传输速度: %3.1f (MB/s)\n",
MB / (elapsedTime / 1000));
test_index++;
elapsedTime = cuda_host_alloc_test(SIZE, false);
printf("[%d] 使用 cudaHostAlloc 创建[%.2f]MB内存空间并从GPU传输到主机100次耗时: %3.1f ms\n",
test_index,MB/100,elapsedTime);
printf("\t传输速度: %3.1f (MB/s)\n",
MB / (elapsedTime / 1000));
test_index++;
// 测试 cudaAllocManaged
elapsedTime = cuda_alloc_managed_test(SIZE, true);
printf("[%d] 使用 cudaMallocManaged 创建[%.2f]MB内存空间并从主机传输到GPU 100次耗时: %3.1f ms\n",
test_index,MB/100,elapsedTime);
printf("\t传输速度: %3.1f (MB/s)\n",
MB / (elapsedTime / 1000));
test_index++;
elapsedTime = cuda_alloc_managed_test(SIZE, false);
printf("[%d] 使用 cudaMallocManaged 创建[%.2f]MB内存空间并从GPU传输到主机100次耗时: %3.1f ms\n",
test_index,MB/100,elapsedTime);
printf("\t传输速度: %3.1f (MB/s)\n",
MB / (elapsedTime / 1000));
}
//------------------------------------------------------------------------------------------------------------------
可以使用以下命令快速把上述源代码编译为可执行文件test_mem :
$ nvcc ./cuda_memcpy_test.cu -o ./test_mem
（全文完）
版权声明：本文为ysh原创文章，遵循CC 4.0 BY-SA版权协议，转载请附上原文出处链接和本声明。