Skip to content

CUDA coding standard

These rules extend the C++ coding standard for code compiled with CUDA. The VMC examples show one project implementation; the rules state which parts should carry across UWHPC projects.

Language and backend boundary

Section titled “Language and backend boundary”
  • Must: require CUDA C++20 or newer for new CUDA code.
  • Must: make the CUDA backend an explicit build-time feature.
  • Must: keep backend-specific compilation behind a single project feature definition rather than scattered compiler guesses.
  • Should: provide a CPU implementation for portable numerical code when the algorithm and maintenance cost permit it.
  • Must: keep public behavior consistent across CPU and CUDA backends; test both where supported.
  • Must: when a function supports both backends, keep one signature and select the body inside it with #ifdef VMC_CUDA_BACKEND; do not declare two file-scope definitions of the same function.

VMC detects CUDA, sets CUDA C++20, and defines VMC_CUDA_BACKEND on the library: CMakeLists.txt, CMakeLists.txt. Its .cu files select a kernel path with #ifdef VMC_CUDA_BACKEND and retain a host implementation after #else, for example JastrowPade::value. SlaterPlaneWave::add_derivatives shows the same rule at the member-function level: one signature whose body branches on the backend, with the CUDA kernel in an anonymous namespace above it: src/slater_plane_wave/add_derivatives.cu.

File template

Section titled “File template”

A ported component is two files: a backend-agnostic header and an implementation whose backend is selected at compile time. Copy the skeletons below and replace each // <...> placeholder. The rules that shape each block are detailed in the sections that follow.

Header — one declaration per operation, identical for both backends:

component.cuh
#pragma once


#include "../utilities/math.cuh"   // real_t, RESTRICT, CUDA_CHECK, vmc:: helpers
// <other dependency headers>


class Component {
private:
  // <members_, each suffixed with a trailing underscore>


public:
  explicit Component(/* <ctor args> */) noexcept;


  // The same signature serves the CPU and CUDA definitions.
  [[nodiscard]] real_t operation(const Particles& particles) const noexcept;
};

Implementation — kernels in an anonymous namespace under the backend guard, then one member definition whose body branches on the backend:

component.cu
#include "component.cuh"


#ifdef VMC_CUDA_BACKEND
#include <cstddef>
namespace {


__global__
void cudaOperation(
  std::size_t n,                 // sizes and scalar constants first
  const real_t* RESTRICT in_x,   // input pointers next
  real_t* RESTRICT out           // output pointers last
) {
  const auto [i]{vmc::cudaThreadIdx<1>()};
  if (i >= n) { return; }


  // <kernel body>
}


} // namespace
#else
// <CPU-only includes, e.g. <cstring>; omit this #else block if none are needed>
#endif


real_t Component::operation(const Particles& particles) const noexcept {
#ifdef VMC_CUDA_BACKEND
  // Host wrapper: launch config, launch mirroring the signature, checks.
  dim3 operationThreads(256);
  dim3 operationBlocks(
    vmc::cudaNumBlocks(/* <n> */, operationThreads.x)
  );
  cudaOperation<<<operationBlocks, operationThreads>>>(
    // <same argument grouping as the kernel signature>
  );
  CUDA_CHECK(cudaGetLastError());
  CUDA_CHECK(cudaDeviceSynchronize());


  // <read result and return>
#else
  // <CPU implementation>
#endif
}

For a filled-in version of this exact shape, see src/slater_plane_wave/add_derivatives.cu.

Kernel shape and bounds

Section titled “Kernel shape and bounds”
  • Must: derive a thread’s global index from CUDA launch coordinates or one shared, reviewed helper.
  • Must: guard every global index before reading or writing user data.
  • Must: guard each dimension of a multidimensional launch.
  • Should: return early for inactive work rather than nesting the kernel body.
  • Must: derive grid dimensions from the problem size; do not assume exact divisibility by the block size.

VMC example:

