The CUDA Handbook ▸ Chapter 3. Software Architecture

Prefer to read without ads? Become a member — from $10/month — and support the work. Already a member? Log in to read ad-free on this device.

3.12 The CUDA Runtime and CUDA Driver API

The CUDA runtime (“CUDART”) facilitates the language integration that makes CUDA so easy to program out of the gate. By automatically taking care of tasks such as initializing contexts and loading modules, and especially by enabling kernel invocation to be done in-line with other C++ code, CUDART lets developers focus on getting their code working quickly. A handful of CUDA abstractions, such as CUDA modules, are not accessible via CUDART.

In contrast, the driver API exposes all CUDA abstractions and enables them to be manipulated by developers as needed for the application. The driver API does not provide any performance benefit. Instead, it enables explicit resource management for applications that need it, like large-scale commercial applications with plug-in architectures.

The driver API is not noticeably faster than the CUDA runtime. If you are looking to improve performance in your CUDA application, look elsewhere.

Most CUDA features are available to both CUDART and the driver API, but a few are exclusive to one or the other; Table 3-9 summarizes the differences.

Feature CUDART Driver API
Device memory allocation * *
Pinned host memory allocation * *
Memory copies * *
CUDA streams * *
CUDA events * *
Graphics interoperability * *
Texture support * *
Surface support * *
cuMemGetAddressRange() *
Language integration *
“Fat binary” *
Explicit JIT options *
Simplified kernel invocation *
Explicit context and module management *
Context migration *
Float16 textures *
Memset of 16- and 32-bit values *
Compiler independence *

Table 3-9. CUDA Runtime v. Driver API Features

Between the two APIs, operations like memcpy tend to be functionally identical, but the interfaces can be quite different. The stream APIs are almost identical:

CUDART Driver API
cudaStream_t stream; CUstream stream;
cudaError_t status = cudaStreamCreate( &stream ); CUresult status = cuStreamCreate( &stream, 0 );
... ...
status = cudaStreamSynchronize( stream ); status = cuStreamSynchronize( stream );

The event APIs have minor differences, with CUDART providing a separate cudaEventCreateWithFlags() function if the developer wants to specify a flags word (needed to create a blocking event):

CUDART Driver API
cudaEvent_t eventPolling; CUevent eventPolling;
cudaEvent_t eventBlocking; CUevent eventBlocking;
cudaError_t status = cudaEventCreate( &eventPolling ); CUresult status = cuEventCreate( &eventPolling, 0 );
cudaError_t status = cudaEventCreateWithFlags( &eventBlocking, cudaEventBlockingSync ); CUresult status = cuEventCreate( &eventBlocking, CU_EVENT_BLOCKING_SYNC );
... ...
status = cudaEventSynchronize( event ); status = cuEventSynchronize( event );

The memcpy functions are the family where the interfaces are the most different, despite identical underlying functionality. CUDA supports three variants of memory (host, device and CUDA array), all permutations of participating memory types, and 1D, 2D or 3D memcpy; so the memcpy functions must contain either a large family of different functions, or a small number of functions that support many types of memcpy.

The simplest memcpy's in CUDA copy between host and device memory, but even those function interfaces are different: CUDART uses void * for the types of both host and device pointers and a single memcpy function with a direction parameter, while the driver API uses void * for host memory, CUdeviceptr for device memory, and three separate functions (cuMemcpyHtoD(), cuMemcpyDtoH(), cuMemcpyDtoD()) for the different memcpy directions. Here are equivalent CUDART and driver API formulations of the three permutations of host<->device memcpy:

CUDART Driver API
void *dptr; CUdeviceptr dptr;
void *hptr; void *src;
void *dptr2; CUdeviceptr dptr2;
status = cudaMemcpy( dptr, hptr, size, cudaMemcpyHostToDevice ); status = cuMemcpyHtoD( dptr, hptr, size );
status = cudaMemcpy( hptr, dptr, size, cudaMemcpyDeviceToHost ); status = cuMemcpyDtoH( hptr, dptr, size );
status = cudaMemcpy( dptr, dptr2, size, cudaMemcpyDeviceToDevice ); status = cuMemcpyDtoD( dptr, dptr2, size );

