The CUDA Handbook ▸ Chapter 5. Memory

5.8 Memory Copy (memcpy)

CUDA has three different memory types – host memory, device memory, and CUDA arrays – and a full complement of functions to copy between them. For host↔︎device memcpy, an additional set of functions provide asynchronous memcpy between pinned host memory and device memory or CUDA arrays. Additionally, a set of peer-to-peer memcpy functions enable memory to be copied between GPUs.

The CUDA runtime and the driver API take very different approaches. For 1D memcpy, the driver API defined a family of functions with type-strong parameters. The host-to-device, device-to-host, and device-to-device memcpy functions are separate:

CUresult cuMemcpyHtoD(CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount);
CUresult cuMemcpyDtoH(void *dstHost, CUdeviceptr srcDevice, size_t ByteCount);
CUresult cuMemcpyDtoD(CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount);

In contrast, the CUDA runtime tends to define functions that take an extra “memcpy kind” parameter that depends on the memory types of the source and destination pointers:

enum cudaMemcpyKind
{
  cudaMemcpyHostToHost = 0,
  cudaMemcpyHostToDevice = 1,
  cudaMemcpyDeviceToHost = 2,
  cudaMemcpyDeviceToDevice = 3,
  cudaMemcpyDefault = 4
};

For more complex memcpy operations, both APIs use descriptor structures to specify the memcpy.

5.8.1 Synchronous Versus Asynchronous Memcpy

Because most variations of memcpy (dimensionality, memory type) are orthogonal to whether the memory copy is asynchronous, this section will discuss the difference in some detail and later sections will include minimal coverage of synchronous memcpy.

By default, any memcpy involving host memory is synchronous: the function does not return until after the operation has been performed¹⁷. Even when operating on pinned memory, such as memory allocated with cudaMallocHost(), synchronous memcpy routines must wait until the operation is completed because the application may rely on that behavior¹⁸.

When possible, synchronous memcpy should be avoided for performance reasons – even when streams are not being used, keeping all operations asynchronous improves performance by enabling the CPU and GPU to run concurrently. If nothing else, the CPU can set up more GPU operations such as kernel launches and other memcpy’s while the GPU is running! If CPU/GPU concurrency is the only goal, there is no need to create any CUDA streams – calling an asynchronous memcpy with the NULL stream will suffice.

While memcpy’s involving host memory are synchronous by default, any memory copy not involving host memory (device↔︎device, or device↔︎array) is asynchronous – the GPU hardware internally enforces serialization on these operations, so there is no need for the functions to wait until the GPU has finished before returning.

Asynchronous memcpy functions have the suffix Async(). For example, the driver API function for asynchronous host→device memcpy is cuMemcpyHtoDAsync() and the CUDA runtime function is cudaMemcpyAsync().

The hardware that implements asynchronous memcpy has evolved over time. The very first CUDA-capable GPU (the GeForce 8800 GTX) did not have any copy engines, so asynchronous memcpy only enabled CPU/GPU concurrency. Later GPUs added copy engines that could perform 1D transfers while the SMs were running, and still later, fully-capable copy engines were added that could accelerate 2D and 3D transfers, even if the copy involved converting between pitch layouts and the block-linear layouts used by CUDA arrays. Additionally, early CUDA hardware only had one Copy Engine, while more recent CUDA hardware sometimes has 2. More than two Copy Engines wouldn’t necessarily make sense: since a single Copy Engine can saturate the PCI Express bus in one direction, only two Copy Engines are needed to maximize both bus performance and concurrency between bus transfers and GPU computation.

The number of Copy Engines can be queried by calling cuDeviceGetAttribute() with CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT, or calling cudaGetDeviceProperties() and examining cudaDeviceProp::asyncEngineCount.

5.8.2 Unified Virtual Addressing

Unified Virtual Addressing enables CUDA to make inferences about memory types based on address ranges: because CUDA tracks which address ranges contain device addresses versus host addresses, there is no need to specify the cudaMemcpyKind parameter to the cudaMemcpy() function. The driver API added a cuMemcpy() function that similarly infers the memory types from the addresses:

CUresult cuMemcpy(CUdeviceptr dst, CUdeviceptr src, size_t ByteCount);

The CUDA runtime equivalent, not surprisingly, is called cudaMemcpy():

cudaError_t cudaMemcpy( void *dst, const void *src, size_t bytes );

5.8.3 CUDA Runtime

Table 5-10 summarizes the memcpy functions available in the CUDA runtime; the interactive picker below answers the same question directly.

Table 5-10. Memcpy Functions (CUDA Runtime)

1D and 2D memcpy functions take base pointers, pitches, and sizes as required; the 3D memcpy routines take a descriptor structure cudaMemcpy3Dparms, defined as follows:

struct cudaMemcpy3DParms
{
  struct cudaArray *srcArray;
  struct cudaPos srcPos;
  struct cudaPitchedPtr srcPtr;
  struct cudaArray *dstArray;
  struct cudaPos dstPos;
  struct cudaPitchedPtr dstPtr;
  struct cudaExtent extent;
  enum cudaMemcpyKind kind;
};

Table 5-11. cudaMemcpy3DParms structure members.

The cudaPos and cudaExtent structures are defined as follows:

struct cudaExtent {
  size_t width;
  size_t height;
  size_t depth;
};
struct cudaPos {
  size_t x;
  size_t y;
  size_t z;
};

5.8.4 Driver API

Table 5-12 summarizes the driver API's memcpy functions; the interactive picker below answers the same question directly.

Table 5-12. Memcpy Functions (Driver API)

cuMemcpy3D() is designed to implement a strict superset of all previous memcpy functionality: any 1D, 2D or 3D memcpy may be performed between any of host, device or CUDA array memory, and any offset into either the source or destination may be applied. The WidthInBytes, Height and Depth members of the input structure, CUDA_MEMCPY3D, define the dimensionality of the memcpy: Height==0 implies a 1D memcpy and Depth==0 implies a 2D memcpy. The source and destination memory types are given by the srcMemoryType and dstMemoryType structure elements, respectively.

Structure elements that are not needed by cuMemcpy3D() are defined to be ignored. For example, if a 1D host->device memcpy is requested, the srcPitch, srcHeight, dstPitch and dstHeight elements are ignored. If srcMemoryType is CU_MEMORYTYPE_HOST, the srcDevice and srcArray elements are ignored. This API semantic, coupled with the C99 “designated initializer” language feature that defines unnamed members to be zero-initialized, enables memory copies to be described very concisely. (One portability note: C99 permits designated initializers in any order, but C++20 – and therefore nvcc – requires them to appear in declaration order, which is why the src members are initialized before the dst members below.) Most memcpy functions can be implemented in a few lines of code, e.g.:

CUresult
my_cuMemcpyHtoD( CUdeviceptr dst, const void *src, size_t N )
{
    CUDA_MEMCPY3D cp = {
        .srcMemoryType = CU_MEMORYTYPE_HOST,
        .srcHost = src,
        .dstMemoryType = CU_MEMORYTYPE_DEVICE,
        .dstDevice = dst,
        .WidthInBytes = N };
    return cuMemcpy3D( &cp );
}