Chapter A10: Accelerated EnvelopeBuild Implementation using OpenCL"

In glmbayes: Bayesian Generalized Linear Models (IID Samples)

knitr::opts_chunk$set(
  collapse = TRUE,
  comment = "#>"
)

library(glmbayes)

Chapter A10: Accelerated EnvelopeBuild Implementation using OpenCL

1. Introduction

The glmbayes package constructs envelopes for accept–reject sampling by evaluating the negative log posterior (f2) and its gradient (f3) at each point of a tangency grid. The grid size grows with dimension (e.g., (3^\ell) points for (\ell) coefficients), and each grid point can be evaluated independently. These evaluations are embarrassingly parallel and are a natural target for GPU acceleration.

The optional OpenCL implementation accelerates these evaluations by running them on a GPU (or OpenCL-capable CPU). OpenCL is vendor-neutral and works across NVIDIA, AMD, and Intel hardware. When OpenCL is available and use_opencl = TRUE, the package uses GPU-accelerated evaluation inside EnvelopeEval, which is called by EnvelopeBuild and EnvelopeDispersionBuild. If OpenCL is not available or use_opencl = FALSE, the CPU path (f2_f3_non_opencl) is used instead. The package runs correctly in either case.

This chapter describes the OpenCL implementation: the two-layer design (kernel wrapper and kernel runner), how the OpenCL program is assembled from ported nmath/rmath functions, and how the pilot and safeguards integrate with the envelope build. For installation and availability checks, see Chapter 16.

2. Architecture Overview

The OpenCL path consists of two layers:

| Layer | Function | Location | Role | |-------|----------|----------|------| | Wrapper | f2_f3_opencl | kernel_wrappers.cpp | Flatten R inputs, select kernel, load sources, call runner, reshape outputs | | Runner | f2_f3_kernel_runner | kernel_runners.cpp | Platform/device setup, buffer creation, kernel launch, readback |

2.1 Call Path

EnvelopeEval (EnvelopeEval.cpp)
    |-- if use_opencl && pilot/check passes:
    |       f2_f3_opencl() -> f2_f3_kernel_runner()
    \-- else:
            f2_f3_non_opencl() -> CPU famfuncs (f2_f3_binomial_logit, etc.)

EnvelopeEval receives the grid G4 (coefficients at each tangency point), the design matrix, prior parameters, and family/link. It chooses the GPU or CPU path based on use_opencl and the outcome of the pilot (when applicable). The GPU path calls f2_f3_opencl, which assembles the OpenCL program and invokes f2_f3_kernel_runner; the CPU path calls f2_f3_non_opencl, which dispatches to the appropriate C++ family function.

3. Program Construction

OpenCL kernels are not stored as single monolithic files. Instead, the package builds the program by concatenating several source components in a fixed order. This mirrors a C/C++ build where headers and libraries are included before the main source.

3.1 Assembly order (production)

Runtime GPU assembly uses load_likelihood_subgradient_program in src/kernel_loader.cpp (see inst/cl/README.md). Prelude, shims, and selective nmath/ are loaded from nmathopencl; entry kernels from glmbayes (inst/cl/src/f2_f3_*.cl).

For interactive exploration, mirror that split with opencltools:

all_src = opencltools::load_kernel_source("OPENCL.cl", package = "nmathopencl")
        + opencltools::load_kernel_library("libR_shims", package = "nmathopencl")
        + … (R_ext_types, R_shims, R_ext_runtime, R_ext_internals, System)
        + opencltools::load_library_for_kernel(
              system.file("cl/src/f2_f3_binomial_logit.cl", package = "glmbayes"),
              system.file("cl/nmath", package = "nmathopencl"),
              depends_tag = "all_depends_nmath")
        + opencltools::load_kernel_source("src/f2_f3_binomial_logit.cl", package = "glmbayes")

1. OPENCL.cl and shim layers — from nmathopencl.
2. nmath subset — stems from @all_depends_nmath on the entry kernel; files from nmathopencl inst/cl/nmath/.
3. Kernel file — model-specific kernel from glmbayes inst/cl/src/.