For 2D and 3D memcpy's, the driver API implements a handful of functions that take a descriptor struct and support all permutations of memcpy, including lower-dimension memcpy's. For example, if desired, cuMemcpy3D() can be used to perform a 1D host->device memcpy instead of cuMemcpyHtoD():

CUDA_MEMCPY3D cp = {0};
cp.dstMemoryType = CU_MEMORYTYPE_DEVICE;
cp.dstDevice = dptr;
cp.srcMemoryType = CU_MEMORYTYPE_HOST;
cp.srcHost = host;
cp.WidthInBytes = bytes;
cp.Height = cp.Depth = 1;
status = cuMemcpy3D( &cp );

CUDART uses a combination of descriptor structs for more-complicated memcpy's (e.g. cudaMemcpy3D()) while using different functions to cover the different memory types.

Like cuMemcpy3D(), CUDART's cudaMemcpy3D() function takes a descriptor struct that can describe any permutation of memcpy, including inter-dimensional memcpy's (e.g. performing a 1D copy to or from the row of a 2D CUDA array, or copying 2D CUDA arrays to or from slices of 3D CUDA arrays). Its descriptor struct is slightly different in that it embeds other structures; the two APIs’ 3D memcpy structures are compared side-by-side in Table 3-10.

CUDART Driver API
struct cudaMemcpy3DParms typedef struct CUDA_MEMCPY3D_st {
{ size_t srcXInBytes;
struct cudaArray *srcArray; size_t srcY;
struct cudaPos srcPos; size_t srcZ;
struct cudaPitchedPtr srcPtr; size_t srcLOD;
struct cudaArray *dstArray; CUmemorytype srcMemoryType;
struct cudaPos dstPos; const void *srcHost;
struct cudaPitchedPtr dstPtr; CUdeviceptr srcDevice;
struct cudaExtent extent; CUarray srcArray;
enum cudaMemcpyKind kind; void *reserved0;
}; size_t srcPitch;
struct cudaPos size_t srcHeight;
{ size_t dstXInBytes;
size_t x; size_t dstY;
size_t y; size_t dstZ;
size_t z; size_t dstLOD;
}; CUmemorytype dstMemoryType;
struct cudaPitchedPtr void *dstHost;
{ CUdeviceptr dstDevice;
void *ptr; CUarray dstArray;
size_t pitch; void *reserved1;
size_t xsize; size_t dstPitch;
size_t ysize; size_t dstHeight;
}; size_t WidthInBytes;
struct cudaExtent size_t Height;
{ size_t Depth;
size_t width; } CUDA_MEMCPY3D;
size_t height;
size_t depth;
};

Table 3-10. 3D Memcpy Structures

Usage of both 3D memcpy functions is similar. They are designed to be zero-initialized, and developers set the members needed for a given operation. For example, performing a host->3D array copy may be done as follows:

CUDART Driver API
struct cudaMemcpy3DParms cp = {0}; CUDA_MEMCPY3D cp = {0};
cp.srcPtr.ptr = host; cp.srcMemoryType = CU_MEMORYTYPE_HOST;
cp.srcPtr.pitch = pitch; cp.srcHost = host;
cp.dstArray = hArray; cp.srcPitch = pitch;
cp.extent.width = Width; cp.srcHeight = Height;
cp.extent.height = Height; cp.dstMemoryType = CU_MEMORYTYPE_ARRAY;
cp.extent.depth = Depth; cp.dstArray = hArray;
cp.kind = cudaMemcpyHostToDevice; cp.WidthInBytes = Width;
status = cudaMemcpy3D( &cp ); cp.Height = Height;
cp.Depth = Depth;
status = cuMemcpy3D( &cp );

For a 3D copy that covers the entire CUDA array, the source and destination offsets are set to 0 by the first line and don't have to be referenced again. Unlike parameters to a function, the code only needs to reference the parameters needed by the copy; and if the program must perform more than one similar copy (e.g., to populate more than one CUDA array or device memory region), the descriptor struct can be reused.