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). The launch thread count must be a multiple of 32 (the inner dot uses a warp reduction).

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|.

Shared scratch: pass s_mem of pcg_scratch_bytes<T,state_size,knot_points>(threads) elements (5 padded vectors + the warp-dot scratch). Five scalars live in static __shared__. Requires blockDim.x be a multiple of 32 (the warp dot).

Functions

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

Shared-memory element count needed by glass::pcg.

Host- or device-callable. Multiply by sizeof(T) for the dynamic-shared-mem bytes to pass at launch. 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 – Shared scratch of pcg_scratch_bytes<T,...>(blockDim.x) elements.

  • 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).