__global__
void cudaValue(
  std::size_t num_particles,
  real_t L, real_t a, real_t b,
  const real_t* RESTRICT p_x, const real_t* RESTRICT p_y, const real_t* RESTRICT p_z,
  real_t* RESTRICT jastrow_pade
) {
  const std::size_t i{blockIdx.x * blockDim.x + threadIdx.x};
  if (i >= num_particles) { return; }


  const std::size_t j{blockIdx.y * blockDim.y + threadIdx.y};
  if (j >= num_particles) { return; }

Source: src/jastrow_pade/value_delta_value.cu.

For two-dimensional work, VMC computes each grid dimension with a ceiling division helper:

dim3 valueThreads(16, 16);
dim3 valueBlocks(
  vmc::cudaNumBlocks(num_particles, valueThreads.x),
  vmc::cudaNumBlocks(num_particles, valueThreads.y)
);

Source: src/jastrow_pade/value_delta_value.cu.

Kernel layout and formatting

Section titled “Kernel layout and formatting”
  • Must: place __global__ on its own line, directly above the kernel’s return type and name.
  • Must: name a device kernel cuda followed by the operation in PascalCase, for example cudaValue or cudaBuildRow.
  • Must: group kernel parameters by role, one group per line: sizes and scalar constants first, input pointers next, output pointers last; keep related pointers on the same line.
  • Should: lay out the launch (<<<...>>>) argument list in the same grouping as the kernel signature so the call reads as a mirror of the declaration.
  • Must: name launch-configuration variables after the kernel with the cuda prefix dropped and Threads or Blocks appended (cudaUpdateTrigRow → updateTrigRowThreads, updateTrigRowBlocks).
  • Must: qualify member-function calls inside a member function with this->.
  • Must: put function attributes and CUDA qualifiers on their own line above the return type.

VMC example — kernel signature, grouped parameters, and namespace layout:

#ifdef VMC_CUDA_BACKEND
#include <cstddef>
namespace {


__global__
void cudaValue(
  std::size_t num_particles,
  real_t L, real_t a, real_t b,
  const real_t* RESTRICT p_x, const real_t* RESTRICT p_y, const real_t* RESTRICT p_z,
  real_t* RESTRICT jastrow_pade
) {
  const std::size_t i{blockIdx.x * blockDim.x + threadIdx.x};
  if (i >= num_particles) { return; }

Source: src/jastrow_pade/value_delta_value.cu.

VMC example — launch-configuration naming and a launch that mirrors the signature:

dim3 updateTrigRowThreads(256);
dim3 updateTrigRowBlocks(
  vmc::cudaNumBlocks(num_k, updateTrigRowThreads.x)
);
cudaUpdateTrigRow<<<updateTrigRowBlocks, updateTrigRowThreads>>>(
  num_k, particle * ROW_STRIDE,
  particles.pos().x_[particle], particles.pos().y_[particle], particles.pos().z_[particle],
  this->k_vector().x_, this->k_vector().y_, this->k_vector().z_,
  this->sin_cache(), this->cos_cache()
);

Source: src/slater_plane_wave/update_restore_trig.cu.

VMC puts attributes and qualifiers on their own line above the return type:

CUDA_CALLABLE [[nodiscard]]
inline real_t sqrt(real_t arg) noexcept {

Source: src/utilities/math.cuh.

Minimal kernel arguments

Section titled “Minimal kernel arguments”
  • Should: pass a primitive input and derive trivial quantities inside the kernel rather than passing many precomputed values; for example pass const real_t L and compute the wrap constants in the kernel.
  • Should: pass the single value a kernel needs rather than a full array plus an index when only one element is read.
  • Should: keep getter-derived caching on the CPU path; it is evaluated once on the GPU path and only dilutes the kernel.

VMC passes the box length and derives the wrap constants inside the kernel:

  real_t L, real_t a, real_t b,
  ...
  const real_t neg_L{-1.0_r * L};
  const real_t half_L{0.5_r * L};
  const real_t neg_half_L{-1.0_r * half_L};

Source: src/jastrow_pade/value_delta_value.cu. cudaUpdateTrigRow likewise takes the moved particle’s position by value rather than the whole array: src/slater_plane_wave/update_restore_trig.cu.

Launches and errors

Section titled “Launches and errors”
  • Must: check every CUDA kernel launch with cudaGetLastError().
  • Must: check every CUDA Runtime and numerical-library call.
  • Must: include the source file, source line, and runtime message in fatal CUDA diagnostics.
  • Should: centralize checks in a macro or function that evaluates its argument exactly once.
  • Should: synchronize only at a required host-consumption or phase boundary; document extra synchronizations used for debugging or deterministic failure localization.

VMC example:

cudaValue<<<valueBlocks, valueThreads>>>(
  num_particles,
  box_length_, this->a(), this->b(),
  particles.pos().x_, particles.pos().y_, particles.pos().z_,
  jastrow_pade[0]
);


CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
return 0.5_r * *jastrow_pade[0];

Source: src/jastrow_pade/value_delta_value.cu.

The checked-call macro evaluates the call once, prints location and the CUDA message, and aborts:

#define CUDA_CHECK(call) \
  do { \
    cudaError_t cuda_check_result_{(call)}; \
    if (cuda_check_result_ != cudaSuccess) { \
      std::fprintf( \
        stderr, \
        "CUDA error at %s:%d: %s\n", \
        __FILE__, \
        __LINE__, \
        cudaGetErrorString(cuda_check_result_) \
      ); \
      std::abort(); \
    } \
  } while (0)

Source: src/utilities/macros.cuh.

Memory and data access

Section titled “Memory and data access”
  • Must: put device or managed allocations under RAII ownership.
  • Must: pair each allocation API with its matching deallocation API in one abstraction.
  • Should: use structure-of-arrays storage for kernels that access the same field across many elements.
  • May: mark non-aliasing pointer parameters with a portable RESTRICT abstraction, but only when the caller satisfies that non-aliasing contract.
  • Must: initialize storage before kernels consume it.
  • Should: for a single device scalar used as a reduction target, pass the buffer’s [0] element as the kernel’s output pointer and read the result as *buf[0]; do not wrap the scalar in extra storage. VMC’s checked launch passes jastrow_pade[0] and returns *jastrow_pade[0] (value_delta_value.cu).

VMC routes managed allocation through a backend allocator and owns it through a custom std::unique_ptr deleter:

void* ptr{};
CUDA_CHECK(cudaMallocManaged(&ptr, size));


return ptr;
struct SoADeleter {
  template <typename T>
  void operator()(T* ptr) const {
    backend_free(ptr);
  }
};

Source: src/utilities/aligned_soa.cuh.

Particles stores positions and derivatives as separate aligned arrays: src/particles/particles.cuh.

Host/device portability

Section titled “Host/device portability”
  • Must: isolate host/device annotations behind a project macro when the same function is compiled for both backends.
  • Should: wrap mathematical operations whose host and device spellings or precision behavior differ.
  • Must: branch on __CUDA_ARCH__ only for device-compilation behavior, not as the project’s general “CUDA enabled” switch.
  • Must: keep backend wrappers small, inline, and unit-testable on the host.
  • Must: do not mark a function CUDA_CALLABLE when it calls device-only APIs. CUDA_CALLABLE expands to __host__ __device__, so split the function with #if defined(__CUDA_ARCH__) into a __device__ body and a host body instead.
  • Should: keep noexcept on a wrapper that also covers the host path, and remove noexcept from a GPU-only function.
  • Should: define a trivial special member with = default rather than an empty user-provided body; an empty {} constructor makes the type non-trivial. Initialize every member when you do so.

VMC’s CUDA_CALLABLE macro expands to __host__ __device__ only in a CUDA build:

#ifdef VMC_CUDA_BACKEND
  #define CUDA_CALLABLE __host__ __device__
#else
  #define CUDA_CALLABLE
#endif

Source: src/utilities/macros.cuh.

Its math wrapper chooses a device implementation only under __CUDA_ARCH__ and uses std::sqrt on the host:

CUDA_CALLABLE [[nodiscard]]
inline real_t sqrt(real_t arg) noexcept {
#if defined(VMC_CUDA_BACKEND) && defined(__CUDA_ARCH__)
  #ifdef FP_64
    return ::sqrt(arg);
  #elif defined(VMC_FAST_MATH)
    return __fsqrt_rn(arg);
  #else
    return sqrtf(arg);
  #endif
#else
  return std::sqrt(arg);
#endif
}

Source: src/utilities/math.cuh.

WalkerRNG shows the device/host split and the defaulted constructor: its member functions that call cuRAND are guarded with #if defined(__CUDA_ARCH__) and marked __device__ (no CUDA_CALLABLE), and the constructor is = default:

WalkerRNG() = default;


#if defined(__CUDA_ARCH__)
[[nodiscard]] __device__
real_t rand_uniform() {

Source: src/utilities/random.cuh.

Numerical modes

Section titled “Numerical modes”
  • Must: make precision and approximate-math modes explicit build options.
  • Must: reject incompatible numerical modes during configuration or compilation.
  • Must: test every supported precision/backend combination with tolerances appropriate to that combination.
  • Must: document when an approximate intrinsic changes accuracy guarantees.
  • Should: default to the most conservative supported numerical mode.

VMC defaults to FP64, makes fast approximate math opt-in, and forces FP32 when fast math is selected: CMakeLists.txt. The math header independently rejects the invalid FP64/fast-math combination: src/utilities/math.cuh.

CUDA testing requirements

Section titled “CUDA testing requirements”
  • Must: run the same high-level numerical invariants on CPU and CUDA builds.
  • Must: include problem sizes that are not multiples of the selected block dimensions, exercising bounds guards.
  • Must: test empty, minimum-size, degenerate, and representative large inputs where the API permits them.
  • Should: use a CUDA memory/race analysis tool in CI or a documented release validation job.
  • Should: record the GPU architecture and CUDA toolkit version with benchmark or validation results.