https://developer.nvidia.com/blog/how-optimize-data-transfers-cuda-cc/
优化数据传输:
- 减少host 和device 之间数据传输。
- 更高的带宽。
- 将许多小传输批量处理为一个较大的传输性能更好,因为它消除了每次传输的大部分开销。
- 主机和设备之间的数据传输有时会与内核执行和其他数据传输重叠。
int main()
{
const unsigned int N = 1048576;
const unsigned int bytes = N * sizeof(int);
int* h_a = (int*)malloc(bytes);
int* d_a;
cudaMalloc((int**)&d_a, bytes);
memset(h_a, 0, bytes);
cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(h_a, d_a, bytes, cudaMemcpyDeviceToHost);
return 0;
}
$ nvcc profile.cu -o profile_test
$ nvprof ./profile_test
nvcc编译nvprof运行测试
$ 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
在每一次的memcpy都会测量一次,返回平均值、最小值、最大值。
Minimizing Data Transfers
不仅要考虑GPU和CPU执行要花的时间,还需要考虑在PCI-e总线上传输数据要花的时间,特别是在我们最初将代码移植到CUDA时。由于CUDA的异质编程模型同时使用CPU和GPU,所以代码可以一次移植到一个内核中。在移植的最初阶段,数据传输可能会占据整个执行时间。将花在数据传输上的时间与花在内核执行上的时间分开是值得的。
Pinned Host Memory
主机(CPU)数据分配在默认情况下是可分页的。GPU无法直接从可分页的主机内存访问数据,因此,当调用从可分页的主机内存到设备内存的数据传输时,CUDA驱动程序必须首先分配一个临时的page-lock或“pinned”的host array,将host data复制到pinned array,然后将数据从 pinned array 传输到 device memory,如下所示。
如图所示,pinned memory用作从device 到host传输的暂存区。通过直接在pinned memory中分配host arrays,我们可以避免在可分页pageable 和pinned host arrays 之间进行传输的开销。使用cudaMallocHost()或cudaHostAlloc()在CUDA C / C ++中分配pinned host memory ,并使用cudaFreeHost()取消分配。pinned memory 分配可能会失败,因此您应一直检查错误。
#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.
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) );
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);
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;
}
Device: NVS 4200M
Transfer size (MB): 16
Pageable transfers
Host to Device bandwidth (GB/s): 2.308439
Device to Host bandwidth (GB/s): 2.316220
Pinned transfers
Host to Device bandwidth (GB/s): 5.774224
Device to Host bandwidth (GB/s): 5.958834
例子中pinned 传输的速度是pageable 传输速度的两倍以上。
Device: GeForce GTX 680
Transfer size (MB): 16
Pageable transfers
Host to Device bandwidth (GB/s): 5.368503
Device to Host bandwidth (GB/s): 5.627219
Pinned transfers
Host to Device bandwidth (GB/s): 6.186581
Device to Host bandwidth (GB/s): 6.670246
CPU更好,pageable 速度更快。
您不应过度分配固定的内存。这样做会降低整体系统性能,因为它会减少操作系统和其他程序可用的物理内存量。多少是很难事先告知的,因此,对于所有优化,请测试您的应用程序和它们在其上运行的系统以获得最佳性能参数。
Batching Small Transfers
由于与每个传输相关的开销,因此最好将许多小传输分批处理成一个传输。这可以通过使用temporary array(最好是pinned的)并将其与要传输的数据打包在一起来轻松实现。
对于二维数组传输,可以使用cudaMemcpy2D()。
cudaMemcpy2D(dest, dest_pitch, src, src_pitch, w, h, cudaMemcpyHostToDevice)
这里的参数是指向第一个目标元素和目标数组的间距的指针,一个指向第一个源元素和源数组的间距,要传递的子矩阵的宽度和高度以及memcpy类型的指针。还有一个cudaMemcpy3D()函数用于传输三级数组节。
建议您不要使用带有CUDA事件或其他计时器的工具代码来测量每次传输所花费的时间,我建议您使用nvprof、命令行CUDA剖析器,或者可视化剖析工具,比如NVIDIA visual profiling tools(也包含在CUDA Toolkit中)。