Trailing Synchronization (TRAILING_SYNC)#

Uniformity rule (project-wide design decision): every public GLASS op takes a final boolean template parameter ``TRAILING_SYNC`` that defaults to ``true``, on both the glass:: and glass::cgrps:: surfaces. It controls whether the function ends on a barrier (__syncthreads() for the block surface, cooperative_groups::sync for the cgrps surface) so the result is valid for ALL threads. One consistent mental model: a GLASS op acts like a standalone kernel unless you opt out.

Default vs. opt-out#

// Default — function returns with all threads at a block-wide barrier.
// Safe to read the result from any thread in the block immediately after.
glass::nvidia::block::gemm_strided_batched_1d<float, 4, 4, 4, BATCH, TC>(
    1.f, A, B, 0.f, C);

// Opt-out — caller is responsible for syncing before reading any output
// not written by the current thread. Pass false when fusing the GEMM with
// subsequent block-wide work that ALREADY does its own barrier (e.g. a
// parallel_loop that begins with __syncthreads()), so two back-to-back
// syncs collapse into one.
glass::nvidia::block::gemm_strided_batched_1d<
    float, 4, 4, 4, BATCH, TC,
    /*B_STRIDE=*/N*K, /*C_STRIDE=*/M*K,
    layout::col_major, layout::col_major, layout::col_major,
    /*TRAILING_SYNC=*/false>(1.f, A, B, 0.f, C);
__syncthreads();  // emit your own here, fused with any other barrier

The default of true makes the common case correct without thinking — GLASS functions act as if they were standalone kernels. The opt-out exists for hot kernels (e.g. GRiD’s end_effector_pose_gradient_inner) that chain a GEMM with a SIMT parallel_loop and want to collapse the two syncs into one.

Where it applies#

Every public block-scope operation accepts TRAILING_SYNC. For operations with a separable final barrier—including the factor/solve and eigensolver families—false genuinely elides it. For pivoted or multi-exit algorithms whose last barrier is fused into the final algorithm step (currently ldlt/ldlt_solve, getrf, and inv_pivoted), the parameter is accepted for interface uniformity but documented as a no-op at the declaration. The compile-time single-RHS posv<T,N> and potrs<T,N> spellings are the one template-ambiguity exception; use their NRHS=1 overload to select TRAILING_SYNC=false.

Interior barriers (between algorithm phases) are required for correctness and are never gated — only the final trailing barrier is.

glass::thread:: operations do not expose the knob at all — a single thread has no cooperating peers to synchronize with. glass::warp:: operations that end on a warp-level publish point DO take TRAILING_SYNC and gate their final __syncwarp() on it (gemm_reduced, gemv_reduced, syrk_reduced, tensor_contract, axpy_strided, copy_strided, banded load_block/store_block, among others); purely lockstep/shuffle-based warp ops need no trailing barrier and omit the parameter. warp::riccati_gain accepts it for interface uniformity but documents it as a no-op — the composed warp solver always ends at a sync boundary.

Pointer qualifiers and in-place operands#

Callers may pass __restrict__ pointers into GLASS, including into the in-place solves that read back what the caller or an earlier step wrote. GLASS does not depend on the qualifier being absent: every scalar one thread publishes for the others to consume after a barrier is reloaded through cooperative_detail::load (a volatile load in the cooperating scopes, an ordinary load in the thread tier), because an optimizing compiler may otherwise forward the caller’s stale register value across the barrier when the pointer is restricted (issue #44, 2026-10). The reload is not a substitute for synchronization: the barriers and the TRAILING_SYNC ownership rules above remain the contract, and restricted pointers must still not alias one another. The restricted-caller regression (test/test_restrict.py) checks this contract for the block primitives on every receipt, natively and through the PTX-JIT lane.

GLASS declares no __restrict__ on its own parameters. The qualifier was on the gemm and symm operands until 2026-10-31, when a read operand selected at runtime between a thread-local array and shared memory made nvcc (11.4 and 13.2, on sm_72, sm_87 and sm_120) drop the inlined writer’s stores to that array, and the single-thread product returned nondeterministic values. The program was well defined, so the only repair on the library side was to stop supplying the qualifier. Removing it changed no instruction count in the measured gemm kernels, so nothing was traded away. The selected-operand regression (test/test_select.py) keeps that case fixed in both precisions.

Testing#

Tests for the surface live at test/cuda/test_trailing_sync.cu + test/test_trailing_sync.py. They verify that:

  1. Both TRAILING_SYNC=true and =false specializations compile and link.

  2. The two variants produce numerically identical output when the =false caller emits its own trailing __syncthreads().

The cuBLASDx-backed L3 case is covered (gated on GLASS_BENCH_CUBLASDX / MATHDX_ROOT).