GPU backends

FORGE has three backends: CPU, Metal (Apple Silicon) and CUDA (NVIDIA). They share the same operator classes and solver. Each backend is compiled into its own library, ForgeCommon_cpu, ForgeCommon_metal or ForgeCommon_cuda, with its own copy of every program, and the choice is made at compile time with #ifdef and if constexpr rather than at run time. That is why a GPU build installs the programs twice, and why you pick a backend by choosing which directory is on your PATH.

Only single precision runs on a GPU. double instantiations of the same classes always use the CPU code.

What runs where

Part	CPU	Metal	CUDA
NUFFT (Gnufft<float>)	FFTW and OpenMP gridding	Gridding, deapodization, padding and cropping kernels, FFTs through MPSGraph	Gridding kernels and cuFFT
DFT (Gdft<float>, GdftR2<float>)	OpenMP loops	Metal kernels	CUDA kernels
SENSE coil loop	CPU	Coil weighting with Accelerate, then all coils through one batched GPU NUFFT	GPU, with coil maps uploaded once and kept on the device
Penalty finite differences	CPU	CPU (Accelerate)	CUDA kernels
Solver vector arithmetic (forgeCol)	Armadillo	Accelerate, and Metal for large complex vectors	CUDA kernels, with vectors resident on the GPU for the whole solve
Cartesian FFT (Gfft)	FFTW	Accelerate vDSP	FFTW

The CPU backend uses OpenMP throughout, and on macOS also Apple Accelerate for vector arithmetic.

Metal

Apple Silicon GPUs share memory with the CPU. FORGE takes advantage of this by allocating forgeCol and forgeMat buffers aligned to the 16 KB page size, so Metal can use the same memory without a copy. A whole NUFFT forward or adjoint runs on the GPU with its intermediate buffers kept there.

For vector arithmetic in forgeCol, each operation on float or forgeComplex<float> first asks the Metal layer whether to take it:

forgeCol operator (for example %= or /=)
  -> if constexpr (element type is float or forgeComplex<float>)
    -> pg_metal::try_metal_*(...)
      -> at least PG_METAL_MIN_SIZE elements and a Metal device exists?
           page-aligned buffers -> zero-copy Metal dispatch
           otherwise            -> Metal dispatch with copies
      -> otherwise return false
  -> CPU fallback (Accelerate or a plain loop)

Small vectors stay on the CPU because launching a GPU kernel costs more than the arithmetic. Complex multiply and divide, scalar multiply and real-times-complex multiply go to the GPU; additions, sums, dot products and norms use Accelerate on the CPU. pcSENSE also has a fused multiply-add kernel for its adjoint.

The .metal shaders are compiled at build time and embedded in the library, so there are no shader files to ship.

CUDA

NVIDIA GPUs usually have their own memory, so the CUDA backend avoids round trips to the host. A forgeCol can live on the device (putOnGPU(), getFromGPU(), on_gpu()), and the solver uploads its vectors once and keeps them there for the whole solve. Operators accept device-resident vectors and return device-resident results, so in the inner loop data moves between host and device only when an operator has no GPU path. Temporary allocations use cudaMallocAsync on one stream, which keeps them cheap.

NUFFT pipelines of the same grid shape share one cuFFT plan and a pool of grid scratch buffers, so the many per-shot and per-coil operators in a pcSENSE reconstruction don't each allocate their own.

The library is compiled for compute capabilities 7.5 to 12.0 (Turing to Blackwell). It is tested on a DGX Spark (GB10, unified memory). Batched execution of the per-coil NUFFTs was slower there and is off by default (FORGE_CUDA_BATCHED_OPERATORS); it may help on discrete GPUs.

Batching coils to fit memory

pcSenseTimeSeg needs memory for every coil, shot and time segment at once. On both GPU backends it splits the coils into batches sized to the device's memory budget, processes one batch at a time, and sums the results, so a large reconstruction runs more slowly rather than failing. FORGE_METAL_WORKING_SET_MB and FORGE_CUDA_WORKING_SET_MB set the budget (see Run on a GPU).

Testing the backends against each other

The CPU path is the reference. metal_tests and cuda_tests compare GPU operators and solver results with the CPU versions and with an exact DFT, and check that the GPU paths actually ran. See Run the tests.