【Zero-Copy】一些关于 零拷贝 的问题

什么是零拷贝

从CUDA 9.x和NVIDIA Xavier架构开始，引入了通过硬件保证的CPU和GPU之间的缓存一致性机制。这种机制允许CPU和GPU更高效地共享物理内存，通过使用所谓的“零拷贝”内存编程模型来实现。这里的“零拷贝”并不意味着完全不需要手动分配CPU或GPU内存，而是指在CPU和GPU之间共享数据时，可以减少数据的复制操作，即不需要将数据从CPU内存拷贝到GPU内存（或反向），从而减少了延迟和带宽消耗。

零拷贝技术通过避免在内存中多次复制数据来解决这一问题，实现了数据传输的直接访问和操作。具体来说，零拷贝技术允许网络数据包或文件数据直接从输入/输出设备传输到最终的内存位置，或者使得CPU和GPU能够共享相同的物理内存地址空间，这样数据就可以不经过中间复制步骤直接被处理。

零拷贝和共享内存，统一内存的区别

零拷贝内存（Zero-Copy Memory）和共享内存（Shared Memory）都是提高程序性能、优化内存使用的技术，尤其在高性能计算和并行计算领域中被广泛使用。尽管它们的目标相似，即减少不必要的数据拷贝以提升效率，但它们的应用场景、实现方式和优化点有所不同。

零拷贝内存

零拷贝技术主要用于减少或消除CPU和I/O之间数据传输时的冗余拷贝操作，以减少延迟和降低CPU负载。在GPU计算场景中，零拷贝内存通常指的是允许GPU直接访问CPU的物理内存，从而避免了将数据从主机内存拷贝到GPU内存的需要。

• 应用场景：适用于数据频繁在CPU和GPU之间传输的情况，特别是对于小批量数据处理。
• 优点：减少数据传输延迟，降低CPU和内存带宽消耗。
• 局限性：直接访问可能不经过GPU的高速缓存，对于大规模数据处理可能不是最高效的方式。

共享内存（在GPU计算中）

共享内存是GPU上的一种快速内存，位于每个线程块中，可以被该线程块内的所有线程访问。它用于线程间的数据共享和减少对全局内存的访问次数。

• 应用场景：适用于同一线程块内的线程需要频繁交换数据或重复访问某些数据的计算任务。
• 优点：提供比全局内存更高的访问速度，有助于减少内存访问延迟，提升并行计算性能。
• 局限性：容量有限，需要精心设计数据访问模式以避免银行冲突。

统一内存（Unified Memory）

统一内存（Unified Memory）是CUDA编程模型中的一个高级特性，旨在简化GPU编程中内存管理的复杂性。统一内存为CPU和GPU提供了一个共享的内存视图，使得开发者无需手动管理数据在CPU和GPU之间的传输。与零拷贝内存和共享内存相比，统一内存提供了不同的优势和使用场景。

区别总结

• 统一内存提供了一个简化的编程模型，通过自动数据迁移来实现CPU和GPU之间的内存透明共享，减轻了开发者在内存管理上的负担，但可能会引入额外的性能开销。
• 共享内存是GPU内部线程块中的快速内存，用于加速线程块内的数据交换和减少全局内存访问，需要开发者精心设计数据访问模式。
• 零拷贝内存专注于减少CPU和GPU之间的数据拷贝操作，适用于小批量数据处理，需要开发者手动管理内存。

什么时候考虑用零拷贝

1. 数据量小且频繁的数据交换

当应用涉及到频繁的小数据量交换时，使用零拷贝能够减少内存复制的开销，因为直接在原地操作数据避免了额外的内存分配和释放过程。

2. 数据由多个GPU核心小批量处理

如果数据被多个GPU核心以小批量方式处理，使用零拷贝能够避免重复拷贝同一数据到GPU内存中的情况，每个核心可以直接访问CPU内存中的数据。

3. 读多写少的操作

对于主要是读操作的应用，使用零拷贝避免了数据的复制，因为数据大部分时间是被读取而不是被修改。当数据修改操作较少时，同步更新共享内存中的数据不会造成太大的开销。

4. 系统内存充足的情况

在系统内存资源充足的情况下，将部分内存用作CPU和GPU的共享内存，可以有效利用资源，减少GPU内存的压力。

5. 避免PCIe带宽瓶颈

在数据需要通过PCIe总线在CPU和GPU之间传输时，使用零拷贝可以减少PCIe带宽的占用，特别是在PCIe带宽是性能瓶颈的情况下。

注意事项：

虽然零拷贝技术在特定场景下可以提高性能，但也有其局限性。例如，当数据量非常大时，直接在共享内存上操作可能会因为缓存一致性同步的额外开销而降低性能。此外，对共享内存的直接访问可能不会利用GPU的高速缓存，可能导致较低的内存访问效率。因此，用的时候要根据应用的具体需求来决定是否使用零拷贝技术。

代码示例

零拷贝

// 设置标志以启用零拷贝访问
cudaSetDeviceFlags(cudaDeviceMapHost);

// 主机数组（CPU指针）
float* h_in  = NULL;
float* h_out = NULL;

// 处理 h_in 输入数据

// 使用CUDA分配调用分配主机内存
cudaHostAlloc((void **)&h_in,  sizeIn,  cudaHostAllocMapped);
cudaHostAlloc((void **)&h_out, sizeOut, cudaHostAllocMapped);

// 设备数组（GPU指针，这里是通过映射获得的，看起来像是CPU指针）
float *d_out, *d_in;
// 从主机内存获取设备指针。不需要额外的内存分配或内存拷贝操作
cudaHostGetDevicePointer((void **)&d_in,  (void *) h_in , 0);
cudaHostGetDevicePointer((void **)&d_out, (void *) h_out, 0);

// 启动GPU内核
kernel<<<blocks, threads>>>(d_out, d_in);

// 不需要将d_out的内容拷贝回主机
// 可以继续使用h_out在主机上进行处理

统一内存

#include <iostream>
#include <math.h>

// CUDA内核函数，用于将两个数组的元素相加
__global__
void add(int n, float *x, float *y)
{
  // 计算当前线程的全局索引
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  // 计算总的线程跨度
  int stride = blockDim.x * gridDim.x;
  // 以stride为步长进行数组加法，以确保所有元素都被处理
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 定义元素个数为2的20次方
  float *x, *y;

  // 分配统一内存 -- 这些指针可以从CPU或GPU访问
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // 在主机上（CPU）初始化x和y数组
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // 在GPU上启动内核，处理1M个元素
  int blockSize = 256; // 每个块的线程数
  int numBlocks = (N + blockSize - 1) / blockSize; // 计算需要的块数
  add<<<numBlocks, blockSize>>>(N, x, y);

  // 等待GPU完成处理，再在主机上访问数据
  cudaDeviceSynchronize();

  // 检查错误（所有值应该是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;

  // 释放内存
  cudaFree(x);
  cudaFree(y);

  return 0;
}