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::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::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#

As of the 2026-06-25 uniformity wave, TRAILING_SYNC is on every public op that has a separable trailing barrier — all of L1 (reduce/dot/elementwise/ axpy/copy/scal/swap/clip/set_const/ident/transpose/ infnorm/asum/nrm2 + the _fast/_lowmem reduction variants), L2 (gemv/ger + gemv_strided/gemv_segmented/gemv_reduced), and L3 (gemm + gemm_strided/gemm_batched_indexed/gemm_reduced/syrk/ syr2k/syrk_reduced/tensor/congruence) — plus the cuBLASDx-backed glass::nvidia:: paths.

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

Documented exceptions (the flag is intentionally absent):

  • Algorithm-terminated-by-a-barrier ops — potrf / inv / trsm / trsv / posv / ldlt. Their final step is itself a barrier inside the factorization loop, so there is no separable trailing barrier to gate; they always end synced.

  • ``glass::warp::`` lockstep ops — a single warp runs in lockstep, so a trailing “sync” is a no-op (__syncwarp at most); the flag would be vacuous.

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