Historical supporting material. The current article uses the single RMS-scaled softmax workload. This page retains an earlier experiment or claim audit.
Agent-loop discussion — October 19, 2026
The two closing discussion sentences distinguish the engineering cost of exploring implementations from their execution cost. KernelBench reports improved kernel-generation results with iterative execution and profiling feedback; extending that mechanism to megakernel schedules is explicitly an inference, not a demonstrated megakernel result. The paragraph also acknowledges that smaller fused kernels benefit from the same tooling and retains end-to-end comparison as the deciding evidence. No new experiment or general megakernel speedup is claimed; the measurements and figures are unchanged.
Factual and narrative review — October 19, 2026
The latest pass checks the current three-part article for a mid-level engineering audience. Six primary sources were consulted again: CUDA Graphs, L2 caching, Cooperative Groups, Programmatic Dependent Launch, Hazy Research's Llama megakernel, and Triton's stable-softmax tutorial. The captured CUDA implementation, timing harness, compiler output, graph metadata, and saved validation were inspected directly.
The machine-readable review records the manuscript hash before the agent-loop addition, every quoted latency, 12 freshly recomputed timing cases, all 24 dependency traces, graph node counts, resource metadata, figure hashes, and source URLs. All 1,764 timing records match saved summaries; 1,372 local raw artifacts match the full-run manifest, and the 369 remote-only files are explicitly accounted for. The previously saved 360 output and 768 intermediate checks reconcile with the case reports and unchanged tolerances. Large binary arrays were not downloaded or numerically revalidated in this pass; no GPU rerun was necessary.
Scroll horizontally to see the full table.
| Review finding | Final wording or treatment |
|---|---|
| One host call is easy to confuse with one GPU kernel | The opening now poses that exact distinction; replay retains the four supplied kernels. |
| Timing scope appeared after the first result | Host completion, synchronization, warmup/setup exclusions, and the 21-round scope now precede the numbers. |
| cuDNN comparisons needed an explicit submission mode | Both quoted cuDNN values are graph replay: 56.08 µs at 4096×4096 and 7.98 µs at 4096×128. No cross-mode comparison is implied. |
| Local fusion needed a clearer bridge to megakernels | Section 3 moves from one block per row to consumers with cross-block producers. |
| Combining tile exponential sums requires rescaling | B now gives the stable combination formula z = sum_t(z_t * exp(m_t-m)), matching the captured CUDA source and the independent reference. |
| Ready does not imply immediately scheduled | The queue discussion distinguishes dependency eligibility from worker availability and contention. |
| The queued graph contains two kernels | Queue initialization and the persistent kernel are explicitly included in timing. This is not presented as a one-kernel end-to-end measurement. |
| Equal workload size was not defined in the article | The caption now gives balanced/ragged row lengths and equal element/task totals. All saved metadata agree. |
| Trace results must not be confused with completion timings | The earlier B starts are identified as an instrumented trace; the slowdown is explicitly from uninstrumented measurements. |
| The toy does not establish a general megakernel ranking | Its negative result remains limited to this centralized queue. The conclusion calls for comparison with good fused and library implementations. |
| Separate kernels can expose concurrency | PDL is qualified by producer-data availability and nonguaranteed overlap; independent streams/branches remain acknowledged as untested alternatives. |
The narrative now follows the title and opening question directly: repeated submission work → intermediate storage within a block → dependency scheduling across blocks → the conditions that justify a megakernel. It keeps the three-section structure, four-way baseline comparison, subtitle, and February 22 article date. Original experimental data and figures are unchanged.
Current article recheck — October 19, 2026
The current draft follows three topics: CUDA Graphs, fusion, and megakernels, preceded by an ordinary-launch baseline. The October 18 checks below concern the earlier examples. The new figures use accepted softmax run 20261019T175825Z; no GPU experiment was repeated for this editorial change.
Scroll horizontally to see the full table.
| Claim or design choice | Evidence and scope |
|---|---|
| Baseline means no graph replay or fusion | Separate kernels in one ordinary stream. Normal compiler optimizations remain enabled. This is an educational decomposition, not an optimized-library baseline. |
| Stable FP32 softmax | Finite-input formula checked against the Triton tutorial and captured CUDA source. FP32 is the arithmetic policy, not an added cast workload. |
| Graph replay retains supplied kernel bodies and arrays | CUDA Graphs guide; source and graph DOT files retain four row kernels or three tiled kernels. |
| All quoted latency values and plotted quantiles | Recomputed all 12 cases from raw JSONL and compared with saved summaries. All match. Fresh recheck. Host completion includes synchronization and excludes setup/transfers. Quantiles are not confidence intervals. |
| Local row fusion avoids a full global exponential array | Source inspection and compiler/runtime resource metadata; no reported local-memory allocation. Reduction scratch uses shared memory. No claim that logical array accesses equal HBM traffic. |
| Fusion can help and can lose to cuDNN | All six shapes and both submission modes retained in full results. The narrow-row counterexample remains in the article. Timings do not isolate storage versus launch savings. |
| The four-way scheduling plot uses the same algorithm | The baseline, graph, phase-fused, and queue paths share A/B/C helpers and global intermediates. Both persistent paths share the compiled kernel and worker count. Queue initialization is included. Source. |
| Balanced and ragged cases have equal total work | 128 rows × 8192 elements = 1,048,576 elements. Ragged groups: 36,864 + 7 × 4096 = 65,536 elements per eight rows. Both have 1024 tile producers, 128 row combines, and 1024 tile consumers. |
| Cooperative fusion retains global phase barriers | Source calls grid synchronization between A/B/C phases; Cooperative Groups documentation. |
| 126 of 128 B tasks begin before the last A ends | Rechecked fixture 2 of the ragged queue graph; every B follows its own producers. All 24 traces pass coverage/visit/dependency checks. Traces are instrumented and do not measure utilization. |
| A task queue is a megakernel mechanism, not its definition | The toy is explicitly labeled. Hazy Research uses precomputed instruction schedules. The toy slowdown does not establish performance of model-scale megakernels. |
| Graphs, fusion, and megakernels can combine | Both persistent kernels were tested under ordinary launches and graph replay. This is a teaching sequence, not a universal performance ranking. |
| Full phase boundaries are specific to the controls | Independent streams/branches and PDL are qualified alternatives, not tested here. |
| 360 outputs, 768 intermediates, 24 traces | Saved independent validation with FP64 references and predeclared tolerances. The current pass rechecks hashes, timings, and traces; large saved arrays were not downloaded again. |
| H100 suffices for these examples | Actual H100 run completed; hardware metadata, compiler logs, graph captures, and sanitizer outputs are included in selected evidence. No cross-GPU speedup generalization. |
| February 22 article date versus October 19 measurements | Both are labeled separately. No publication year inferred from the original X date. |
The new figure provenance maps every label to its exact case and variant. It records hashes separately from the immutable raw-run manifest. The site selection manifest distinguishes included artifacts from the full remote archive. The current draft is 1,180 whitespace-delimited words, versus approximately 1,441 before the three-part restructuring.
Audience-level edit
The article now assumes a mid-level engineering reader. Definitions of kernels, streams, thread blocks, softmax output properties, and the FP32 expansion have been removed. FP32 remains as an experiment parameter. The workload formula, baseline, graph-versus-fusion distinction, synchronization controls, measurement scopes, and all numerical results are retained. This is an editorial change; no experiment or figure was rerun.
Blog draft recheck — October 18, 2026
This pass checks the shortened site edition. The earlier 71-entry audit below covers the original X article, not a claim count for the current blog post. Current CUDA documentation and the cited research were consulted again; every retained numerical result was recomputed from the saved H100 data. No new GPU run was needed.
The revised article removes the historical benchmark-number detour and the illustrative SM-wave arithmetic. It also removes the inference that the pointwise timings establish how much of the improvement comes from memory work. The opening now explains the common goal and the ambiguity of “one launch,” followed by definitions of the execution and storage concepts used in the examples.
Scroll horizontally to see the full table.
| Current claim or assumption | Finding and scope | Evidence |
|---|---|---|
| Both approaches can submit multiple operations through one host call | Supported for a host-launched graph and a single fused kernel; one graph call may execute multiple kernels. This is a motivating comparison, not a guarantee of speedup. | CUDA Graphs; MPK |
| Host/device, kernel, thread block, and SM definitions | Supported; resource limits govern how many blocks can be resident. | Programming model |
| Ordinary same-stream ordering does not require the host to wait between launches | Supported; other streams may expose independent work. | Asynchronous execution |
| Capture, instantiation, replay, non-kernel nodes, and dependencies | Supported; graph replay retains the supplied kernel functions. Capturing calls is distinct from compiler fusion. | Graph structure and execution |
| Graph benefits include host work and device dispatch; setup must be recovered | Supported as possible performance effects, not fixed microsecond costs. | NVIDIA launch study |
| PDL is a qualified exception to full-completion ordering | Supported from compute capability 9.0. The consumer must establish producer-result availability, and execution overlap is not guaranteed. The toy code does not enable PDL. | PDL; benchmark source |
| Registers, shared memory, and global arrays have different scopes and lifetimes | Supported. The article does not claim all cross-block communication must use global memory, or that a launch boundary requires HBM access. | Storage model; L2 |
| Fusion and persistent worker reuse are distinct implementation choices | Supported by the cited designs and the two toy implementations. No operation-count threshold or universal interpreter requirement is asserted. | Hazy; MPK; benchmark source |
| Fusion can permit local consumers while preserving internal waits and resource costs | Supported as an opportunity. Register/shared-memory use constrains residency, and the measured fused scheduling case is a counterexample to guaranteed lower latency. | Resource limits; recomputed results |
| Pointwise arithmetic preserves rounding after each FP32 operation | CUDA uses explicit round-to-nearest add/multiply intrinsics. Saved PTX retains those operations; all stored pointwise outputs match FP32 CPU references. | CUDA source; PTX; recheck |
| 3/1 kernel counts, 3/1 host submissions, 24N/8N logical bytes | Source and graph metadata agree. Byte counts count array accesses only; they are not DRAM measurements. Capture/setup and synchronization are outside the submission counts. | Graph metadata; source |
| Same-code stream/graph pairs test submission; fusion changes multiple factors | Supported by inspection of launch lambdas and pointers. Reworded to avoid attributing the latency difference exclusively to traffic or dispatch. | source |
| Four sizes, H100 80GB HBM3, 132 SMs, 50 MiB L2, largest input 64 MiB | Device fields and sizes verified. 16,777,216 × 4 bytes = 67,108,864 bytes = 64 MiB. | metadata; source |
| Captured pointers stay fixed while inputs change for correctness checks | Source and saved fixtures agree. Timing uses a fixed fixture within each configuration; the article does not claim timing over all fixtures. | source; CPU reference checks |
| 21 samples per configuration, shuffled variant order, exact timing boundaries | Verified against source and all 588 samples. Completion includes host submission and stream synchronization. CUDA-event batches can include host-starvation gaps. | source; samples |
| All eight displayed pointwise medians and the pointwise plot's percentile statistics | All 28 configurations recomputed from raw samples and matched saved summaries. No performance change or fresh GPU allocation is implied. | recheck; figure provenance |
| 597 jobs, 256 threads, 256/8,192 producer iterations, 512 consumer iterations, 132 queue workers | Verified against code and metadata. Long jobs are indices divisible by eight, including index zero. Queue reset is included before every invocation, including as a graph memset node. | source; metadata |
| All scheduling variants retain global intermediate accesses | Source uses volatile loads/stores and block barriers; PTX retains both. This does not hold cache behavior or scheduler costs constant. Volatile is not presented as a synchronization primitive. | source; PTX |
| One active block per SM is an artificial occupancy limit | Verified occupancy API assertions for each kernel. The reservation is 116,992 shared-memory bytes; the text now says at most one block, not guaranteed utilization of every SM. | source; metadata |
| 0 versus 596 early B starts out of 597 jobs in graph traces | Recomputed for both saved fixtures. Stage timestamps measure recorded intervals; they do not independently prove instruction utilization or isolate idle causes. | all 24 trace checks; trace files |
| Ragged host medians of 206.85, 212.52, and 205.79 µs | Recomputed from 21 samples each. The wording describes this run and avoids claiming a statistically established queue speedup. | recheck |
| Balanced and ragged times cannot isolate imbalance | Producer iteration totals differ: 611,328 balanced versus 748,032 ragged, before multiplying by 256 threads. This limitation is now explicit in the article. | source; recheck |
| 272 exact comparisons, 24 trace checks, 588 timing samples | Re-ran CPU references on the full saved binary data: 112 preflight + 112 benchmark + 48 trace comparisons. Recomputed trace coverage/order and all timing groups. | recheck; independent analyzer |
| Site evidence is a selected bundle | Clarified in the article: all timing samples and selected data are on the site; full binary arrays and raw profiler files remain in the research archive. | bundle manifest; evidence index |
| Choosing graphs, fusion, persistence, or a combination | Presented as measurement guidance. No universal ranking, model speedup, or proof of cross-block synchronization correctness is claimed. | toy results; CUDA Graphs |
The machine-readable recheck records the manuscript hash, fresh numerical checks, and source references. Figures retain the original data provenance. Profiling limitations and the contaminated first-replay setup field remain documented in detailed results; neither is used to support a new performance claim.
Reading order and subtitle refinement
The opening question is now a subtitle. The current reading order is: three operations and stream ordering; graph replay; local fusion and its measured costs; blocks, SMs, and persistent workers in the scheduling example; model-scale tradeoffs and PDL; then an explicit answer about when further fusion is worth building. PDL remains a qualified exception to ordinary dependencies, introduced after the examples so it does not interrupt the basic explanation. The measured values, figures, and scope limitations are unchanged.
Concision and confidence pass
The latest pass rechecked the ten cited primary sources, inspected the benchmark and emitted PTX, and recomputed all 28 timing groups from 588 samples plus coverage and dependency ordering for all 24 traces. Every statistic matches the previous recheck. The original exact CPU-reference validation remains applicable; neither the experiment code nor its saved data changed, and no new GPU run was made.
The article removes repeated transitions and the speculative suggestion that fusion will offer little benefit once submission is addressed. It retains the limits on causal attribution, HBM traffic, instrumented traces, artificial occupancy, and model-level extrapolation. The long-job indices are explicit. The opening question, easy-to-advanced progression, and February 22 display date are unchanged.
Original X article audit
71 substantive claims and assumptions audited against primary sources. Verdicts describe the original wording; grouped rows cover closely related bullets. Coverage includes every prose section, numerical example, tradeoff list, and conclusion. The cover illustration is outside this text audit.
The X item is already published. The inspected UI offers unpublishing to drafts before revision; the corrected local manuscript leaves the published version live. The UI's February 22 date does not display a year, so no publication year is inferred.
The central corrections concern universal launch constants, CPU versus device timing, global memory versus HBM, remaining tails and synchronization, the B200 runtime breakdown, percentage denominators, and historical benchmark scope. Source snapshots are retained locally for research with hashes; citations point to the originals.
Scroll horizontally to see the full table.
| ID | Original section and excerpt | Verdict | Verified finding | Primary evidence |
|---|---|---|---|---|
| C01 | Opening: They operate at different layers of the GPU execution stack | qualify | Useful distinction: submission versus device implementation. Layers interact and techniques can be combined. | graphs, mpk |
| C02 | Execution layer: Application layer — Python/C++ framework logic (PyTorch, vLLM) | retain | A useful explanatory abstraction, not an official exhaustive CUDA taxonomy. | streams |
| C03 | Execution layer: Runtime submission layer — kernel launch and dependency scheduling | retain | Submission and scheduling are useful concepts; some scheduling also occurs on device. | graphs, mpk |
| C04 | Execution layer: Kernel boundary — each individual global function launch | qualify | A kernel launch creates a grid; dependencies and execution scope matter more than a four-layer classification. | memory, streams |
| C05 | Execution layer: In-kernel execution — threads, warps, shared memory, registers | retain | Correct examples of device execution and storage concepts. | memory |
| C06 | Execution layer: CUDA Graph -> runtime submission layer | qualify | Also affects device dispatch; graph nodes and dependencies expose a complete workflow to CUDA. | graphs, launch |
| C07 | Execution layer: MegaKernels -> kernel boundary and in-kernel layers | retain | Correct conceptual emphasis, with explicit dependency and resource management. | mpk |
| C08 | CUDA Graph: The normal way to run GPU code is to submit kernels one at a time from the CPU. | qualify | Describes conventional host stream launches; CUDA also supports graph and device-side launch mechanisms. | graphs, streams |
| C09 | CUDA Graph: Each submission goes through the CUDA driver and lands on a GPU stream. | qualify | Describe runtime/driver submission without asserting a distinct direct driver transition per framework operation. | graphs |
| C10 | CUDA Graph: it triggers the entire workflow at once | qualify | One replay enqueues a workflow; nodes execute when their dependencies permit, not simultaneously. | graphs |
| C11 | CUDA Graph: reducing the idle time ("jitter") between kernels | qualify | Can reduce submission and dispatch gaps; jitter specifically means variability and is not a synonym for all idle time. | launch |
| C12 | CUDA Graph: about 2.1 us on a regular stream, and about 1.3 us with CUDA Graph replay | qualify | Attributed dummy-kernel observation; distinguish the observed device idle/setup interval from CPU API return time. | hazy, launch |
| C13 | CUDA Graph: That's 38% decrease. | retain | (2.1−1.3)/2.1 = 38.095%; a reduction in that reported overhead, not end-to-end model latency. | hazy |
| C14 | CUDA Graph: do not: Fuse kernels | retain | Replay does not itself fuse kernel bodies. Framework compilation may fuse before capture. | graphs |
| C15 | CUDA Graph: Eliminate global-memory traffic between kernels | retain | Replay does not automatically remove the captured kernels' intermediate accesses. | graphs |
| C16 | CUDA Graph: Remove kernel-boundary synchronization | correct | Default completion edges preserve ordering; graphs can also encode PDL programmatic edges. Avoid an unconditional prohibition. | graphs, pdl |
| C17 | CUDA Graph: Keep intermediate data in registers across kernels | retain | No general register handoff across separate kernels. | memory |
| C18 | CUDA Graph: Each kernel inside a CUDA Graph is still an independent kernel. | qualify | Separate kernel nodes remain distinct kernels, but they can have explicit dependencies. A graph also has non-kernel node types. | graphs |
| C19 | CUDA Graph: The graph makes dispatching cheaper. | retain | Can amortize setup and improve repeated dispatch; retain the setup costs and measured scope. | graphs, launch |
| C20 | PDL: can launch a dependent kernel early so it can do independent work sooner | retain | Supported on compute capability 9.0+. Concurrent execution is opportunistic. | pdl |
| C21 | PDL: dependent work still has to synchronize (cudaGridDependencySynchronize) before consuming producer results | qualify | The required data-availability check applies; the guide also permits other means of verification. Never rely on concurrent progress. | pdl |
| C22 | PDL: this can still leave pipeline stalls for fine-grained dependencies | retain | A possibility, not an invariant measured PDL cost. | pdl |
| C23 | Boundary overhead: Three overhead sources stack at each boundary | correct | They may matter and overlap; do not treat them as three independent additive constants. | launch, streams |
| C24 | Boundary overhead: 1.3-2.1 us of scheduling overhead | correct | Do not reuse one microbenchmark as a universal scheduling cost. | launch, hazy |
| C25 | Boundary overhead: With in-order execution on a stream, boundaries force whole-kernel completion semantics. | qualify | Correct for ordinary ordered launches; mention PDL and separate-stream concurrency. | streams, pdl |
| C26 | Boundary overhead: 512 thread blocks on a 148-SM B200 leaves a partial final wave with only 68/148 SMs active (80 idle) | qualify | 512 mod 148 = 68 and 148−68 = 80. This requires an idealized one-resident-block-per-SM wave model; real occupancy and durations vary. | hazy, memory |
| C27 | Boundary overhead: Intermediate tensors are written to HBM, then read by the next kernel | correct | Use global memory; L2 may serve the accesses. DRAM traffic requires counters rather than an assumption. | cache |
| C28 | Boundary overhead: Even with PDL-style overlap, dependent sections still wait on data visibility/completion. | retain | Data dependencies still require correctness synchronization. | pdl |
| C29 | Llama estimate: around ~100 kernels for Llama-1B forward | qualify | Keep attribution, model/configuration, and date; not a timeless launch count for every serving stack. | hazy |
| C30 | Llama estimate: ~130 us of launch overhead per forward pass as a back-of-the-envelope estimate | qualify | 100×1.3 = 130 µs, but this transfers a dummy-kernel observation. It is not a validated model-runtime decomposition; remove from the main argument. | hazy, launch |
| C31 | Llama estimate: vLLM/SGLang reaching only ~50% of available H100 memory bandwidth in this workload | qualify | Historical attributed result, not a statement about current versions or an independently reproduced counter measurement. | hazy |
| C32 | Megakernel definition: one global CUDA function that executes what would traditionally be many kernels | retain | A useful working definition; persistent execution and large fusion are related rather than universally identical terms. | mpk |
| C33 | Megakernel definition: kernel_1: RMSNorm ... kernel_5: RMSNorm + MLP | qualify | Illustrative partition only; actual stacks fuse and split these operations differently. | mega |
| C34 | Megakernel definition: removes inter-kernel boundaries and their associated launch/tail/bubble overhead | correct | Removes those launches; reducing tail and pipeline costs requires a suitable internal schedule and communication design. | mpk, mla |
| C35 | Megakernel definition: For a 16-layer Llama-1B path ... ~100 boundaries removed | qualify | A rough configuration-specific scale. K separate launches fused to one removes K−1 inter-launch boundaries; do not present an exact count without a trace. | hazy |
| C36 | Implementation: Fusing 2-3 ops is standard fusion. Fusing ~100 ops is a different problem. | qualify | Useful scale intuition, not a formal threshold between fusion and megakernels. | mpk |
| C37 | Implementation: A common pattern is an on-GPU interpreter | qualify | Established for the cited Hazy designs; neither all fused kernels nor all persistent kernels require an interpreter. | mla, mega |
| C38 | Implementation: each SM executes a precomputed instruction stream | qualify | Specific to Hazy's scheduling design; CUDA ordinarily assigns blocks to SMs, not one abstract interpreter per SM by definition. | mega |
| C39 | Implementation: Schedule construction happens on CPU once and is reused. | qualify | Schedules can be reused; input shape, batch composition, and runtime work can require updates or dynamic scheduling. | mla, mpk |
| C40 | ThunderKittens: provides the low-level kernel framework used by Hazy's MegaKernel work | retain | Supported by the official repositories and ThunderMLA implementation description. | tk, mega, mla |
| C41 | ThunderMLA: 20-35% faster than DeepSeek FlashMLA across tested workloads | qualify | Headline range, not every reported case: the same post includes a 7.6% speedup case and uses throughput-style speedup percentages. Use case-specific time ratios. | mla |
| C42 | ThunderMLA: 41 us vs 52 us on SXM H100 | retain | Verified representative configuration. 21.15% latency reduction; 52/41 = 1.2683× speedup. Include prompt lengths, query count, and tensor parallelism. | mla |
| C43 | Look Ma: full Llama-3.2-1B forward megakernel | retain | Scope remains the cited single-sequence 16-bit inference experiment, not arbitrary models or training. | mega |
| C44 | Look Ma: 78% memory bandwidth utilization and over 1.5x speedup vs baselines on H100 | qualify | Keep historical attribution and workload. The reported baseline comparisons differ; not a generic expected speedup. | hazy |
| C45 | Look Ma: code is open-sourced | retain | The official Hazy Megakernels repository is accessible. | mega |
| C46 | MPK: mega-kernelizes PyTorch programs with a few lines of code changes | qualify | Author-reported interface for supported models/configurations, not a guarantee for arbitrary PyTorch programs. | mpk |
| C47 | MPK: up to 1.0-1.7x latency improvement vs vLLM/SGLang | correct | Use the paper's 1.0–1.7× single-batch serving-performance range; say up to 1.7×, not up to a range. Keep evaluation scope. | mpkv1, mpk |
| C48 | MPK: Qwen3-8B on A100 from 14.5 ms to 12.5 ms per-token decode latency | retain | Present in v1 and v2 §6.3; 13.79% latency reduction and 1.16× speedup. This system-level result has multiple causes. | mpkv1, mpk |
| C49 | Critical difference: CUDA Graphs eliminate: CPU submission overhead per kernel | correct | Graphs amortize per-node host work into replay, which still has nonzero enqueue and device costs. | graphs, launch |
| C50 | Critical difference: Kernel boundary setup/teardown overhead | qualify | Eliminates repeated launches across fused stages; one launch and internal task transitions remain. | mpk |
| C51 | Critical difference: Cross-kernel tail effects | correct | Can avoid whole-kernel waits for unrelated ready work; remaining load imbalance and the final tail are not erased. | mla, mpk |
| C52 | Critical difference: Full intermediate round-trips between separately launched kernels | correct | Only eliminated for values actually kept local/reused. Cross-block communication may remain global; separate accesses may hit cache. | memory, cache |
| C53 | Critical difference: Many memory-pipeline bubbles caused by kernel boundaries | qualify | Possible when an implementation actually pipelines stages; requires evidence from the chosen schedule. | mla, mpk |
| C54 | Counter-intuitive: megakernels win by improving memory behavior. True in many cases. | qualify | Replace an unquantified prevalence claim with specific mechanisms and measured examples. | mpk |
| C55 | Counter-intuitive: asynchronous barrier overhead around ~60 ns each | qualify | Hazy's observed low-level overhead on its B200 implementation, not a universal CUDA primitive latency. | hazy |
| C56 | Counter-intuitive: sync plus setup/other overhead summing to about ~120 us | correct | 40+80 = 120 µs in the 600-µs B200 breakdown. Original paragraph omits the architecture and nearby context suggests H100. | hazy |
| C57 | Counter-intuitive: megakernels do not remove all global-memory traffic | retain | Correct; storage and synchronization follow the implementation and dependency scope. | memory, mpk |
| C58 | Counter-intuitive: Large cross-SM dependencies and broad reductions still require global-memory communication. | qualify | Often true for these implementations; cluster/distributed shared-memory and cooperative mechanisms make 'require' too universal. | memory |
| C59 | Counter-intuitive: What they remove is the repeated boundary-induced waste | qualify | Launches are removed; costs can reappear as internal communication, scheduling, or resource pressure. | mpk |
| C60 | Combination: A megakernel can be launched via CUDA Graphs. | retain | It can be a kernel node, subject to normal capture/graph restrictions. | graphs |
| C61 | Combination: if ... one kernel, graph benefit is usually small | qualify | Plausible workload hypothesis, not an established universal rule. Measure the one-kernel stream/graph pair. | launch |
| C62 | Combination: CUDA Graph value remains highest when many kernel boundaries still exist. | qualify | Topology, host work, setup amortization, and node costs matter; launch count alone cannot rank graph benefit. | graphs, launch |
| C63 | Tradeoffs: Minimal code changes; Keeps modular kernels; Native support in common stacks; Easier debugging/maintenance | qualify | Graphs preserve kernel modularity, but capture-safe control flow, memory lifetime, changing parameters, and debugging can require engineering. Minimal/easier are context-dependent judgments. | graphs |
| C64 | Tradeoffs: Removes cross-kernel tail effects; Enables deeper cross-op pipelining; Raises bandwidth utilization; Can reduce per-token latency | qualify | These are opportunities under suitable implementation/workload conditions, not guaranteed properties of a single launch. | mla, mpk |
| C65 | Tradeoffs: Harder to implement/debug; Portability/tuning ... harder; Synchronization mistakes ... dangerous; bottlenecks ... shift | qualify | Reasonable engineering considerations; present them as tradeoffs rather than measured universal comparisons. | mpk, memory |
| C66 | Recommendation: For most teams: start with CUDA Graphs, profile, then invest ... only if boundary overhead is clearly dominant. | judgment | Keep profiling as guidance. Avoid an unsupported claim about most teams or a rule excluding fusion benefits from memory reuse. | launch, mpk |
| C67 | Analogy: CUDA Graphs = better dispatching. Same stations, cheaper scheduling. | qualify | Useful mnemonic, but it omits device scheduling, DAG concurrency, and non-kernel nodes. | graphs |
| C68 | Analogy: Megakernels = express route. Remove stops between segments. | qualify | Can remove external boundaries while retaining internal waits. Replace with the precise implementation comparison. | mpk |
| C69 | Conclusion: Both reduce end-to-end time | correct | Both can reduce time; either can fail to improve or regress a workload. | launch, mpk |
| C70 | Conclusion: CUDA Graphs and megakernels are not substitutes for the same thing. | qualify | They change different aspects and can combine, while both may reduce some overlapping launch/dispatch overhead. | graphs, mpk |
| C71 | Conclusion: Choosing well starts with identifying which execution layer is actually bottlenecking | retain | Retain as practical advice, with measurement scope and actual implementation evidence. | launch |
Counts: correct 11, judgment 1, qualify 38, retain 21.
retain means supported within stated scope, not independently reproduced. qualify narrows wording. correct replaces a misleading or false generalization. judgment distinguishes advice from an empirical fact.
The corrected manuscript incorporates the findings and two controlled examples. The protocol separates planned tests from measured outcomes. Historical serving results are attributed reports; the new toy measurements cannot independently reproduce or validate them.