Chapter 07: Kernels — Writing and Using OpenCL Kernel Files

What is a kernel?

An OpenCL kernel is the entry point that executes on the GPU. It is declared with the __kernel qualifier and is compiled at runtime (by the OpenCL driver) from a source string assembled by the host program. Each invocation of the kernel processes one work-item — typically one element of an output array — in parallel with all other work-items.

In nmathopencl, kernels live in inst/cl/src/. Each file defines exactly one __kernel void function. The file name follows the pattern <function>_kernel.cl; for example, dnorm_kernel.cl implements the vectorized density computation for the normal distribution.

Anatomy of a simple kernel

// dnorm_kernel.cl
// Vectorized wrapper kernel for the public Mathlib dnorm interface.

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel void dnorm_kernel(
    __global const double* x,
    const double mu,
    const double sigma,
    const int give_log,
    __global double* out,
    const int n
) {
    int i = get_global_id(0);
    if (i >= n) return;

    out[i] = dnorm4(x[i], mu, sigma, give_log);
}

Key elements:

The if (i >= n) return; guard handles the case where the global work size is rounded up to a multiple of the local work-group size, leaving excess work-items that must not write out-of-bounds.

Standard argument layout

All kernels in inst/cl/src/ follow a consistent argument layout designed to work with openclPort::opencl_dbl_scalar_kernel_runner:

kernel(double arg0, double arg1, ..., double argN,
       __global double* out, int n_out)

Scalar double parameters come first (in any order needed by the underlying nmath function), followed by the output buffer and its length. Input vectors are also represented as __global const double* and precede the scalar arguments when present.

The *_ex_kernel.cl variants

Some distributions expose a secondary interface. The bessel_*_ex_kernel.cl files, for example, provide a version that returns both the function value and a sign/exponent component simultaneously. These exist to mirror the extended interfaces in nmath (bessel_i vs bessel_i with expo argument) and are distinguished by the _ex_ infix.

Kernel file index

The ~130 kernel files in inst/cl/src/ cover every function exposed by the nmathopencl R API:

Program assembly

Kernels are never compiled in isolation. The host-side runner assembles a complete OpenCL program string by concatenating:

1. The nmath library source (via opencltools::load_kernel_library("nmath", package = "nmathopencl"))
2. The kernel file (via opencltools::load_kernel_source("src/<name>_kernel.cl", package = "nmathopencl"))

This concatenated string is then passed to clBuildProgram. Because opencltools::load_kernel_library() performs topological sorting, the library functions are always defined before the kernel refers to them. Chapter 08 covers this assembly process in detail.

Writing a custom kernel

To write a kernel that calls an nmathopencl nmath function:

1. Identify the nmath function you need (e.g., dgamma).
2. Check its @provides tag in the corresponding .cl file to confirm the exact C signature.
3. Write a __kernel void function following the standard layout.
4. Load the nmath library subset your function needs using the @all_depends metadata as a guide (see Chapter 08).
5. Append your kernel source after the library source in the assembled program string.

Chapter 10 works through this process end-to-end for the glmbayes GLM log-posterior kernels.