转载

NVIDIA Tegra TK/X系列板子的零拷贝(zero copy)问题

Introduction

本文系翻译，原文地址：

Zero Copy on Tegra K1

之前写过的关于cuda 零拷贝的文章：

CUDA零复制内存

CUDA锁页内存和零复制

参考里面的两篇论坛上的讨论给出的观点也是零拷贝在unified memory上是没有发生copy行为的，强烈推荐看他们的讨论，可能需要翻墙。

零拷贝（Zero Copy）已经加入CUDA Toolkit很长时间了。但是很少有应用使用这个特性，因为随着GPU的发展，显存容量已经很大了。现有的使用零拷贝的应用主要是数据库处理的应用，因为它有及其严格的内存要求。

零拷贝是一个映射主机内存到GPU上，使得CUDA内核可以直接通过 PCIe 访问，而不用明确的内存传输。使用零拷贝后，读数据是从PCIe而不是从全局显存了，所以现在读取的速度是限制于PCIe的速度(最多16GB/s)而不是全局内存的速度(大约200GB/s)了。因此，零拷贝实际上对大多数应用没有真实的性能提升。

然而，随着Jetson TK1板子的面世，零拷贝变得非常有用。TK1板子有2GB的物理内存，ARM CPU和 NVIDIA GPU共享。如果执行一个 cudaMemcpy 命令，从Host拷贝到GPU，在TK1板子上，数据只不过是从内存的一个地方拷贝到了同一个内存的另外一个地方。在这种情况下，零拷贝就变得很必要了。

标准的CUDA Pipeline

// Host Arrays float* h_in  = new float[sizeIn]; float* h_out = new float[sizeOut];  //Process h_in  // Device arrays float *d_out, *d_in;  // Allocate memory on the device cudaMalloc((void **) &d_in,  sizeIn )); cudaMalloc((void **) &d_out, sizeOut));  // Copy array contents of input from the host (CPU) to the device (GPU) cudaMemcpy(d_in, h_in, sizeX * sizeY * sizeof(float), cudaMemcpyHostToDevice);  // Launch the GPU kernel kernel<<<blocks, threads>>>(d_out, d_in);  // Copy result back cudaMemcpy(h_out, d_out, sizeOut, cudaMemcpyDeviceToHost);  // Continue processing on host using h_out

零拷贝的CUDA pipeline

// Set flag to enable zero copy access cudaSetDeviceFlags(cudaDeviceMapHost);  // Host Arrays float* h_in  = NULL; float* h_out = NULL;  // Process h_in  // Allocate host memory using CUDA allocation calls cudaHostAlloc((void **)&h_in,  sizeIn,  cudaHostAllocMapped); cudaHostAlloc((void **)&h_out, sizeOut, cudaHostAllocMapped);  // Device arrays float *d_out, *d_in; // Get device pointer from host memory. No allocation or memcpy cudaHostGetDevicePointer((void **)&d_in,  (void *) h_in , 0); cudaHostGetDevicePointer((void **)&d_out, (void *) h_out, 0);  // Launch the GPU kernel kernel<<<blocks, threads>>>(d_out, d_in);  // No need to copy d_out back // Continue processing on host using h_out

在Tegra上用零拷贝看起来代码简单一点。kernel函数的代码不变。You can also allocate host memory using the cudaHostAlloc call without using zero copy access. This allows the fast pinned memory transfer when doing a cudaMemcpy.