CUDA编程之内存零拷贝技术

一、实现原理

        零拷贝内存通过将‌主机锁页内存‌直接映射到设备地址空间，实现CPU与GPU共享内存，避免显式数据拷贝‌。锁页内存通过cudaHostAlloc或cudaHostRegister分配，确保物理地址固定且不被操作系统换页，从而支持DMA（直接内存访问）‌。

二、核心函数

• ‌内存分配‌

  cudaHostAlloc(&host_ptr, size, cudaHostAllocMapped | cudaHostAllocWriteCombined); // 分配锁页内存‌
  cudaHostGetDevicePointer(&device_ptr, host_ptr, 0); // 获取设备端指针‌

• ‌内存释放‌：‌

  cudaFreeHost(host_ptr); // 必须使用CUDA接口释放‌

三、适用场景

       零拷贝和普通的cudaMemcpy一样也是要走pci-e总线的，只不过cudaMemcpy是一次性全部copy过去，而零拷贝是用的时候自动在后台通过pci-e总线传输。一般建议只使用一次的数据以及少量的返回数据可以使用零拷贝，其他情况建议copy到显存使用，显存DRAM的带宽要比pci-e的带宽高出一个量级。

• ‌小规模数据频繁访问‌：
  如实时图像处理场景，GPU可直接读取主机内存中的输入数据，减少显存占用‌。
• ‌主机-设备交互频繁‌：
  需多次读写同一内存区域时，零拷贝可替代cudaMemcpy，减少通信开销‌。

四、代码

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

#define N 1024*1024

__global__ void doubleData(float *data, int size) {

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if(i<size){
     data[i]=10;
    }
}

int main()
{
    // 零拷贝内存示例：GPU直接操作主机内存  
    float *host_data, *device_data;  
    cudaHostAlloc(&host_data, N*sizeof(float), cudaHostAllocMapped); // 分配锁页内存‌
    cudaHostGetDevicePointer(&device_data, host_data, 0); // 获取设备指针‌
    int block = 1024;
    int grid = (N+1023)/block;
    doubleData<<<grid, block>>>(device_data, N); // GPU内核直接访问主机内存‌
    cudaDeviceSynchronize(); // 确保内核执行完成  
    cudaFreeHost(host_data); // 释放内存‌
    
     return 0;
}