# 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) --- # Agent Instructions: Querying This Documentation If you need additional information that is not directly available in this page, you can query the documentation dynamically by asking a question. Perform an HTTP GET request on the current page URL with the `ask` query parameter: ``` GET https://legacy.cookielau.com/archives/10-hardware/0-gpu/010-optimize_data_transfers_in_cuda.md?ask= ``` The question should be specific, self-contained, and written in natural language. The response will contain a direct answer to the question and relevant excerpts and sources from the documentation. Use this mechanism when the answer is not explicitly present in the current page, you need clarification or additional context, or you want to retrieve related documentation sections.