100 DAYS OF GPU Course preview
Day 58Module 6
draft

Programmatic dependent launch

Programmatic dependent launch, or PDL, needs compute capability 9.0 or newer. You can still complete the timeline exercise from the shipped report if you do not have a compatible GPU.

Day 48 measured a kernel launch, and CUDA graphs reduced its cost. A dependency still adds a wait: in one stream, kernel B's first instruction waits for kernel A's last block to retire, even when B's opening work reads nothing A wrote.

PDL moves that wait from the kernel boundary to the first line that needs the data. This lesson adds PDL to a two-kernel pipeline and measures the overlap on a timeline.

Where the serialized gap comes from

Stream order applies at grid boundaries. Kernel B cannot begin until every block of kernel A has finished. This rule may enforce more waiting than the data dependency needs.

A's blocks do not all finish at once. B may be able to compute indices, load its own inputs, and fill registers before it reads A's output.

PDL splits the dependency between two calls. The primary kernel calls cudaTriggerProgrammaticLaunchCompletion() when it is far enough along that the dependent kernel may be scheduled. The secondary kernel calls cudaGridDependencySynchronize() at the point where it first reads the primary's output.

The driver "can launch the secondary kernel when all primary thread blocks have launched and executed cudaTriggerProgrammaticLaunchCompletion", and the secondary's wait blocks "until all primary kernels the secondary kernel is dependent on have completed and flushed results to global memory" (https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html , checked 2026-09-01). A primary that never calls the trigger still works: the trigger implicitly fires when its last block exits, which is exactly the old stream rule.

The host opts in per launch, with the attribute shown in the code section, on the same cudaLaunchKernelEx that day 28 used for cooperative launches. What PDL may allow is the secondary's prologue running while the last primary blocks finish. The guide states that the concurrency "is opportunistic and not guaranteed" because the secondary's blocks still need free SMs.

Diagram: one dependent pair, three schedules. Three horizontal bands, each a GPU timeline with a bar for kernel A above a bar for kernel B; a dashed line marks where B first reads A's output. Band 1, stream order: B starts only after A's ragged tail fully retires. Caption "two kernels, one gap: B's independent prologue waits too." Band 2, PDL with free SMs: B's bar starts under A's tail; the dashed read line sits after A ends. Caption "the prologue hides under the tail; the read still waits." Band 3, PDL on a saturated card: B's bar starts where A's blocks begin retiring, overlap near zero. Caption "opportunistic: no free SM, no overlap." Alt text: "Three schedules of two dependent kernels. Stream order serializes them fully. Programmatic dependent launch slides the second kernel's prologue under the first one's tail, unless zero SMs are free."

The trigger does not make memory visible

A __syncthreads() or a recorded event orders execution and makes memory visible. PDL assigns those jobs to separate calls. Treating the trigger as both calls creates a data race.

The trigger only permits the driver to schedule the dependent kernel. It says nothing about which writes have reached global memory. The cudaGridDependencySynchronize() call waits for the primary grid to complete and flush its writes.

Moving the trigger earlier only changes scheduling if every read of the primary's output remains after the wait. Reading that output before the wait creates a data race with no error message.

One attribute bit, two schedules

Full program in code/day58-pdl/pdl_pipeline.cu. Three safeguards protect the measurement.

One launch helper, one differing bit. Both timed modes run the same two kernels through the same code path; the baseline sets programmaticStreamSerializationAllowed to 0 and the PDL mode sets it to 1. Nothing else differs, so nothing else can explain a gap.

Bit-identical outputs before any timing. The baseline is checked against a CPU reference, then the PDL run must memcmp equal to the baseline. The attribute changes scheduling and must change nothing else; a mismatch is a failed run, not a footnote.

Events time it, NVTX names it. The chain of 64 pairs is timed with CUDA events behind a warm-up that covers both kernels, and the verify, baseline and pdl ranges make the day 41 timeline readable.

The primary computes mid, triggers, then does epilogue work the secondary never reads:

