<h3>Introduction</h3>
Zero Copy memory was introduced in CUDA quite a while back (~2009). However, there were very few takers for the capability because with time GPU memory has become reasonably large. The only takers were mainly folks with extremely large memory requirements, such as database processing.

Zero Copy memory is a way to map host memory and access it directly over PCIe without doing an explicit memory transfer. It is a way to allow CUDA kernels to directly access host memory. So instead of reading data from global memory (limit ~200GB/s), data would be read over PCIe and be limited by PCIe bandwidth (upto 16GB/s). Hence there was no real advantage.

However, with the release of the Jetson TK1 which uses the Tegra K1 processor, this has become really useful. The Jetson board has 2GB of physical memory that is shared by host and the CUDA GPU. If a cudaMemcpy Host->Device is done on the Jetson, the memory is simply being copied to a new location on the same physical memory and retrieving a CUDA pointer from it.In such a scenario, Zero Copy memory is perfect.

<h3>Standard CUDA Pipeline</h3>
Lets take a look at a standard pipeline in CUDA:

<pre>
// 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

</pre>

<h3>Zero Copy Memory CUDA Pipeline</h3>
Now lets look at a pipeline using Zero Copy memory

<pre>
// Set flag to enable zero copy memory
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_a, (void *) a, 0);
cudaHostGetDevicePointer((void **)&d_b, (void *) b, 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
</pre>

As you can see, it is considerably simpler to use zero copy memory on the Tegra from a coding perspective. The kernel code remains the same. You can also allocate host memory using the cudaHostAlloc call without using zero copy memory. This allows the fast pinned memory transfer when doing a cudaMemcpy.

<h3>Results</h3>
So how does it perform?

When I ran a transpose kernel using zero copy vs using standard pipeline, I achieve a bandwidth of 5.8 GB/s (23ms) vs 3.0 GB/s (45ms) for a 4096 x 4096 matrix. The bandwidth of a device-to-device copy kernel on the Tegra is ~6.6 GB/s.

The results are tremendous. There a a few caveats though. If you are running multiple kernels without modifying the data on the host side, it may be wise to run the standard pipeline. There is no right answer. Since it is so simple to modify the code for memory transfer vs zero copy pointers, it would be best to run both techniques and benchmark them.

<h3>Code</h3>
You can find the entire code from my transpose exercise here: <a href="http://sandbox.arrayfire.com/wp-content/uploads/2014/10/transpose.zip">transpose</a>: 
Note: If you wish to run this on a x86 system, make sure you change the compute version in the makefile from 32 to 30 (for Kepler).