Optimize_Data_Transfers_in_CUDA

Optimize Data Transfers in CUDA C/C++

2020 Jan 27th CookieLau

Source: https://devblogs.nvidia.com/how-optimize-data-transfers-cuda-cc/

[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 查看所有可以记录的指标

详细信息可查阅官方文档：Link

Pinned Host memory （固定内存，锁页内存）

在 host 上分配的数据内存可能是可以换页的(pageable)，这就有可能会导致数据的缺页或者是脏页，所以当 cuda 需要进行 data transfer 的时候，会进行一个先在 host 上分配 pinned memory，将 pageable memory 上的数据拷贝到 pinned memory 上，再从 pinned memory 拷贝数据到 device 端的操作，如下图所示：

所以如果我们需要在 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

建议通过每次传输数据多一些来 尽量减少传输的次数。这就很 是具体情况而定了。

对于 高阶数组而言，为了达到更好的访存效果，我们有专门的内存分配和拷贝函数（这里只讨论到二维矩阵，因为三维及以上的矩阵拷贝时需要自定义额外的大量参数，而且使用也不是非常频繁，不在此讨论，有兴趣可以看 Link）：

• 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 列 的元素的地址为：

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

    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

2. 博客园 | 二维数组 cudaMallocPitch() 和三维数组 cudaMalloc3D() 的使用

3. 相关函数的cuda文档位置