4.6. Fine-Grained Synchronization#
Fine-grained synchronization lets a consumer determine when each tile of its input is ready by inspecting the data itself, rather than by waiting on a flag or on the completion of the producer grid. In this chapter, a tile is a subset of the input data that a consumer thread block loads and uses as a unit. Each consumer thread block can therefore use its tile as soon as the tile is ready, while the producer continues working on other tiles. This can reduce the handoff latency in chains of short, dependent kernels.
The protocol marks data that has not been written yet with a reserved bit pattern called a sentinel. The buffer is first initialized with the sentinel. The producer replaces it with final output values whose bit patterns never match the sentinel, and the consumer checks for the sentinel when it loads its input tile. The producer must also follow the store rules of Section 4.6.4.3, so that a tile in which the check finds no sentinel can be treated as complete.
This chapter describes the implementation with TMA data validity reporting, which is available on devices of compute capability 10.7. A bulk-asynchronous copy can check for the sentinel while it transfers data from global to shared memory. It reports the result through a shared memory barrier initialized with the layout::v1 format, which distinguishes copy completion from data readiness. If the copy reports the sentinel, the consumer must issue the copy again; the hardware does not retry automatically.
The protocol can coordinate separate kernels, or thread blocks within a fused or persistent kernel, provided that the producer can make progress while the consumer waits. It also applies between GPUs when the producer writes into a buffer in the consumer GPUâs memory, using memory operations whose scope covers both GPUs (see Section 3.2.3). This chapter covers a producer and a consumer on the same GPU. The example in this chapter runs an initialization kernel, a producer kernel, and a consumer kernel on one GPU, with programmatic dependent launch to let the last two overlap. The performance benefit depends on that overlap, but correctness does not (see Section 4.6.5).
4.6.1. Definitions#
The following terms are used throughout this chapter.
Tile: A tile is a subset of the input data, such as a block of a matrix, that a consumer thread block loads and uses as a unit. This general meaning differs from that of a tile in the CUDA Tile programming model (see Section 1.2.2.3.1).
Sentinel: A sentinel is a reserved bit pattern that marks data that the producer has not written yet. The buffer is initialized with the sentinel, and the producer never stores the sentinel as a result. The report mechanism determines the sentinel (see Section 4.6.4.1).
Report Mechanism:
A report mechanism determines whether and how a bulk-asynchronous copy from global to shared memory inspects the data that it copies. The copy specifies it with a qualifier in PTX or an argument in cuda::ptx. A validity report mechanism is any report mechanism other than .mbarrier::report::disabled. It specifies the sentinel and the granularity at which the copy inspects the source data (see Table 15).
Layout v1:
Layout v1 is the layout::v1 format of the PTX mbarrier object. A barrier in this format stores a payload report and lets its primary and conditional phases advance separately. A copy with a validity report mechanism must be tracked by a layout v1 barrier.
Report-On Operation: A report-on operation records information in the payload report of a shared memory barrier. A copy with a validity report mechanism performs one on the barrier that tracks it when an element that the copy inspects matches the sentinel.
Payload Report: A payload report is information that a layout v1 barrier stores for each primary phase, consisting of a report predicate and a report value. It is zero when the phase starts and becomes non-zero when an operation tracked by the phase performs a report-on operation.
Primary Phase and Conditional Phase:
Every PTX mbarrier object, whatever its layout, has a primary phase and a conditional phase. The primary phase advances when the pending arrival count and the transaction count both reach zero. In a layout v1 barrier, the conditional phase advances with it only if the payload report is zero (see Section 4.6.3.1).
Report Predicate: The report predicate is part of the payload report and is returned by a wait on the primary phase. It is meaningful only when the wait returns true. It is then true if an operation tracked by the completed phase performed a report-on operation, such as a copy that found the sentinel.
4.6.2. Background#
A producer kernel writes data, such as a tensor, to global memory, and a consumer kernel reads it. Table 14 summarizes the ways to order the two kernels.
With plain stream ordering, the consumer is not launched until the producer has completed and all of its writes have been made visible. This is the simplest mechanism and also the coarsest: the consumer waits for the last producer thread block to complete, and the consumerâs launch latency and setup work are fully exposed.
Programmatic dependent launch lets the consumer launch early and run its setup work while the producer is still executing. The dependency is still resolved at the granularity of the whole grid, however: cudaGridDependencySynchronize() in the consumer returns only when the producer grid has completed and all of its memory operations have been made visible.
A flag protocol resolves dependencies between individual thread blocks while both kernels run. The producer writes a tile, executes a device-scope release fence so that the tile is visible before the flag (see Section 3.2.3), and then writes the flag. The consumer polls the flag with acquire semantics and only then loads the tile. The critical path therefore contains the fence plus two dependent memory round trips.
Fine-grained synchronization keeps the per-tile granularity of a flag protocol but removes both the fence and the flag. The consumerâs load that finds no sentinel is also the load that delivers the data, so the handoff costs one memory round trip instead of two. In exchange, the buffer must be filled with the sentinel before every use, and the producer must follow a small set of rules (see Section 4.6.4.3).
Mechanism |
Consumer Waits For |
Granularity |
Requirements |
|---|---|---|---|
Stream ordering |
Completion and memory flush of the producer grid, then launch of the consumer grid |
Grid |
None |
Programmatic dependent launch |
Completion of the whole producer grid and visibility of its writes, observed by |
Grid |
Consumer launched with programmatic stream serialization |
Flag protocol |
A flag written by the producer after a release fence |
Tile |
Producer able to make progress while the consumer waits; one flag per tile |
Fine-grained synchronization |
The data itself no longer holding the sentinel |
Tile |
Producer able to make progress while the consumer waits; sentinel-initialized buffer; producer follows the rules of Section 4.6.4.3 |
4.6.3. TMA Data Validity Reporting#
To use fine-grained synchronization, the consumer loads each tile with a bulk-asynchronous copy from global to shared memory that specifies a validity report mechanism. The copy is either a one-dimensional cp.async.bulk or a cp.async.bulk.tensor that uses a tensor map (see Section 4.13.2.2). If an element that the copy inspects matches the sentinel bit pattern, a report-on operation is performed on the shared memory barrier that tracks the copy. The barrier records this in a payload report that the waiting threads can observe together with the copyâs completion. The data is still written to shared memory; the report only tells the kernel that it must not be used yet.
The .report_mechanism qualifier is supported only on copies from global to shared memory. It requires PTX ISA 9.4 or later and the family-specific sm_107f or the architecture-specific sm_107a compilation target of compute capability 10.7, or a later target of the same family (see Section 5.1.2.2). The cuda::ptx namespace of libcu++ exposes it through additional overloads of cuda::ptx::cp_async_bulk and cuda::ptx::cp_async_bulk_tensor that take a report mechanism argument. Table 15 lists the available mechanisms.
PTX Qualifier |
|
Sentinel |
Inspection |
|---|---|---|---|
|
|
none |
none |
|
|
|
a sampled element per aligned 16 bytes |
|
|
|
a sampled element per aligned 16 bytes |
|
|
|
a sampled element per aligned 16 bytes |
|
|
|
a sampled element per aligned 16 bytes |
|
|
|
every element |
With the per-16-byte mechanisms, the copy compares a sampled element of each aligned 16-byte chunk of the source with the pattern. The producer must therefore write each chunk as a unit (see Section 4.6.4.3). The per_element::ff mechanism compares every element, so it does not depend on the producer writing whole 16-byte chunks.
4.6.3.1. Barrier Layout and Phases#
A copy that uses a validity report mechanism must be tracked by a shared memory barrier initialized with the layout::v1 format of the PTX mbarrier object. Every barrier has a primary phase and a conditional phase. The default format, layout::v0, tracks the pending arrival count and the transaction count described in Section 4.11.6, and its two phases always advance together. Layout v1 additionally stores a payload report for each primary phase and allows the two phases to advance separately:
The primary phase advances when the pending arrival count and the transaction count both reach zero. For a copy tracked by the barrier, primary completion means that all bytes have been written to shared memory.
The conditional phase advances together with the primary phase only if the payload report of the completing phase is zero, that is, if no operation tracked by that phase performed a report-on operation. If a copy reported the sentinel, the primary phase advances but the conditional phase does not. Each primary phase starts with a payload report of zero.
The mbarrier.test_wait and mbarrier.try_wait instructions select which phase to observe with the .phase_type::primary or .phase_type::conditional qualifier. Waiting on the primary phase additionally returns a report predicate that is true when the payload report was non-zero. The predicate only indicates that at least one copy tracked by that phase matched its pattern; it does not identify the copy. Copies whose outcomes must be told apart should therefore be tracked by separate barriers.
The two phases map directly onto the two roles in a consumer:
The thread that is responsible for re-issuing the copy waits on the primary phase. If the report predicate is set, the data was not ready. The thread then arms the barrier again with the expected byte count, re-issues the copy, and waits on the next primary phase with the flipped parity.
The threads that compute on the data wait on the conditional phase. Its parity does not change when a copy reports the sentinel, so these threads wait for the same parity however many times the copy is re-issued. They are released once, when a copy completes with valid data.
In libcu++, the barrier is created with cuda::ptx::mbarrier_init(cuda::ptx::layout_v1, ...) (see the libcu++ documentation of mbarrier_init). The wrappers cuda::ptx::mbarrier_try_wait_parity and cuda::ptx::mbarrier_test_wait_parity take cuda::ptx::mbarrier_phase_primary with a bool& that receives the report predicate, or cuda::ptx::mbarrier_phase_conditional without it. cuda::barrier objects use layout v0 and cannot be used with a validity report mechanism.
Note
A layout v1 barrier accepts at most 511 (29 - 1) expected arrivals per phase, compared with 220 - 1 for layout v0. This is the same layout and the same payload report that compute fabric transport operations use to report errors. The behavior of a program is undefined if it combines fabric operations with copies that use a validity report mechanism.
4.6.4. Protocol#
This section describes the protocol as it is implemented with TMA data validity reporting on devices of compute capability 10.7. The protocol consists of five steps, carried out by the code that manages the buffer, the producer, and the consumer:
Choose a sentinel value that valid data never contains.
Initialize the buffer with the sentinel before the consumer starts to load it.
Produce the data without ever writing the sentinel, and write each location once. With a per-16-byte report mechanism, write each aligned 16-byte chunk as a unit.
Load each tile in the consumer with a bulk-asynchronous copy that reports validity, and re-issue the copy as long as the sentinel is reported.
Re-initialize the buffer before it carries the next round of data.
4.6.4.1. Sentinel Value#
The report mechanism determines the sentinel (see Table 15). For the floating-point formats with a sign bit in that table, the per-16-byte mechanisms use negative zero. Negative zero compares equal to positive zero, so a producer can replace it with positive zero, and only operations that inspect the sign of zero can observe the substitution. The per_element::ff mechanism uses the all-ones pattern, which is the NaN encoding of the 8-bit scale-factor formats ue8m0 and ue5m3. These formats have no sign bit and therefore no negative zero.
The check compares raw bit patterns regardless of the element type, so the 0x80000000 pattern serves any 32-bit type and the 0xFF pattern any 8-bit type. For integer data, the sentinel must be a bit pattern that the producer can be shown never to emit. With the per_element::ff mechanism, the pattern for 4-bit elements is 0xF, and every element is compared with it. Buffers of 4-bit metadata are therefore initialized with whole 0xFF bytes, so that both 4-bit elements of every byte hold the pattern. Table 16 lists a sentinel for each element type.
Element Type |
Sentinel Value |
Bit Pattern |
Element Width |
|---|---|---|---|
32-bit floating point ( |
-0.0 |
|
32 bits |
16-bit floating point ( |
-0.0 |
|
16 bits |
8-bit floating point ( |
-0.0 |
|
8 bits |
4-bit floating point ( |
-0.0 |
|
4 bits |
8-bit scale factors without a sign bit ( |
NaN / all ones |
|
8 bits |
4-bit metadata that never holds |
All ones |
|
4 bits |
Integer data |
A value the producer never emits |
|
32, 16, or 8 bits |
4.6.4.2. Buffer Initialization#
Every location whose data the consumer waits for must be initialized with the sentinel before the producer writes it. The initialization must also be visible to the consumer before the consumer loads the location. Initialization is a normal set of stores, followed by an ordering between the initializing code and the consumer. The first initialization is commonly placed in a separate kernel or in the producer.
In a separate kernel. A separate kernel earlier in the same stream, or an earlier node in the graph, writes the sentinel, and stream ordering or the graph dependency guarantees visibility. This keeps initialization code out of the producer and the consumer and lets one kernel clear several buffers, but it adds a launch and makes the framework responsible for scheduling the clear.
In the producer. The producer writes the sentinel to its own output before it triggers the dependent launch. The threads of a producer thread block write the sentinel to the blockâs output tile and synchronize. One thread then executes a release fence followed by cudaTriggerProgrammaticLaunchCompletion():
// Every thread of the block writes the sentinel to the block's output tile.
write_sentinel_to_output_tile();
__syncthreads(); // all sentinel stores of the block precede the fence and the trigger
if (threadIdx.x == 0) {
cuda::atomic_thread_fence(cuda::memory_order_release, cuda::thread_scope_device);
cudaTriggerProgrammaticLaunchCompletion();
}
// Compute and store the results; they replace the sentinel.
A release fence that precedes griddepcontrol.launch_dependents in program order synchronizes with the start of the dependent grid. This guarantee requires that both kernels run in the same memory synchronization domain and the fence has device scope, or that the fence has system scope. Because the trigger takes effect for the whole thread block at its first execution, the threads that wrote the sentinel synchronize before the trigger, so that every sentinel store is ordered before the fence.
The sentinel stores, the fence, and the trigger delay the consumerâs launch. They also delay the producerâs own work, unless an otherwise idle warp performs them. When one warp writes the sentinel while other warps compute the output values that replace it, the warps must synchronize so that each output store is ordered after the sentinel store to the same location. A separate initialization kernel lets the producer trigger at its start instead, as in the example.
If the sentinel is written with bulk-asynchronous stores instead, the thread that issued them must commit them with cp.async.bulk.commit_group and wait for their completion with cp.async.bulk.wait_group before the block synchronizes. The .read variant of that instruction only waits until the source has been read, which is not sufficient.
If the sentinel is written with ordinary stores and the results with bulk-asynchronous stores, the two kinds of store use different proxies. Each thread then also executes a proxy fence (cuda::ptx::fence_proxy_async) after its sentinel stores, before the block synchronizes (see async proxy).
Each thread block of a persistent producer can perform the work of several thread block indices, for example in a grid-stride loop. It first writes the sentinel to the output tiles of all of these indices. It then synchronizes, executes the release fence, and triggers once, rather than once per index. If a thread block obtains further indices after it has triggered, for example through cluster launch control, their output tiles must already hold the sentinel, for example from a separate initialization kernel.
4.6.4.3. Producer Rules#
The producer must follow three rules: never write the sentinel, write each location once, and, with a per-16-byte report mechanism, write each aligned 16-byte chunk as a unit. The first rule keeps the consumer from waiting indefinitely; the other two let a tile in which the check finds no sentinel be treated as complete.
Canonicalization. A producer that emitted the sentinel as a legitimate value can make the consumer wait indefinitely, so results must be canonicalized before they are stored. With negative zero as the sentinel, a negative-zero result is replaced with positive zero. Adding positive zero with __fadd_rn() does this, as in the example, because negative zero plus positive zero is positive zero in round-to-nearest mode. An equality comparison with 0.0f, followed by a select, also replaces negative zero with positive zero, because the comparison is true for both zeros. For the sm_107f and sm_107a compilation targets, the PTX cvt instruction accepts a .pzo (positive zero only) qualifier. This qualifier performs the same replacement as part of a narrowing floating-point conversion, for example from float to half or to an 8-bit format.
With the NaN pattern as the sentinel, the producer must likewise never store a NaN. A clamp with fmaxf and fminf removes NaN, because these functions return the numeric operand when the other operand is NaN. With an integer sentinel, the producer must never store that value, for example by saturating results away from it.
Write once. Each location must be written once, with its final value. A location that is first written with a partial result and later updated, as in a split-K reduction with atomic adds, does not work with this protocol. The consumer could accept the partial value.
Whole chunks. The producer writes each aligned 16-byte chunk with a single 128-bit store, for example the PTX st.relaxed.gpu.global.b128 instruction that the example wraps in a small helper. A bulk-asynchronous store with the .relaxed.gpu qualifier and the .b128 type also writes each aligned chunk with a single 128-bit operation. With a per-16-byte report mechanism, writing each chunk as a unit is a requirement. The check can inspect as little as one element per chunk, so a chunk assembled from several narrower stores could be reported valid while some of its elements still hold the sentinel.
4.6.4.4. Consumer Check#
The consumer loads each tile with a bulk-asynchronous copy that uses a validity report mechanism, and re-issues the copy as long as the report predicate is set (see Section 4.6.3.1). The threads that wait on the conditional phase with acquire semantics, as in the example, can use the tile as soon as that phase completes, without any further fence. This covers only the data that the copy loaded. If the consumer also depends on other data written by the producer, that data still needs conventional synchronization, such as cudaGridDependencySynchronize() or a release and acquire pair.
4.6.4.5. Buffer Reuse#
Before the buffer can carry the next round of data, it must be re-initialized, and the re-initialization must not race with consumers that are still reading the previous round. Rotating between two or three buffers keeps the clear off the critical path, because one buffer can be cleared while the others are written and read.
The consumer can also re-initialize each tile after it has consumed the tile, for example in persistent or iterative kernels. This is safe only when no other consumer thread block can still be reading the tile, for example when a single consumer thread block reads each tile. The re-initialization must also be ordered before the next producerâs writes and become visible before the next consumer loads the tile, which can require a proxy fence (see async proxy).
4.6.4.6. Partial Tiles and Padding#
Every element that a copy with validity reporting reads must eventually hold a value other than the sentinel, or the consumer can wait indefinitely. Elements that the producer never writes, such as the part of a last tile that extends past the valid data or padding inserted for alignment, can be handled in three ways:
The elements are initialized with a value other than the sentinel.
The producer writes them with a canonical value such as positive zero.
The copy is restricted to the valid data, for example through the bounds of a tensor map.
When the valid data ends inside an aligned 16-byte chunk, for example because a row of a tensor is not a multiple of 16 bytes long, the rest of that chunk needs care. Its bytes beyond the valid data must also eventually hold a value other than the sentinel, even when a tensor-map bound excludes them from the copy. The simplest way is for the producer to write the whole chunk as a unit, with a canonical value beyond the valid data. With a per-16-byte mechanism, initializing only those bytes with a non-sentinel value may not be sufficient, because the chunk could then be accepted before its valid data is written.
A one-dimensional copy with the .ignore_oob qualifier cannot use a validity report mechanism (see cp.async.bulk).
4.6.5. Producer and Consumer Concurrency#
A consumer thread block that waits for a tile cannot make progress until the producer has written that tile. The producer must therefore be able to make progress while consumer thread blocks wait, and the performance benefit appears only when the two kernels overlap. This chapter pairs fine-grained synchronization with programmatic dependent launch, which guarantees the producerâs progress and allows the two kernels to overlap:
The producer calls
cudaTriggerProgrammaticLaunchCompletion()as early as its initialization strategy allows (see Section 4.6.4.2), so that the consumer can be scheduled while the producer is still computing.The consumer is launched with the
cudaLaunchAttributeProgrammaticStreamSerializationattribute, or with a programmatic edge from thecudaGraphKernelNodePortProgrammaticport in a CUDA graph (see Section 4.5.3).The consumer loads the produced data before it calls
cudaGridDependencySynchronize(). That call waits for the entire producer grid, which is the coarse dependency that fine-grained synchronization replaces for the loaded data. The consumer must still make the call, for example after it has consumed the loaded data, where the call is off the critical path of the handoff.
Progress guarantee. The consumer grid is not launched until every producer thread block has either executed the trigger or exited. Every producer thread block is therefore already resident, or has finished, when the first consumer thread block starts to load data. Consumer thread blocks cannot keep any producer thread block from being scheduled.
Separate streams. When the producer and the consumer are in different streams, a programmatic event provides the same progress guarantee. The producer is launched with the cudaLaunchAttributeProgrammaticEvent attribute, recording an event created with cudaEventDisableTiming, and the consumerâs stream waits for that event. When the two streams are captured into a CUDA graph, this dependency becomes a programmatic edge between the two kernel nodes (see Section 4.5.3). If the attributeâs triggerAtBlockStart field is non-zero, the event can trigger as soon as every producer thread block has started. The sentinel must then be written before the producer starts, for example by a separate kernel. The consumer still calls cudaGridDependencySynchronize() after it has consumed the loaded data. Without such a dependency, the two launches are unordered, and a consumer that occupies the GPU first may starve the producer (see Section 5.9.3.1).
Overlap. Whether the two kernels actually overlap remains opportunistic, as it is for programmatic dependent launch in general. The protocol is correct without overlap, because a consumer re-issues its copies until the data is present. Overlap can happen on SMs that the last wave of the producer leaves idle, or on SMs with room for thread blocks of both kernels. The example takes two further steps that make the overlap more likely:
The example loads all three kernels before the sequence is launched, with
cudaFuncGetAttributes(), as recommended in Section 4.9.5. With lazy loading, which is enabled by default, a kernel is otherwise loaded at its first launch, which can serialize kernel execution.The example sets
cudaFuncAttributePreferredSharedMemoryCarveoutto the same value for the producer and the consumer (see Section 3.2.6). The L1 cache and the shared memory of an SM share one unified data cache whose split is configured per kernel, and kernels with different shared memory configurations can serialize behind its reconfiguration. The driver treats the carveout as a hint.
4.6.6. Example#
The following example implements the protocol for float data with three kernels in one stream. An initialization kernel fills the intermediate buffer with the sentinel, a producer kernel doubles an input array into that buffer, and a consumer kernel reverses the elements of each 1024-element tile of the result. Each producer thread writes one 16-byte chunk, and each consumer thread block loads one tile into shared memory with a one-dimensional bulk-asynchronous copy that reports validity.
The sentinel, the canonicalization of results, and the initialization kernel are defined as follows:
// The sentinel is negative zero. Its bit pattern is 0x80000000 for a 32-bit
// float; valid producer output never contains this pattern (see canonicalize()).
constexpr uint32_t SENTINEL_BITS = 0x80000000u;
// The producer must never emit the sentinel. __fadd_rn() adds positive zero in
// round-to-nearest mode, which turns a negative-zero result into positive zero.
// No other value becomes the sentinel, and a NaN stays a NaN.
__device__ __forceinline__ float canonicalize(float v)
{
return __fadd_rn(v, 0.0f);
}
// Fill a buffer with the sentinel, using 16-byte stores.
__global__ void init_sentinel_kernel(float4 *buf, int num_chunks)
{
const float s = __uint_as_float(SENTINEL_BITS);
const float4 sentinel = make_float4(s, s, s, s);
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < num_chunks; i += gridDim.x * blockDim.x)
{
buf[i] = sentinel;
}
}
Because a separate kernel writes the sentinel (see Section 4.6.4.2), nothing has to precede the trigger, and the producer triggers the dependent launch at its start. It then canonicalizes its results and writes each chunk once with a single relaxed 128-bit store at device scope:
// Relaxed, device-scope 128-bit store of one aligned 16-byte chunk.
__device__ __forceinline__ void store_chunk_relaxed_gpu(float4 *p, const float4 &v)
{
__int128 bits;
memcpy(&bits, &v, sizeof(bits));
asm volatile("st.relaxed.gpu.global.b128 [%0], %1;" : : "l"(p), "q"(bits) : "memory");
}
__global__ void producer_kernel(const float4 *__restrict__ in, float4 *__restrict__ out)
{
// Allow the consumer to launch as early as possible. The sentinel was
// written by a separate kernel, so nothing has to precede the trigger here.
cudaTriggerProgrammaticLaunchCompletion();
const int chunk = blockIdx.x * blockDim.x + threadIdx.x;
float4 v = in[chunk];
v.x = canonicalize(2.0f * v.x);
v.y = canonicalize(2.0f * v.y);
v.z = canonicalize(2.0f * v.z);
v.w = canonicalize(2.0f * v.w);
// Each aligned 16-byte chunk is written once, with a single 128-bit store,
// as the per-16-byte validity check requires.
store_chunk_relaxed_gpu(&out[chunk], v);
}
The consumerâs copy is written in inline PTX with the .relaxed.gpu qualifier and the .b128 type, and the barrier operations use the cuda::ptx wrappers of libcu++. Like the producerâs stores, the copy therefore accesses global memory with strong, naturally aligned 128-bit operations at device scope. A strong write through the generic proxy eventually becomes observable by a strong access to the same location through the async proxy, when both are at device scope. A relaxed operation is a strong operation with neither release nor acquire semantics.
In the consumer, thread 0 initializes a layout v1 barrier, issues the copy with the 32-bit negative-zero report mechanism, and re-issues it as long as the report predicate indicates that the tile was not ready. All threads wait for the conditional phase with parity 0 before they read the tile, and they call cudaGridDependencySynchronize() after they have written their results:
// Relaxed, device-scope bulk-asynchronous copy from global to shared memory
// that compares a sampled element of each aligned 16-byte chunk with the
// 32-bit negative-zero sentinel and reports whether any matched. The copy
// accesses global memory with 128-bit operations and completes on the
// layout v1 barrier `bar`.
__device__ __forceinline__ void copy_tile_report_validity(void *smem_dst, const void *gmem_src, uint32_t bytes,
uint64_t *bar)
{
asm volatile("cp.async.bulk.relaxed.gpu.shared::cta.global.mbarrier::complete_tx::bytes"
".mbarrier::report::validity::per_16bytes::80000000.b128 [%0], [%1], %2, [%3];"
:
: "r"(static_cast<uint32_t>(__cvta_generic_to_shared(smem_dst))), "l"(gmem_src), "r"(bytes),
"r"(static_cast<uint32_t>(__cvta_generic_to_shared(bar)))
: "memory");
}
__global__ void consumer_kernel(const float *produced, float *__restrict__ result)
{
// Source and destination of a one-dimensional bulk copy must be 16-byte
// aligned and the size a multiple of 16 bytes.
__shared__ alignas(128) float tile[TILE_ELEMS];
// A copy with a validity report mechanism must be tracked by a layout v1
// barrier, which stores the payload report.
__shared__ alignas(8) uint64_t bar;
if (threadIdx.x == 0)
{
ptx::mbarrier_init(ptx::layout_v1, &bar, 1); // one arrival per phase: the issuing thread
}
__syncthreads();
const float *src = produced + blockIdx.x * TILE_ELEMS;
if (threadIdx.x == 0)
{
uint32_t primary_parity = 0;
for (;;)
{
// Arm the current primary phase with the expected byte count, then
// issue the copy with validity reporting enabled.
ptx::mbarrier_arrive_expect_tx(ptx::sem_release, ptx::scope_cta, ptx::space_shared, &bar, sizeof(tile));
copy_tile_report_validity(tile, src, sizeof(tile), &bar);
// Wait for the *primary* phase, which completes once all bytes have
// landed. The report predicate tells whether a sampled element still
// held the sentinel.
bool sentinel_seen = false;
while (!ptx::mbarrier_try_wait_parity(ptx::mbarrier_phase_primary, ptx::sem_acquire, ptx::scope_cta,
sentinel_seen, &bar, primary_parity))
{
}
if (!sentinel_seen)
{
break; // valid data: the conditional phase has advanced as well
}
// Not ready: some of the producer's stores to this tile were not yet
// visible to the copy. The primary phase completed anyway, so the next
// copy is tracked by the following primary phase.
primary_parity ^= 1;
}
}
// Every thread waits for the *conditional* phase. It advances only when a
// copy completes with no sentinel reported, so the threads wait for parity 0
// however many times the copy is re-issued.
while (!ptx::mbarrier_try_wait_parity(ptx::mbarrier_phase_conditional, ptx::sem_acquire, ptx::scope_cta, &bar, 0))
{
}
// The tile in shared memory is complete and valid. Consume it: each thread
// writes four elements of the reversed tile.
float *out = result + blockIdx.x * TILE_ELEMS;
for (int i = 4 * threadIdx.x; i < 4 * threadIdx.x + 4; ++i)
{
out[i] = tile[TILE_ELEMS - 1 - i];
}
// A dependent grid must still call cudaGridDependencySynchronize(). Calling
// it here, after the loaded data has been consumed, keeps the wait off the
// critical path of the handoff.
cudaGridDependencySynchronize();
}
Arrival count. The arrival count of the barrier is one: only the issuing thread arrives, through mbarrier_arrive_expect_tx, and it does so once per primary phase, including for every re-issued copy.
On the host, the three kernels are preloaded, and the producer and the consumer are given the same shared memory carveout (see Section 4.6.5). The three kernels are then launched into one stream. Stream ordering makes the sentinel visible before the producer starts, and only the consumer is launched with programmatic stream serialization, so that it can start while the producer is running:
// Load all three kernels before the sequence is launched. With lazy module
// loading (the default), a kernel is otherwise loaded at its first launch,
// which can serialize kernel execution and keep the producer and the
// consumer from overlapping.
cudaFuncAttributes func_attr;
CUDA_CHECK(cudaFuncGetAttributes(&func_attr, init_sentinel_kernel));
CUDA_CHECK(cudaFuncGetAttributes(&func_attr, producer_kernel));
CUDA_CHECK(cudaFuncGetAttributes(&func_attr, consumer_kernel));
// The L1 cache and the shared memory of an SM share one unified data cache,
// and the driver configures the split per kernel. Kernels with different
// shared memory configurations can serialize behind the reconfiguration.
// The producer uses no shared memory and the consumer does, so both request
// the same carveout.
CUDA_CHECK(cudaFuncSetAttribute(producer_kernel, cudaFuncAttributePreferredSharedMemoryCarveout,
cudaSharedmemCarveoutMaxShared));
CUDA_CHECK(cudaFuncSetAttribute(consumer_kernel, cudaFuncAttributePreferredSharedMemoryCarveout,
cudaSharedmemCarveoutMaxShared));
// 1. Initialize the intermediate buffer with the sentinel. Stream order
// guarantees this completes, and is visible, before the producer starts.
init_sentinel_kernel<<<NUM_TILES, THREADS, 0, stream>>>(d_mid, NUM_ELEMS / 4);
// 2. Launch the producer normally.
producer_kernel<<<NUM_TILES, THREADS, 0, stream>>>(d_in, d_mid);
// 3. Launch the consumer with programmatic dependent launch so that it can
// start while the producer is still running.
cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attribute[0].val.programmaticStreamSerializationAllowed = 1;
cudaLaunchConfig_t config = {};
config.gridDim = NUM_TILES;
config.blockDim = THREADS;
config.stream = stream;
config.attrs = attribute;
config.numAttrs = 1;
CUDA_CHECK(cudaLaunchKernelEx(&config, consumer_kernel, (const float *)d_mid, d_out));
4.6.7. Warp-Specialized Kernels#
In the example, one thread issues the copy, waits on the primary phase, and re-issues the copy as long as the sentinel is reported. A warp-specialized kernel with a multi-stage pipeline (see Section 4.11.7) can instead add a retry warp:
The load warp issues the copies of each stage, as it would without fine-grained synchronization.
The retry warp waits on the primary phase of each stage, tracking a separate parity for each stage. When a stageâs report predicate is set, the retry warp re-issues all copies for that stage.
The compute warps only wait on the conditional phase. Only the warp that issues or re-issues a copy arrives on the barrier that tracks it.
When a stageâs report predicate is not set, the retry warp also arrives on the barrier that signals that the stage is ready to be filled (a
consumer_releaseoperation, see Section 4.12). Without this arrival, the load warp could refill the stage before the retry warp has checked the previous copies, and the retry warp could miss a report.
With the .cta_group::2 qualifier of cp.async.bulk.tensor, one barrier can track copies into both thread blocks of a CTA pair in a cluster, but only the thread block that holds the barrier can wait on it. If the other thread block re-issues its own copies or uses the data, it must be told when the copies have completed and whether the sentinel was reported, for example through distributed shared memory.
4.6.8. Guidelines#
The following points cover when the technique pays off and common sources of errors.
Applicability. Fine-grained synchronization can shorten the handoff between kernels. It can benefit chains of short kernels whose consumer thread blocks can start while the producerâs last thread blocks are still running. Long, throughput-bound kernels can gain little, and they still pay for the initialization.
Leaked sentinels. The consumer can hang if a location that its copies depend on keeps the sentinel, for example unwritten padding (see Section 4.6.4.6) or a NaN result when NaN is the sentinel (see Section 4.6.4.3). While debugging, a retry loop that traps after a generous number of re-issues turns such a hang into an error that the host can observe.
Missing initialization. A buffer that is not initialized with the sentinel before each use can cause wrong results instead of a hang. The copy reports nothing for stale data, and the consumer accepts it as ready.
