Preconditioned Conjugate Gradient (glass::pcg)#

A single-block preconditioned conjugate gradient solver for a block-tridiagonal symmetric-positive-definite system S x = b, with a block-tridiagonal preconditioner Pinv applied as z = Pinv·r. One CUDA block solves one system (launch one block per independent solve); it uses only __syncthreads() — no cooperative groups.

S and Pinv use the [L|D|R] block-tridiagonal layout of Block-tridiagonal Ops (glass::bdmv / glass::bdsv), and all vectors use the same padded (knot_points + 2)·state_size layout. Size the dynamic shared memory with glass::pcg_scratch_bytes<T, state_size, knot_points>(threads); it returns bytes, ready to pass as the launch argument. Any thread count is legal (the fast dot reduction bounds its shuffle mask to the active lanes of a ragged last warp); multiples of 32 are the fast path and reproduce the historical results bit-for-bit.

See Block-tridiagonal Solves for the layout and a worked walkthrough.

Block-wide preconditioned conjugate gradient (glass::pcg).

Solves a single block-tridiagonal SPD system S x = b inside ONE CUDA block, with a block-tridiagonal preconditioner Pinv applied as z = Pinv r. This is the single-block (one-block-per-problem) PCG — launch one block per independent solve. It uses only __syncthreads() (no cooperative groups); for the cooperative grid-wide variant see the backlog glass::cgrps::grid.

Layouts (see glass::bdmv): S and Pinv are block-tridiagonal [L|D|R] row-major strips; all vectors use the padded layout (knot_points + 2) * state_size (one state_size pad block on each end). The caller seeds x with an initial guess (zeros are fine).

Convergence is tested on the preconditioned residual rho = rᵀ z: |rho| < abs_tol + rel_tol * |rho_init|.

Allocate pcg_scratch_bytes<T,state_size,knot_points>(threads) bytes of dynamic shared memory and pass its base as s_mem (5 padded vectors + the warp-dot scratch). Five scalars live in static __shared__. Any launch thread count is legal (the warp-dot reduction bounds its shuffle mask to the active lanes of a ragged last warp); multiples of 32 keep the best throughput and are bit-identical to the historical results. Like dot_fast itself, results are reduction-order-dependent across DIFFERENT thread counts (each count is individually deterministic).

namespace glass

Matrix memory layout for the cuBLASDx-backed glass::nvidia::block:: wrappers.

Maps directly to cublasdx::Arrangement<>: col_major (Fortran/cuBLAS default) and row_major (C-style). Passed per matrix as the LA/LB/LC template arguments of gemm/gemv/row_strided_*. The numeric values (0/1) are part of the public ABI: the DEFINE_NVIDIA_*_LAYOUT* macros take integer literals and static_cast them back to this enum in their specializations.

namespace block

Functions

template<typename T, uint32_t state_size, uint32_t knot_points>
inline constexpr std::size_t pcg_scratch_bytes(uint32_t threads)

Dynamic shared-memory byte count needed by glass::pcg.

Host- or device-callable. Pass the return value directly as the kernel’s dynamic-shared-memory byte count. It covers the 5 padded work vectors plus the warp-dot scratch (ceil(threads/32)); the 5 scalars are static __shared__.

Template Parameters:

  • T – Scalar type.

  • state_size – Block dimension.

  • knot_points – Number of block-rows.

Parameters:

threads – Launch thread count (blockDim.x).

Returns:

Bytes of dynamic shared memory required.

template<typename T, uint32_t state_size, uint32_t knot_points, bool TRAILING_SYNC = true>
void pcg(T *x, T *S, T *Pinv, T *b, T *s_mem, uint32_t max_iters, T rel_tol, T abs_tol, uint32_t *iters)

Solve S x = b by preconditioned conjugate gradient, one block (glass::pcg).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • state_size – Block dimension (= BlockSize of the banded layout).

  • knot_points – Number of block-rows.

Parameters:

  • x – In/out padded solution, length (knot_points+2)*state_size (seed with an initial guess; result written back here).

  • S – Block-tridiagonal SPD system, [L|D|R] row-major strips.

  • Pinv – Block-tridiagonal preconditioner, [L|D|R] row-major strips.

  • b – Padded right-hand side.

  • s_mem – Base of a dynamic-shared allocation of pcg_scratch_bytes<T,...>(blockDim.x) bytes.

  • max_iters – Maximum CG iterations.

  • rel_tol – Relative tolerance on the preconditioned residual.

  • abs_tol – Absolute tolerance on the preconditioned residual.

  • iters – Output: iteration count written by thread 0 (may be null-safe only if the caller guarantees a valid pointer; pass a valid device address).