The CUDA Handbook ▸ Chapter 7. Kernel Execution

7.1 Launching Kernels

When using the CUDA runtime, a kernel launch is specified using the familiar triple-angle-bracket syntax:

Kernel<<<gridSize, blockSize, sharedMem, Stream>>>( Parameters… )

Kernel specifies the kernel to launch;

gridSize specifies the size of the grid in the form of a dim3() structure;

blockSize specifies the dimension of each threadblock as a dim3();

sharedMem specifies additional shared memory¹ to reserve for each block; and

Stream specifies the stream in which the kernel should be launched.

The dim3() structure used to specify the grid and block sizes has 3 members (x, y and z) and, when compiling with C++, a constructor with default parameters such that the y and z members default to 1. From vector_types.h:

struct __device_builtin__ dim3{    unsigned int x, y, z; #if defined(__cplusplus)     __host__ __device__ dim3(        unsigned int vx = 1,        unsigned int vy = 1,        unsigned int vz = 1) : x(vx), y(vy), z(vz) {}     __host__ __device__ dim3(uint3 v) : x(v.x), y(v.y), z(v.z) {}     __host__ __device__ operator uint3(void) {        uint3 t;        t.x = x;        t.y = y;        t.z = z;        return t;    } #endif /* __cplusplus */ };

Kernels can be launched via the driver API using cuLaunchKernel(), though that function takes the grid and block dimensions as discrete parameters rather than dim3():

CUresult cuLaunchKernel (
CUfunction kernel,
    unsigned int gridDimX,
    unsigned int gridDimY,
    unsigned int gridDimZ,
    unsigned int blockDimX,
    unsigned int blockDimY,
    unsigned int blockDimZ,
    unsigned int sharedMemBytes,
    CUstream hStream,
    void **kernelParams,
    void **extra
);

As with the triple-angle-bracket syntax, the parameters to cuLaunchKernel() include the kernel to invoke, the grid and block sizes, the amount of shared memory, and the stream. The main difference is in how the parameters to the kernel itself are specified: since the kernel microcode emitted by ptxas contains metadata that describes each kernel’s parameters², kernelParams is an array of void * where each element corresponds to a kernel parameter. Since the type is known by the driver, the correct amount of memory (4 bytes for an int, 8 bytes for a double, etc.) will be copied into the command buffer as part of the hardware-specific command used to invoke the kernel.

The fixed parameter list of <<< >>> and cuLaunchKernel() – grid, block, shared memory, stream – leaves no room for the per-launch options that later architectures introduced. cudaLaunchKernelEx() supplies that room. It takes a cudaLaunchConfig_t, which carries the same four parameters plus a pointer to an array of cudaLaunchAttribute values, each a tagged option: one attribute sets a thread block cluster’s dimensions (Section 7.8), another enables the programmatic serialization that dependent launch relies on (Section 7.9), still others select a stream priority or a preferred shared-memory carveout. The attribute array is the extensible launch interface – the mechanism through which several features later in this chapter reach the launch, and through which future ones will, without changing the function’s signature.

7.1.1 Limitations

All C++ classes participating in a kernel launch must be “plain old data” (POD) with the following characteristics:

• no user-declared constructors,

• no user-defined copy assignment operator,

• no user-defined destructor,

• no non-static data members that are not themselves PODs,

• no private or protected non-static data,

• no base classes, and

• no virtual functions.

Note that non-POD classes may be used in CUDA, or even in CUDA kernels; they simply cannot be used for a kernel launch. In that case, the classes used by a CUDA kernel can be constructed using the POD input data from the launch.

CUDA kernels also do not have return values. They must report their results back via device memory (which must be copied back to the CPU explicitly) or mapped host memory. Because a kernel returns nothing, the only way to detect errors such as an invalid kernel launch is to call cudaGetLastError() (Section 7.1.3).

7.1.2 Caches and Coherency

The GPU contains specialized caches to accelerate computation when reuse occurs: the constant cache is optimized for broadcast of read-only data to the execution units within an SM. This cache is not kept coherent with respect to writes to memory by the GPU; there is no protocol to enforce coherency between it and the L1 or L2 caches that serve to reduce latency and aggregate bandwidth to global memory. (Texturing once read through a separate texture cache with the same limitation, but on Maxwell and later hardware, the texture cache is unified with the L1.) That means two things:

1. when a kernel is running, it must take care not to write memory that it (or a concurrently-running kernel) is also reading through a cache that will not observe the write – such as the constant cache, or the read-only path used by texturing and __ldg(); and

2. the CUDA driver must invalidate these caches before each kernel launch.

7.1.3 Asynchrony and Error Handling

Kernel launches are asynchronous: as soon as a kernel is submitted to the hardware, it begins executing in parallel with the CPU³. This asynchrony complicates error handling: if a kernel encounters an error (for example, if it reads an invalid memory location), the error is reported to the driver (and the application) sometime after the kernel launch. The surest way to check for such errors is to synchronize with the GPU using cudaDeviceSynchronize() or cuCtxSynchronize(). If an error in kernel execution has occurred, the error code “unspecified launch failure” is returned.

Invalid Kernel Launches

It is possible to request a kernel launch that the hardware cannot perform: for example, by specifying more threads per block than the hardware supports. When possible, the driver detects these cases and reports an error rather than trying to submit the launch to the hardware.

The CUDA runtime and the driver API handle this case differently. When an invalid parameter is specified, the driver API’s explicit API calls such as cuLaunchGrid() and cuLaunchKernel() return error codes. But when using the CUDA runtime, since kernels are launched in-line with C/C++ code, there is no API call to return an error code. Instead, the error is “recorded” into a thread-local slot and applications can query the error value with cudaGetLastError(). This same error handling mechanism is used for kernel launches that are invalid for other reasons, such as a memory access violation. I discuss the rationale for this design in cudaGetLastError(): Its Raison D’Etre.

7.1.4 Timeouts

Because the GPU is not able to context-switch in the midst of kernel execution, a long-running CUDA kernel may negatively impact the interactivity of a system that uses the GPU to compose its output. As a result, many CUDA systems implement a “timeout” that resets the GPU if it runs too long without context switching.

On WDDM (Windows Display Driver Model), the timeout is enforced by the operating system. Microsoft has documented how this “Timeout Detection and Recovery” (TDR) works, including the Registry keys that control TDR behavior⁴. TDR can be safely disabled by using the Tesla Compute Cluster (TCC) driver, though the TCC driver is not available for all hardware.

On Linux, the NVIDIA driver enforces a default timeout of 2 seconds.

No timeout is enforced on secondary GPUs that are not being used for display. Developers can query whether a runtime limit is being enforced on a given GPU by calling cuDeviceGetAttribute() with CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, or by examining cudaDeviceProp::kernelExecTimeoutEnabled.

7.1.5 Local Memory

Since local memory is per-thread, and a grid in CUDA can contain thousands of threads, the amount of local memory needed by a CUDA grid can be considerable. The developers of CUDA took pains to preallocate resources, to minimize the likelihood that operations such as kernel launches would fail due to a lack of resources; but in the case of local memory, a conservative allocation simply would have consumed too much memory. As a result, kernels that use a large amount of local memory take longer, and may be synchronous, because the CUDA driver must allocate memory before performing the kernel launch. Furthermore, if the memory allocation fails, the kernel launch will fail due to a lack of resources.

By default, when the CUDA driver must allocate local memory to run a kernel, it frees the memory after the kernel has finished. This behavior additionally makes the kernel launch synchronous. But, this behavior can be inhibited by specifying CU_CTX_LMEM_RESIZE_TO_MAX to cuCtxCreate(), or by calling cudaSetDeviceFlags() with cudaDeviceLmemResizeToMax before the primary context is created. In this case, the increased amount of local memory available will persist after launching a kernel that required more local memory than the default.

7.1.6 Shared Memory

Shared memory is allocated at kernel launch time, and stays allocated for the duration of the kernel’s execution. Besides static allocations that can be declared in the kernel, shared memory can be declared as an unsized extern; in that case, the amount of shared memory to allocate for the unsized array is specified as the third parameter of the kernel launch, or the sharedMemBytes parameter to cuLaunchKernel().