// The primary kernel. One thread owns one element. Consecutive threads read
// and write consecutive elements, so every access coalesces; nothing here
// is memory-clever on purpose.
//
// Phase 1 computes mid[i], which the secondary kernel will read. The
// trigger then tells the driver the dependent launch may begin. Phase 2
// writes aux[i], work the secondary never reads, so it can overlap the
// secondary's prologue. The trigger promises nothing about memory: the
// secondary's cudaGridDependencySynchronize() is what makes mid visible.
__global__ void produceStage(const float* __restrict__ x,
                             float* __restrict__ mid, float* __restrict__ aux,
                             size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;

    if (i < n) {
        mid[i] = fmaChain(x[i], kMulProduce, kAddProduce);
    }

    cudaTriggerProgrammaticLaunchCompletion();

    if (i < n) {
        aux[i] = fmaChain(x[i], kMulEpilogue, kAddEpilogue);
    }
}

The secondary starts with a prologue that only touches its own input, and waits before its first read of mid:

// The secondary kernel. Same launch shape and access pattern as the
// primary. The prologue depends only on y, so it is legal before the
// synchronize; the read of mid[i] is not, and sits after it. Launched
// without the PDL attribute this kernel is still correct: the wait finds
// the primary already finished and costs nothing.
__global__ void combineStage(const float* __restrict__ mid,
                             const float* __restrict__ y,
                             float* __restrict__ out, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;

    float p = 0.0f;
    if (i < n) {
        p = fmaChain(y[i], kMulPrologue, kAddPrologue);
    }

    cudaGridDependencySynchronize();

    if (i < n) {
        out[i] = mid[i] + p;
    }
}

The host side is one attribute on the launch configuration:

    cudaLaunchAttribute attrs[1];
    attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
    attrs[0].val.programmaticStreamSerializationAllowed = allowPdl;

    cudaLaunchConfig_t cfg = {};
    cfg.gridDim = dim3(static_cast<unsigned int>(blocks));
    cfg.blockDim = dim3(kThreadsPerBlock);
    cfg.dynamicSmemBytes = 0;
    cfg.stream = stream;
    cfg.attrs = attrs;
    cfg.numAttrs = 1;

Note. The workload is a synthetic chain of dependent FMAs, chosen so each phase's cost is knowable and the CPU reference matches exactly. Real pipelines have the same shape with better excuses: an attention kernel loading weights, a GEMM epilogue, a decoder step's setup. The schedule is what this program measures, not the arithmetic.

Results

Not yet run. Verification needs a compute capability 9.0 or newer GPU. The transcript, day58-pdl.nsys-rep capture, and CSV exports will go in content/profiles/programmatic-dependent-launch/. Until then, the front matter stays draft with every hardware field null.

Mode Total (ms) Per pair (us) Ratio to baseline
baseline ? ? ?
pdl ? ? ?

What a compatible run should show

  1. The PDL loop is faster, and by less than the prologue's share of a pair. The bet: between 2 and 25 percent. The prologue and epilogue are each roughly a third of a pair's arithmetic, and both grids fill the GPU, so most overlap can occur only as the last primary blocks finish. A saving above the prologue share would reject this explanation.
  2. The trace shows the overlap directly. In cuda_gpu_trace, within the pdl range, combineStage's start timestamp precedes the end of the produceStage launched just before it, and in the baseline range it never does.
  3. The memcmp gate passes. If it fails, the lesson's premise fails with it. Check the wait before the trigger.
  4. Failure condition. If the two modes tie within noise, either PDL never engaged (the trace shows a positive gap everywhere) or the card had no free resources at the trigger (the gap shrinks but the total does not). The trace tells those two apart; the event timer alone cannot.

Compare the ratio rather than the measured milliseconds. SM counts, clocks, and launch latency differ by GPU. The trace must show whether the prologue ran before the primary finished.

Run it yourself

