The CUDA Handbook ▸ Chapter 6. Streams, Events, and Graphs

6.6 Mapped Pinned Memory: Concurrent Copying and Kernel Processing

Mapped pinned memory can be used to overlap PCI Express transfers and kernel processing – especially for device→host copies, where there is no need to cover the long latency to host memory. Mapped pinned memory has stricter alignment requirements than the native GPU memcpy, since they must be coalesced. Uncoalesced memory transactions run 2-6 times slower when using mapped pinned memory.

A naïve port of our concurrencyMemcpyKernelMapped.cu program yields an interesting result: on a cg1.4xlarge instance in Amazon EC2, mapped pinned memory runs very slowly for values of cycles below 64:

For small values of cycles, the kernel takes as long to run as if cycles were greater than 200! Only NVIDIA can discover the reason for this performance anomaly for certain, but it is not difficult to work around: by unrolling the inner loop of the kernel, we create more work per thread, and performance improves.

template<const int unrollFactor>__device__ voidAddKernel_helper( int *out, const int *in, size_t N, int increment, int cycles ){    for ( size_t i = unrollFactor*blockIdx.x*blockDim.x+threadIdx.x;                  i < N;                 i += unrollFactor*blockDim.x*gridDim.x )    {        int values[unrollFactor];         for ( int iUnroll = 0; iUnroll < unrollFactor; iUnroll++ ) {            size_t index = i+iUnroll*blockDim.x;            values[iUnroll] = in[index];        }        for ( int iUnroll = 0; iUnroll < unrollFactor; iUnroll++ ) {            for ( int k = 0; k < cycles; k++ ) {                values[iUnroll] += increment;            }        }        for ( int iUnroll = 0; iUnroll < unrollFactor; iUnroll++ ) {            size_t index = i+iUnroll*blockDim.x;            out[index] = values[iUnroll];        }    }} __device__ voidAddKernel( int *out, const int *in, size_t N, int increment, int cycles, int unrollFactor ){    switch ( unrollFactor ) {        case 1: return AddKernel_helper<1>( out, in, N, increment, cycles );        case 2: return AddKernel_helper<2>( out, in, N, increment, cycles );        case 4: return AddKernel_helper<4>( out, in, N, increment, cycles );    }}

Note that this version of AddKernel() is functionally identical to the previous one⁶ – it just computes unrollFactor outputs per loop iteration. Since the unroll factor is a template parameter, the compiler can use registers to hold the values array and the innermost for loops can be unrolled completely.

For unrollFactor==1, this implementation is identical to that of Listing 6-3. For unrollFactor==2, mapped pinned formulation shows some improvement over the streamed formulation; the tipping point drops from cycles==64 to cycles==48. For unrollFactor==4, performance is uniformly better than the streamed version:

These values are given for 32M integers, so the program reads and writes 128MB of data. For cycles==48, the program runs in 26ms. To achieve that effective bandwidth rate (more than 9GB/s over PCI Express 2.0), the GPU is concurrently reading and writing over PCI Express while performing the kernel processing!

The slowdown for small cycles, and the loop unrolling that works around it, belong to the hardware of the first edition. A modern GeForce RTX 3060 shows no such cliff: the naïve mapped-pinned kernel of Listing 6-3 runs neck-and-neck with the streamed formulation – both about 6ms for 32M integers – at every value of cycles, so the unrolling of Listing 6-6 is no longer needed to make mapped pinned memory competitive.