Optimize Data Transfers in CUDA C/C++
2020 Jan 27th CookieLau
Source:
[toc]
nvprof
运行 cuda 应用时，可以加上 nvprof 来查看每个 device 的指令的运行情况以找到加速的瓶颈。
$ nvprof ./a.out
----------------------------------------------------
======== NVPROF is profiling a.out...
======== Command: a.out
======== Profiling result:
Time(%) Time Calls Avg Min Max Name
50.08 718.11us 1 718.11us 718.11us 718.11us [CUDA memcpy DtoH]
49.92 715.94us 1 715.94us 715.94us 715.94us [CUDA memcpy HtoD]
nvprof 有很多 option 运行选项，可以通过 nvprof --help 查看，主要的有：
1. --print-gpu-trace
2. --print-api-trace
3. --events 记录指定的事件，如 warps_launched(Number of warps launched)，如有多个用英文逗号分隔，查看全部则接 all；可以用 nvprof --query-events 查看所有可以记录的事件
4. --metrics 记录指定的指标(metrics)，如 ipc(Instructions executed per cycle，如有多个用英文逗号分隔，查看全部则接 all；可以用 nvprof --query-metrics 查看所有可以记录的指标
详细信息可查阅官方文档：
Pinned Host memory （固定内存，锁页内存）
所以如果我们需要在 host 和 device 端来回传输数据时，pageable memory 的host端拷贝所带来的开销是不可忽视的，为了减少这部分的开销，我们可以通过 在host端分配数据时直接分配pinned memory。
具体方法为： 1. 创建：cudaMallocHost or cudaHostAlloc
2. 销毁：cudaFreeHost
由于直接分配 pinned memory 是可能失效的，所以我们需要进行 cudaError 的检查：
cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes);
if (status != cudaSuccess)
printf("Error allocating pinned host memory\n");
Example 十分标准的一个测试传输带宽的程序
#include <stdio.h>
#include <assert.h>
// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
// Cuda 函数调用正确性检查的规范形式
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
if (result != cudaSuccess) {
fprintf(stderr, "CUDA Runtime Error: %s\n",
cudaGetErrorString(result));
assert(result == cudaSuccess);
}
#endif
return result;
}
void profileCopies(float *h_a,
float *h_b,
float *d,
unsigned int n,
char *desc)
{
printf("\n%s transfers\n", desc);
unsigned int bytes = n * sizeof(float);
// events for timing
cudaEvent_t startEvent, stopEvent;
checkCuda( cudaEventCreate(&startEvent) );
checkCuda( cudaEventCreate(&stopEvent) );
checkCuda( cudaEventRecord(startEvent, 0) );
checkCuda( cudaMemcpy(d, h_a, bytes, cudaMemcpyHostToDevice) );
checkCuda( cudaEventRecord(stopEvent, 0) );
checkCuda( cudaEventSynchronize(stopEvent) );
// cudaEventSync 要让CPU在此处等待直到memcpy完成，否则CPU会继续执行下面的代码
float time;
checkCuda( cudaEventElapsedTime(&time, startEvent, stopEvent) );
printf(" Host to Device bandwidth (GB/s): %f\n", bytes * 1e-6 / time);
checkCuda( cudaEventRecord(startEvent, 0) );
checkCuda( cudaMemcpy(h_b, d, bytes, cudaMemcpyDeviceToHost) );
checkCuda( cudaEventRecord(stopEvent, 0) );
checkCuda( cudaEventSynchronize(stopEvent) );
checkCuda( cudaEventElapsedTime(&time, startEvent, stopEvent) );
printf(" Device to Host bandwidth (GB/s): %f\n", bytes * 1e-6 / time);
// 之所以 *1e-6 是因为后面的 time 是 ms，所以转换成为 s 需要 1e3，合起来就是 1e9 = 1G
for (int i = 0; i < n; ++i) {
if (h_a[i] != h_b[i]) {
printf("*** %s transfers failed ***\n", desc);
break;
}
}
// clean up events
checkCuda( cudaEventDestroy(startEvent) );
checkCuda( cudaEventDestroy(stopEvent) );
}
int main()
{
unsigned int nElements = 4*1024*1024;
const unsigned int bytes = nElements * sizeof(float);
// host arrays
float *h_aPageable, *h_bPageable;
float *h_aPinned, *h_bPinned;
// device array
float *d_a;
// allocate and initialize
h_aPageable = (float*)malloc(bytes); // host pageable
h_bPageable = (float*)malloc(bytes); // host pageable
checkCuda( cudaMallocHost((void**)&h_aPinned, bytes) ); // host pinned
checkCuda( cudaMallocHost((void**)&h_bPinned, bytes) ); // host pinned
checkCuda( cudaMalloc((void**)&d_a, bytes) ); // device
for (int i = 0; i < nElements; ++i) h_aPageable[i] = i;
memcpy(h_aPinned, h_aPageable, bytes);
memset(h_bPageable, 0, bytes);
memset(h_bPinned, 0, bytes);
// output device info and transfer size
cudaDeviceProp prop;
checkCuda( cudaGetDeviceProperties(&prop, 0) );
printf("\nDevice: %s\n", prop.name);
printf("Transfer size (MB): %d\n", bytes / (1024 * 1024));
// perform copies and report bandwidth
profileCopies(h_aPageable, h_bPageable, d_a, nElements, "Pageable");
profileCopies(h_aPinned, h_bPinned, d_a, nElements, "Pinned");
printf("n");
// cleanup
cudaFree(d_a);
cudaFreeHost(h_aPinned);
cudaFreeHost(h_bPinned);
free(h_aPageable);
free(h_bPageable);
return 0;
}
原文中作者用了两种 CPU 做测试，在性能较弱的 CPU 上 pageable memory 的带宽是 pinned memory 的一半不到，而在好的 CPU 上测试时两者带宽近似。
注意：
这里不是无脑推荐只要在 host 上分配内存都用 pinned memory，因为大量的 pinned memory 的分配会导致整体性能的下降，因为连续分配固定内存会使得其他应用程序的可以使用的内存减少，导致性能的降低。具体什么时候分配需要经过自己的测试得到。
Batch Small Data Transfer
建议通过每次传输数据多一些来 尽量减少传输的次数。这就很 是具体情况而定了。
host cudaError_t cudaMallocPitch(void* devPtr, size_t pitch, size_t widthInBytes, size_t height)
pitch：分配存储器的宽度，以字节为单位 （通常是 256B or 512B）
cudaMallocPitch 可以保证，虽然可能存在内存空间的浪费，但是访存的时候对于每一行数据的读取是对齐的。
这里为什么要单独返回 Pitch 呢？是因为后面的高阶拷贝要用到，用来计算2D数据中的某一个元素的地址：
对于一个给定类型的数组，Row 行 Column 列 的元素的地址为：
T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;
host cudaError_t cudaMemcpy2D(void dst, size_t dpitch, const void src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)
dpitch: dst指向的2D数组中的内存宽度，以字节为单位，是cuda为了读取方便，对齐过的内存宽度，可能大于一行元素占据的实际内存
spitch: src指向的2D数组中的内存宽度，以字节为单位
width: src指向的2D数组中一行元素占据的实际宽度。以字节为单位，等于 width*sizeof(type)
Example
size_t width = 6;
size_t height = 5;
float *h_data, *d_data;
size_t pitch;
// 这里只分配 host 端的内存
h_data = (float *)malloc(sizeof(float)*width*height);
for (int i = 0; i < width*height; i++)
h_data[i] = (float)i;
// 这里通过 cudaMallocPitch 来分配 device 端的内存，并获得 pitch 值
printf("\n\tAlloc memory.");
cudaMallocPitch((void **)&d_data, &pitch, sizeof(float)*width, height);
printf("\n\tPitch = %d B\n", pitch);
// 将数据从 host 端拷贝到 device 端
// 由于 host 上是我们直接连续分配的，所以这个 spitch 就由我们直接写出来
printf("\n\tCopy to Device.\n");
cudaMemcpy2D(d_data, pitch, h_data, sizeof(float)*width, sizeof(float)*width, height, cudaMemcpyHostToDevice);
myKernel << <1, 1 >> > (d_data, height, width, pitch);
cudaDeviceSynchronize();
// 将数据从 device 端拷贝回 host 端，此时 dpitch 是我们的 host
printf("\n\tCopy back to Host.\n");
cudaMemcpy2D(h_data, sizeof(float)*width, d_data, pitch, sizeof(float)*width, height, cudaMemcpyDeviceToHost);
Reference
