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.

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