cudaMallocHost和malloc对于提高性能没有区别

Question

我已经阅读了这个网站。从这里我得知，使用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);
}

using 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的执行有什么确切的区别？

并且随着每次运行，执行时间都会减少。

- user3545251

---

如果只有100000个元素，你几乎看不到差异。你可以运行来自http://docs.nvidia.com/cuda/cuda-samples/#bandwidth-test的“带宽测试”CUDA示例。它运行更复杂的测试，这时候差异应该会变得明显。 - Marco13

内存传输可能只占您所测量时间的一小部分。使用更高分辨率的计时器，仅测量cudaMemcpy调用，您可能会注意到差异。或者，使用CUDA工具包中提供的任何一个分析工具适用于您的平台。 - talonmies

如果我使用一个很大的N值，比如N=1610241024，结果与预期不符。它们不是平方根，而只是“i”的值。为什么会出现这种异常行为？我还尝试过使用malloc()而不是cudaHostMalloc()。 - MuneshSingh

Answer 1

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

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

因此，在许多示例中，仅计时整个执行过程并不显示出很大的差异，因为尽管从固定内存进行cudaMemcpy操作可能更快，但与相应的malloc相比，cudaMallocHost需要更长的时间。

那么这有什么意义呢？

1. 如果您需要对单个缓冲区进行重复传输，您可能会对使用固定内存（即cudaMallocHost）感兴趣。您只需支付一次固定的额外成本，但在每次传输/使用时都会受益。
2. 必须使用固定内存来重叠数据传输操作（cudaMemcpyAsync）和计算活动（内核调用）。请参阅编程指南。

- Robert Crovella

Answer 2

我也发现，仅仅在内存上声明cudaHostAlloc / cudaMallocHost并不会有太大的作用。 为确保，可以使用--print-gpu-trace进行nvprof，查看memcpyHtoD或memcpyDtoH的吞吐量是否良好。 对于PCI2.0，应该获得约6-8gbps。

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

我很惊讶，我以这种方式能够节省相当多的时间，值得一试。

- lppier