> For the complete documentation index, see [llms.txt](https://legacy.cookielau.com/llms.txt). Markdown versions of documentation pages are available by appending `.md` to page URLs; this page is available as [Markdown](https://legacy.cookielau.com/archives/10-hardware/0-gpu/010-optimize_data_transfers_in_cuda.md). # Optimize\_Data\_Transfers\_in\_CUDA Optimize Data Transfers in CUDA C/C++ 2020 Jan 27th CookieLau Source: \[toc] ## nvprof 运行 cuda 应用时，可以加上 nvprof 来查看每个 device 的指令的运行情况以找到加速的瓶颈。 ```bash $ 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` 查看所有可以记录的指标 详细信息可查阅官方文档：[Link](https://docs.nvidia.com/cuda/profiler-users-guide/index.html#nvprof-overview) ## Pinned Host memory （固定内存，锁页内存） 在 host 上分配的数据内存可能是可以换页的(pageable)，这就有可能会导致数据的缺页或者是脏页，所以当 cuda 需要进行 data transfer 的时候，会进行一个先在 host 上分配 pinned memory，将 pageable memory 上的数据拷贝到 pinned memory 上，再从 pinned memory 拷贝数据到 device 端的操作，如下图所示：\ ![](/files/-M78qWgthpY6tYUfiFG4) 所以如果我们需要在 host 和 device 端来回传输数据时，pageable memory 的host端拷贝所带来的开销是不可忽视的，为了减少这部分的开销，我们可以通过 **在host端分配数据时直接分配pinned memory**。 具体方法为： 1. 创建：cudaMallocHost or cudaHostAlloc\ 2\. 销毁：cudaFreeHost 由于直接分配 pinned memory 是可能失效的，所以我们需要进行 cudaError 的检查： ```c cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes); if (status != cudaSuccess) printf("Error allocating pinned host memory\n"); ``` ### Example 十分标准的一个测试传输带宽的程序 ```c #include #include // 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 建议通过每次传输数据多一些来 **尽量减少传输的次数**。这就很 **是具体情况而定了**。 对于 高阶数组而言，为了达到更好的访存效果，我们有专门的内存分配和拷贝函数（这里只讨论到二维矩阵，因为三维及以上的矩阵拷贝时需要自定义额外的大量参数，而且使用也不是非常频繁，不在此讨论，有兴趣可以看 [Link](https://www.cnblogs.com/cuancuancuanhao/p/7805892.html)）： * **host**​ cudaError\_t cudaMallocPitch(void*\* devPtr, size\_t* pitch, size\_t widthInBytes, size\_t height) * devPtr：新开辟的矩阵的头指针 * pitch：分配存储器的宽度，以字节为单位 （通常是 256B or 512B） * width：分配矩阵的列数 * height：分配矩阵的行数 cudaMallocPitch 可以保证，虽然可能存在内存空间的浪费，但是访存的时候对于每一行数据的读取是对齐的。 这里为什么要单独返回 Pitch 呢？是因为后面的高阶拷贝要用到，用来计算2D数据中的某一个元素的地址：\ 对于一个给定类型的数组，Row 行 Column 列 的元素的地址为： ```c 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) * dst: 目的矩阵内存头指针 * dpitch: dst指向的2D数组中的内存宽度，以字节为单位，是cuda为了读取方便，对齐过的内存宽度，可能大于一行元素占据的实际内存 * src：源矩阵内存头指针 * spitch: src指向的2D数组中的内存宽度，以字节为单位 * width: src指向的2D数组中一行元素占据的实际宽度。以字节为单位，等于 width\*sizeof(type) * height: src指向的2D数组的行数 ### Example ```c 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 1. [程序园 | cudaMallocPitch 和 cudaMemcpy2D](http://www.voidcn.com/article/p-tgqeuirj-bgk.html) 2. [博客园 | 二维数组 cudaMallocPitch() 和三维数组 cudaMalloc3D() 的使用](https://www.cnblogs.com/cuancuancuanhao/p/7805892.html) 3. [相关函数的cuda文档位置](https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g32bd7a39135594788a542ae72217775c)