The CUDA Handbook ▸ Chapter 11. Streaming Workloads

11.2 Asynchronous Memcpy

Unless the input and output data can stay resident on the GPU, the logistics of streaming the data through the GPU – copying the input and output data to and from device memory – become the primary consideration. The two tools best suited to improve transfer performance are pinned memory and asynchronous memcpy (which can only operate on pinned memory).

The stream2Async.cu application illustrates the effect of moving the pageable memory of stream1Device.cu to pinned memory, and invoking the memcpy’s asynchronously:

Measuring times with 128M floats (use --N to specify number of Mfloats)
Memcpy( host->device ): 40.05 ms (26809.70 MB/s)
Kernel processing     : 4.88 ms (329866.21 MB/s)
Memcpy (device->host ): 20.38 ms (26343.40 MB/s)

Total time (wall clock): 65.33 ms (24652.07 MB/s)

Listing 11-5 contrasts the difference between the timed portions of stream1Device.cu (which performs synchronous transfers) and stream2Async.cu (which performs asynchronous transfers)¹. In both cases, four CUDA events are used to record the times at the start, after the host→device transfers, after the kernel launch, and at the end. For stream2Async.cu, all of these operations are requested of the GPU in quick succession, and the GPU records the event times as it performs them. For stream1Device.cu, the GPU event-based times are a bit suspect, since any cudaMemcpy() calls must wait for the GPU to complete before proceeding, causing a pipeline bubble before the cudaEventRecord() calls for evHtoD and evDtoH are processed.

Note that despite using the slower, naïve implementation of saxpyGPU() (from Listing 11-2), the wall clock time from this application shows that it completes the computation about a third faster: 65 ms versus 88 ms. The combination of faster transfer performance and asynchronous execution delivers much better performance.

Despite the improved performance, the application output highlights another performance opportunity: some of the kernel processing can be performed concurrently with transfers. The next two sections describe two different methods to overlap kernel execution with transfers.

//// from stream1Device.cu//cudaEventRecord( evStart, 0 );cudaMemcpy( dptrX, hptrX, ..., cudaMemcpyHostToDevice );cudaMemcpy( dptrY, hptrY, ..., cudaMemcpyHostToDevice );cudaEventRecord( evHtoD, 0 );    saxpyGPU<<<nBlocks, nThreads>>>( dptrOut, dptrX, dptrY, N, alpha;cudaEventRecord( evKernel, 0 );cudaMemcpy( hptrOut, dptrOut, N*sizeof(float), cudaMemcpyDeviceToHost );cudaEventRecord( evDtoH, 0 );cudaDeviceSynchronize(); //// from stream2Async.cu//cudaEventRecord( evStart, 0 );cudaMemcpyAsync( dptrX, hptrX, ..., cudaMemcpyHostToDevice, NULL );cudaMemcpyAsync( dptrY, hptrY, ..., cudaMemcpyHostToDevice, NULL );cudaEventRecord( evHtoD, 0 );    saxpyGPU<<<nBlocks, nThreads>>>( dptrOut, dptrX, dptrY, N, alpha;cudaEventRecord( evKernel, 0 );cudaMemcpyAsync( hptrOut, dptrOut, N*sizeof(float), ... , NULL );cudaEventRecord( evDtoH, 0 );cudaDeviceSynchronize();