CUDA统一内存

简介

• 简介
  统一内存使得分配和访问系统中任何处理器上运行的代码都可以使用的数据遍历非常容易，CPU或GPU。

• 使用
  分配统一内存非常简单，只须将malloc，new的调用替换为对**cudaMallocManaged()**的调用，这是一个分配函数，返回可从任何处理器的访问的指针。
  cudaError_t cudaMallocManaged(void **ptr, size_t size);
  当在CPU或GPU上运行的代码访问这种分配的数据（通常称为CUDA管理数据），CUDA系统软件和硬件负责将MIG额定内存也分配给访问处理的内存。重要：Pascal GPU体系结构是第一个通过页面MIG比率引擎对虚拟内存页错误处理和MIG比率提供硬件支持的架构。基于更旧的kezbr架构和更为统一的kezbr形式的支持。

• 举例

  #include <iostream>
  #include <math.h>
  
  // CUDA kernel to add elements of two arrays
  __global__ void add(int n, float *x, float *y) {
      int index = blockIdx.x * blockDim.x + threadIdx.x;
      int stride = blockDim.x * gridDim.x;
      for (int i = index; i < n; i += stride)
          y[i] = x[i] + y[i];
  }
  
  int main(void) {
      int N = 1 << 20;
      float *x, *y;
      // Allocate Unified Memory -- accessible from CPU or GPU
      cudaMallocManaged(&x, N * sizeof(float));
      cudaMallocManaged(&y, N * sizeof(float));
      // initial x and y arrays on the host
      for (int i = 0; i < N; i++) {
          x[i] = 1.0f;
          y[i] = 2.0f;
      }
      // Launch kernel on 1M element on the GPU
      int blockSize = 256;
      int numBlocks = (N + blockSize - 1) / blockSize;
      add<<<numBlocks, blockSize>>>(N, x, y);
      // Wait for GPU to finish before accessing on host
      cudaDeviceSynchronize();
      // Check for error (all values should be 3.0f)
      float maxError = 0.0f;
      for (int i = 0; i < N; i++)
          maxError = fmax(maxError, fabs(y[i] - 3.0f));
      std::cout << "Max error: " << maxError << std::endl;
      // Free memory
      cudaFree(x);
      cudaFree(y);
      return 0;
  }
  

  编译并执行程序，结果如下：
  nvcc test_um.cu
  ./a.out
  对应输出为：Max error: 0

• 参考文献
  《CUDA C编程权威指南》