cudaMallocHost vs malloc 以获得更好的性能显示没有区别

Question

我已经通过了这个site。从这里我得到了使用 cudamallocHost 的固定内存，它提供了比 cudamalloc 更好的性能。然后我使用两个不同的简单程序并测试了执行时间

使用 cudaMallocHost

#include <stdio.h>
#include <cuda.h>

// Kernel that executes on the CUDA device
__global__ void square_array(float *a, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx<N) a[idx] = a[idx] * a[idx];
}

// main routine that executes on the host
int main(void)
{
    clock_t start;
    start=clock();/* Line 8 */
    clock_t finish;
  float *a_h, *a_d;  // Pointer to host & device arrays
  const int N = 100000;  // Number of elements in arrays
  size_t size = N * sizeof(float);
  cudaMallocHost((void **) &a_h, size);
  //a_h = (float *)malloc(size);        // Allocate array on host
  cudaMalloc((void **) &a_d, size);   // Allocate array on device
  // Initialize host array and copy it to CUDA device
  for (int i=0; i<N; i++) a_h[i] = (float)i;
  cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
  // Do calculation on device:
  int block_size = 4;
  int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
  square_array <<< n_blocks, block_size >>> (a_d, N);
  // Retrieve result from device and store it in host array
  cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
  // Print results
  for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
  // Cleanup
  cudaFreeHost(a_h);
  cudaFree(a_d);
  finish = clock() - start;
      double interval = finish / (double)CLOCKS_PER_SEC; 
      printf("%f seconds elapsed", interval);
}

使用 malloc

#include <stdio.h>
#include <cuda.h>

// Kernel that executes on the CUDA device
__global__ void square_array(float *a, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx<N) a[idx] = a[idx] * a[idx];
}

// main routine that executes on the host
int main(void)
{
    clock_t start;
    start=clock();/* Line 8 */
    clock_t finish;
  float *a_h, *a_d;  // Pointer to host & device arrays
  const int N = 100000;  // Number of elements in arrays
  size_t size = N * sizeof(float);
  a_h = (float *)malloc(size);        // Allocate array on host
  cudaMalloc((void **) &a_d, size);   // Allocate array on device
  // Initialize host array and copy it to CUDA device
  for (int i=0; i<N; i++) a_h[i] = (float)i;
  cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
  // Do calculation on device:
  int block_size = 4;
  int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
  square_array <<< n_blocks, block_size >>> (a_d, N);
  // Retrieve result from device and store it in host array
  cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
  // Print results
  for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
  // Cleanup
  free(a_h); cudaFree(a_d);
  finish = clock() - start;
      double interval = finish / (double)CLOCKS_PER_SEC; 
      printf("%f seconds elapsed", interval);
}

在这两个程序的执行过程中，执行时间几乎相似。 执行中有什么问题吗？？ cudamalloc 和 cudamallochost 在执行上的确切区别是什么？？

而且每次运行都会减少执行时间

Answer 1

我也发现仅仅在一块内存上声明 cudaHostAlloc / cudaMallocHost 并没有多大作用。 可以肯定的是，使用 --print-gpu-trace 执行 nvprof 并查看 memcpyHtoD 或 memcpyDtoH 的吞吐量是否良好。对于 PCI2.0，您应该获得大约 6-8gbps。

但是，固定内存是 cudaMemcpyAsync 的必备条件。 在调用 cudaMemcpyAsync 之后，我立即转移了主机上的所有计算。通过这种方式，您可以将异步 memcpys 与主机计算“分层”。

我很惊讶我能够通过这种方式节省大量时间，值得一试。

Answer 2

如果您想查看复制操作的执行时间差异，只需对复制操作进行计时即可。在许多情况下，当底层内存被固定时，仅复制操作的执行时间大约会相差 2 倍。并使您的复制操作足够大/足够长，以使您远高于所使用的任何计时机制的粒度。各种profilers，例如视觉分析器和nvprof，都可以在这里提供帮助。

底层的cudaMallocHost 操作类似于malloc 加上额外的操作系统功能来“固定”与分配相关的每个页面。与仅执行malloc 相比，这些额外的操作系统操作需要额外的时间。请注意，随着分配大小的增加，注册（“固定”）成本通常也会增加。

因此，对于许多示例，仅对整体执行进行计时并没有太大差异，因为虽然cudaMemcpy 操作可能从固定内存中更快，但cudaMallocHost 比相应的malloc 花费更长的时间。

那么有什么意义呢？

1. 当您将从单个缓冲区进行重复传输时，您可能对使用固定内存（即cudaMallocHost）感兴趣。您只需支付固定一次的额外费用，但每次转移/使用都会受益。
2. 需要固定内存才能将数据传输操作 (cudaMemcpyAsync) 与计算活动（内核调用）重叠。请参阅programming guide。