Accelerated Ray Tracing （五）

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中)。