Use a compute capability 9.0 or newer GPU. The -arch=sm_90 build embeds PTX that can compile for newer architectures at run time. An sm_90a build cannot, because it is an architecture-specific target.

If you cannot run the program, start at /setup/learn-cuda-without-a-gpu. The build and profile commands are in the README. Current paid hardware notes remain in FACT-SHEET.md.

Exercise

Move the trigger from after phase 1 to the first line of produceStage, rebuild, and re-measure both modes. Then compare cuda_gpu_trace gaps: for each pair, subtract the primary's end timestamp from the secondary's start timestamp, in both NVTX ranges, before and after your change.

Time: 30 to 45 minutes. Submit: the two pdl/baseline ratios, the sign of the median gap in each range, and one sentence on why the early trigger did or did not pay.

Check: measured, as ratios, because absolute times differ by card. The shipped CSVs support the same work without hardware: a negative median gap inside the pdl range is PDL engaging; a positive one everywhere means it never did. If your early-trigger build changes the memcmp result, you have moved a read above the wait somewhere, and the program says so and exits nonzero.

Hint 1

Each of the two device calls makes one promise. Which call guarantees the values in mid are readable, and does moving the other call change that guarantee at all?

Hint 2

cuda_gpu_trace gives every kernel a start and a duration. The gap you want is the next kernel's start minus this kernel's start plus duration. And ask where the overlap physically runs: if every SM is still occupied by primary blocks, what does an earlier trigger actually buy?

Solution

Moving the trigger to the top is safe because visibility comes from cudaGridDependencySynchronize(), which waits for the whole primary to complete and flush regardless of when the trigger fired. The ratio can only improve or stay flat.

Both grids fill the GPU, so secondary blocks cannot run until primary blocks start retiring. Most of the extra time before the wait is spent queued. The gap column still moves because launching earlier removes launch latency from the critical path even when SM time does not change.

The trigger controls scheduling, and the wait controls data visibility. Call the trigger as early as the algorithm allows. Read the primary's output only after the wait.

Pitfalls

You built with -arch=sm_90a and a Blackwell card refuses to run it. Architecture-specific targets do not JIT forward; that build loads on Hopper and nothing else. PDL needs nothing from sm_90a, so build -arch=sm_90 and let the embedded PTX carry it to CC 10.x and 12.x. Day 78 maps which feature needs which target.

Both kernels have the device calls, and nothing overlaps. The attribute is on the launch, not in the kernel. A secondary launched with plain <<<>>>, or through a config whose attribute value is 0, serializes silently: no error, no warning, correct answers. The timeline is the only witness, which is why the exercise reads cuda_gpu_trace instead of trusting the source.

You read the primary's output right after the trigger. The trigger only affects scheduling and does not make memory visible. A read of mid above the wait races with the primary's stores, and on a lightly loaded card it will pass for weeks. Only the wait makes those writes visible; day 62 covers tools that detect this race.

You expected the whole secondary to overlap. Only the code above cudaGridDependencySynchronize() can run early, and even that concurrency "is opportunistic and not guaranteed" (the programming guide's words). PDL can reduce launch overhead and overlap work with the last primary blocks. If the dependent part dominates the kernel, restructure so more of it is independent, or fuse instead.

Both grids saturate the card, so PDL measures zero. Secondary blocks need somewhere to run. When the primary occupies every SM until its last blocks finish, the overlap is limited to that final period. This is the diagram's third band and the reason prediction 1 is bounded.

You timed one pair with a host clock. Async launches return immediately; a wall clock around them measures the host, and reading it before a sync measures even less. The program times 64-pair chains with events behind a warm-up, and Nsight Systems explains what the events cannot.

Go deeper

Next

Day 59 adds the CPU side of the pipeline: a producer thread feeding a stream while the GPU consumes, which is the same overlap question asked across the PCIe bus instead of between two grids. Later, day 77 reuses this page's cudaLaunchKernelEx attribute machinery to launch thread block clusters, where the "who may start before whom" question moves inside the grid.