3.2 opencltools loaders

• opencltools::load_kernel_source(relative_path, package = …) – Loads a single .cl file via system.file("cl", …).
• opencltools::load_kernel_library(subdir, package = …) – Loads all .cl files in a subdirectory with @provides / @depends topological sort.
• opencltools::load_library_for_kernel() – Minimal nmath subset from @all_depends_nmath tags (see ?opencltools::load_library_for_kernel).

The @provides and @depends tags allow modular, maintainable kernel code. See ?opencltools::load_kernel_library for details.

4. Ported Math Libraries (nmath / rmath / dpq)

The likelihood and gradient computations require statistical functions that match R's behavior. Because OpenCL C does not include these, the package ports a core set from R's nmath and rmath libraries to OpenCL C.

4.1 nmath

| Module | Provides | Purpose | |--------|----------|---------| | nmath.cl | ML_, ME_, ISNAN, R_FINITE, ML_ERROR, etc. | Constants, error handling, validation macros | | log1p.cl | Rlog1p | Log(1+x) for numerical stability | | expm1.cl | expm1 | exp(x)-1 for numerical stability | | bd0.cl | bd0, ebd0, etc. | Poisson/binomial deviance terms | | stirlerr.cl | stirlerr | Stirling error for factorial terms | | lgamma.cl | lgammafn | Log-gamma | | gamma.cl | gammafn | Gamma function | | dbinom.cl | dbinom, dbinom_raw | Binomial density | | dpois.cl | dpois, dpois_raw | Poisson density | | dgamma.cl | dgamma | Gamma density | | dnorm.cl | dnorm4 | Normal density | | pnorm.cl | pnorm5, pnorm_both | Normal CDF (probit) | | pgamma.cl | pgamma | Gamma CDF (inverse link) | | d1mach.cl | d1mach | Machine constants | | chebyshev.cl | chebyshev_eval | Polynomial evaluation |

4.2 rmath

Rmath.cl provides additional constants (M_E, M_PI, M_LN2, etc.) and distribution functions consistent with R's Rmath library.

4.3 dpq

dpq.cl and dpq_prelude.cl provide macros for density and CDF handling with give_log and lower_tail, matching R's DPQ (density, probability, quantile) conventions.

These ports ensure that the OpenCL kernels produce results numerically consistent with the CPU path and with R itself.

5. Family/Link Kernels

Each supported family/link has a dedicated kernel file that implements the f2/f3 logic for that model.

| Family | Link | Kernel file | |--------|------|-------------| | binomial | logit | f2_f3_binomial_logit.cl | | binomial | probit | f2_f3_binomial_probit.cl | | binomial | cloglog | f2_f3_binomial_cloglog.cl | | poisson | log | f2_f3_poisson.cl | | Gamma | inverse | f2_f3_gamma.cl | | gaussian | identity | f2_f3_gaussian.cl |

5.1 Kernel Structure

Each kernel follows the same pattern:

1. Work-item mapping – int j = get_global_id(0); one work item per grid point (j).
2. Prior term – Compute (P(\beta_j - \mu)) and the quadratic form (\tfrac{1}{2}(\beta_j - \mu)' P (\beta_j - \mu)), accumulate into qf[j].
3. Prior gradient – Initialize local gradient with (P(\beta_j - \mu)).
4. Data loop – For each observation (i), compute linear predictor, link, likelihood contribution, and gradient contribution. Use ported functions (dbinom_raw, dpois_raw, lgamma, pnorm5, etc.) as needed.
5. Write outputs – qf[j] = negative log posterior; grad[k*m1 + j] = gradient (column-major layout).

The kernels use a fixed MAX_L2 for local arrays (e.g., 64); this limits the number of coefficients when using OpenCL. See the source for current limits.

6. f2_f3_opencl Flow

The wrapper f2_f3_opencl in kernel_wrappers.cpp performs:

1. Input flattening – Convert R matrices/vectors (x, b, mu, P, alpha, y, wt) to contiguous C++ vectors in the layout expected by the runner.
2. Kernel selection – Map (family, link) to kernel_name and kernel_file (e.g., binomial + logit → f2_f3_binomial_logit, src/f2_f3_binomial_logit.cl).
3. Program assembly – Load OPENCL.cl, rmath, nmath, dpq, and the kernel file; concatenate into all_src.
4. Runner call – Invoke f2_f3_kernel_runner(all_src, kernel_name, l1, l2, m1, ...) with flattened inputs and output buffers.
5. Output reshaping – Copy qf_flat into an R vector; wrap grad_flat as an Armadillo matrix (m1 × l2) for return.

The returned list has components qf (negative log posterior per grid point) and grad (gradient matrix).

7. f2_f3_kernel_runner

The runner in kernel_runners.cpp handles the low-level OpenCL API:

1. Platform and device – clGetPlatformIDs, clGetDeviceIDs (CL_DEVICE_TYPE_DEFAULT).
2. Context and queue – clCreateContext, clCreateCommandQueueWithProperties.
3. Program – clCreateProgramWithSource with the concatenated string, clBuildProgram.
4. Kernel – clCreateKernel with the kernel name.
5. Buffers – Create read-only buffers for X, B, mu, P, alpha, y, wt; write-only for qf and grad. Copy host data for inputs.
6. Launch – clEnqueueNDRangeKernel with global = m1 (one work item per grid point).
7. Readback – clEnqueueReadBuffer for qf and grad.
8. Sanity check – If both outputs are all zeros, throw (likely kernel failure).
9. Cleanup – Release buffers, kernel, program, queue, context.

The runner is GLM-specific and lives in the glmbayes::opencl namespace; the general OpenCL utilities (kernel loading, device enumeration) are in openclPort.

8. Pilot and Safeguards

For large grids (e.g., (m_1 > 50{,}000)), EnvelopeEval runs f2_f3_opencl_pilot before the full evaluation to estimate runtime.

8.1 Pilot Logic

• Warm-up – Run f2_f3_opencl on a small grid slice.
• Calibration – Run on slices of ~1% and ~2% of the grid to estimate fixed cost and per-grid cost.
• Refined estimate – Compute refined_est_total_sec = per_grid_sec × m1.
• 5-minute safeguard – If refined_est_total_sec > 300:
• Interactive session: Prompt "Do you want to continue? [y/N]: ".
• Non-interactive session: Proceed automatically (e.g., CI, batch).

This mirrors the parallel sampling pilots in Chapter A09. The envelope build phase cannot be interrupted once the full OpenCL run starts, so the pilot gives users a chance to decline long runs.

8.2 When the Pilot Runs

The pilot runs when m1_total > 50000 and use_opencl is true. For smaller grids, the full evaluation runs without a pilot.

9. Installation and Availability

OpenCL support is optional. The package compiles and runs without it; in that case, use_opencl is ignored and the CPU path is always used.

To enable OpenCL:

• Install OpenCL headers and development libraries (see Chapter 16).
• Build the package from source with OpenCL enabled (configure script detects OpenCL).
• Ensure an OpenCL runtime (ICD loader, vendor driver) is available at runtime.

Use has_opencl() to check availability and diagnose_glmbayes() for a full diagnostic report. See Chapter 16 for platform-specific installation instructions (Windows, Linux, macOS) and AMD ROCm notes.

10. Cross-References

• Chapter A05 – Simulation pipeline (optimization, standardization, envelope build, sampling)
• Chapter A08 – Envelope-related functions
• Chapter A09 – Parallel sampling (pilot pattern, 5-minute safeguard)
• Chapter 16 – GPU acceleration: installation, has_opencl, diagnose_glmbayes
• ?EnvelopeBuild, ?EnvelopeEval – use_opencl argument
• ?opencltools::load_kernel_library, ?opencltools::load_kernel_source – Program assembly
• ?glmb, ?rglmb – use_opencl argument

glmbayes documentation built on Aug. 5, 2026, 1:07 a.m.