Integrate fast.cu kernels and B200 optimization experiments

code/ch09/README.md

@@ -87,3 +87,10 @@ python -m cli.aisp bench run --targets ch09 --profile minimal
 - `inline_ptx_example.cu` demonstrates how to wrap tcgen05 intrinsics safely with architecture guards.
 - `requirements.txt` includes the shared `requirements_latest.txt` pins for PyTorch 2.9.1 and its compatible Triton release.
 - `optimized_cublaslt_gemm_fp4` is intentionally capability-gated: if cuBLASLt cannot provide the native block-scaled NVFP4 heuristic, the benchmark reports a clean skip instead of silently falling back to a different FP4 mode.
+
+## Related fast.cu experiments
+
+The [fast.cu lab](../labs/fast_cu/README.md) preserves the complete H100 and GB300
+source progressions and adds native harness pairs, including B200 adaptations.
+See its [chapter mapping and optimization guide](../labs/fast_cu/optimization_guide.md)
+for the overlap with this chapter and the SM100 versus SM103 hardware boundaries.

code/ch10/README.md

@@ -88,3 +88,10 @@ python -m cli.aisp bench run --targets ch10 --profile minimal
 - `cufile_gds_example.py` demonstrates integrating GPUDirect Storage into tensor-core pipelines for IO-heavy training loops.
 - `requirements_cufile.txt` holds the optional `cufile` wheel; install it only on hosts with GPUDirect Storage enabled.
 - The CUTLASS-style warp-specialization pair provides a reference implementation aligned with `sm100_mma_array_warpspecialized` for performance comparison.
+
+## Related fast.cu experiments
+
+The [fast.cu lab](../labs/fast_cu/README.md) preserves the complete H100 and GB300
+source progressions and adds native harness pairs, including B200 adaptations.
+See its [chapter mapping and optimization guide](../labs/fast_cu/optimization_guide.md)
+for the overlap with this chapter and the SM100 versus SM103 hardware boundaries.

code/ch11/README.md

@@ -86,3 +86,10 @@ python -m cli.aisp bench run --targets ch11 --profile minimal
 - The README calls out the legacy target names explicitly so book-facing labels still point at the actual copy+elementwise overlap workload and runtime-adaptive scheduling pairs.
 - `warp_specialized_triton.py` provides a Triton analogue for the CUDA concurrency demos so you can compare compiler-generated schedules.
 - `kv_prefetch_pipeline_enhanced_demo.cu` builds on the DSMEM kernels bundled in this directory so you can study the entire pipeline locally.
+
+## Related fast.cu experiments
+
+The [fast.cu lab](../labs/fast_cu/README.md) preserves the complete H100 and GB300
+source progressions and adds native harness pairs, including B200 adaptations.
+See its [chapter mapping and optimization guide](../labs/fast_cu/optimization_guide.md)
+for the overlap with this chapter and the SM100 versus SM103 hardware boundaries.

code/ch14/README.md

@@ -86,3 +86,10 @@ python -m cli.aisp bench run --targets ch14 --profile minimal
 - `requirements.txt` includes the shared `requirements_latest.txt` pins for stable PyTorch and Triton wheels aligned with the CUDA 13 toolchain.
 - For repo-native supporting examples that fill the training hot-path gaps without changing this chapter's primary compile narrative, see `labs/training_hotpath`.
 - `cublas_vs_cutlass` is a supplementary comparison pair. Chapter-native performance claims stay anchored on `model_compile_reduced_precision`, `regional_triton`, and `triton_persistent`.
+
+## Related fast.cu experiments
+
+The [fast.cu lab](../labs/fast_cu/README.md) preserves the complete H100 and GB300
+source progressions and adds native harness pairs, including B200 adaptations.
+See its [chapter mapping and optimization guide](../labs/fast_cu/optimization_guide.md)
+for the overlap with this chapter and the SM100 versus SM103 hardware boundaries.

code/core/scripts/refresh_readmes.py

@@ -48,6 +48,7 @@ class Entry:
     validation_heading: str = "Validation Checklist"
     extra_sections: Sequence[str] = field(default_factory=tuple)
     notes: Sequence[str] = field(default_factory=tuple)
+    trailing_sections: Sequence[str] = field(default_factory=tuple)
 
 
 def _format_markdown(entry: Entry) -> str:
@@ -97,6 +98,12 @@ def _format_markdown(entry: Entry) -> str:
         lines.append("## Notes")
         for note in entry.notes:
             lines.append(f"- {note}")
+    if entry.trailing_sections:
+        lines.append("")
+        for idx, section in enumerate(entry.trailing_sections):
+            if idx:
+                lines.append("")
+            lines.append(section.strip())
     lines.append("")
     return "\n".join(lines)
 
@@ -404,6 +411,7 @@ def chapter_entry(
         "when you need repeatable artifact capture."
     ),
     validation_heading: str = "Validation Checklist",
+    trailing_sections: Sequence[str] = (),
 ) -> Entry:
     """Create a chapter README entry."""
     return Entry(
@@ -418,6 +426,7 @@ def chapter_entry(
         run_intro=run_intro,
         validation_heading=validation_heading,
         notes=notes,
+        trailing_sections=trailing_sections,
     )
 
 
@@ -660,6 +669,17 @@ def lab_entry(
     """
 ).strip()
 
+FAST_CU_RELATED_SECTION = dedent(
+    """\
+    ## Related fast.cu experiments
+
+    The [fast.cu lab](../labs/fast_cu/README.md) preserves the complete H100 and GB300
+    source progressions and adds native harness pairs, including B200 adaptations.
+    See its [chapter mapping and optimization guide](../labs/fast_cu/optimization_guide.md)
+    for the overlap with this chapter and the SM100 versus SM103 hardware boundaries.
+    """
+)
+
 ENTRIES["README.md"] = Entry(
     title="AI Systems Performance Engineering",
     summary=dedent(
@@ -965,6 +985,7 @@ def lab_entry(
             | `labs/cutlass_profiler_kernel_selector/` | CUTLASS profiler-based kernel selection | ch06, ch09 |
             | `labs/decode_optimization/` | Decoder hot-path optimization | ch18, ch19 |
             | `labs/dynamic_router/` | Dynamic prefill/decode routing | ch17, ch19 |
+            | `labs/fast_cu/` | Pinned fast.cu H100 GEMM, B200 reduction and wide-store epilogue adaptations, and the complete GB300 NVFP4 ladder; see the [chapter mapping](fast_cu/optimization_guide.md) | ch09, ch10, ch11, ch14 |
             | `labs/flashattention4/` | FlashAttention-4 pipeline co-design, including a [Colfax diary guide](flashattention4/colfax_optimization_diaries.md) to decode S/P ping-pong and backward hdim-64 TMEM de-aliasing | ch10, ch11, ch18 |
             | `labs/flashattention_gluon/` | FlashAttention experimentation | ch18 |
             | `labs/flashinfer_attention/` | FlashInfer block-sparse attention lab | ch16 |
@@ -1972,6 +1993,7 @@ def lab_entry(
         "`requirements.txt` includes the shared `requirements_latest.txt` pins for PyTorch 2.9.1 and its compatible Triton release.",
         "`optimized_cublaslt_gemm_fp4` is intentionally capability-gated: if cuBLASLt cannot provide the native block-scaled NVFP4 heuristic, the benchmark reports a clean skip instead of silently falling back to a different FP4 mode.",
     ],
+    trailing_sections=[FAST_CU_RELATED_SECTION],
 )
 
 ENTRIES["ch10"] = chapter_entry(
@@ -2077,6 +2099,7 @@ def lab_entry(
         "`requirements_cufile.txt` holds the optional `cufile` wheel; install it only on hosts with GPUDirect Storage enabled.",
         "The CUTLASS-style warp-specialization pair provides a reference implementation aligned with `sm100_mma_array_warpspecialized` for performance comparison.",
     ],
+    trailing_sections=[FAST_CU_RELATED_SECTION],
 )
 
 ENTRIES["ch11"] = chapter_entry(
@@ -2183,6 +2206,7 @@ def lab_entry(
         "`warp_specialized_triton.py` provides a Triton analogue for the CUDA concurrency demos so you can compare compiler-generated schedules.",
         "`kv_prefetch_pipeline_enhanced_demo.cu` builds on the DSMEM kernels bundled in this directory so you can study the entire pipeline locally.",
     ],
+    trailing_sections=[FAST_CU_RELATED_SECTION],
 )
 
 ENTRIES["ch12"] = chapter_entry(
@@ -2554,6 +2578,7 @@ def lab_entry(
         "For repo-native supporting examples that fill the training hot-path gaps without changing this chapter's primary compile narrative, see `labs/training_hotpath`.",
         "`cublas_vs_cutlass` is a supplementary comparison pair. Chapter-native performance claims stay anchored on `model_compile_reduced_precision`, `regional_triton`, and `triton_persistent`.",
     ],
+    trailing_sections=[FAST_CU_RELATED_SECTION],
 )
 
 ENTRIES["ch15"] = chapter_entry(

code/labs/README.md

@@ -66,6 +66,7 @@ python -m cli.aisp bench list-targets --chapter labs/moe_cuda
 | `labs/cutlass_profiler_kernel_selector/` | CUTLASS profiler-based kernel selection | ch06, ch09 |
 | `labs/decode_optimization/` | Decoder hot-path optimization | ch18, ch19 |
 | `labs/dynamic_router/` | Dynamic prefill/decode routing | ch17, ch19 |
+| `labs/fast_cu/` | Pinned fast.cu H100 GEMM, B200 reduction and wide-store epilogue adaptations, and the complete GB300 NVFP4 ladder; see the [chapter mapping](fast_cu/optimization_guide.md) | ch09, ch10, ch11, ch14 |
 | `labs/flashattention4/` | FlashAttention-4 pipeline co-design, including a [Colfax diary guide](flashattention4/colfax_optimization_diaries.md) to decode S/P ping-pong and backward hdim-64 TMEM de-aliasing | ch10, ch11, ch18 |
 | `labs/flashattention_gluon/` | FlashAttention experimentation | ch18 |
 | `labs/flashinfer_attention/` | FlashInfer block-sparse attention lab | ch16 |

code/labs/fast_cu/README.md

@@ -0,0 +1,86 @@
+# Lab - fast.cu kernel optimization studies
+
+This lab brings [Pranjal Shankhdhar's fast.cu](https://github.com/pranjalssh/fast.cu)
+into the repository: all twelve Hopper GEMM snapshots, four reduction variants,
+and all ten GB300 NVFP4 snapshots live in [upstream/](upstream/). Repository-native
+`baseline_*` / `optimized_*` pairs add explicit hardware gates, current-stream
+execution, caller-owned random inputs, and full output verification.
+
+The [optimization guide](optimization_guide.md) maps the examples and lessons to
+existing chapters and labs. The [source manifest](upstream_manifest.json) pins
+revision `2dfe5e26aecfd9e5f27bf9d5837deea01acda24b` and the SHA256 of every upstream
+file. The upstream MIT license is preserved in [upstream/LICENSE](upstream/LICENSE).
+
+## Native benchmark pairs
+
+| Target in `labs/fast_cu` | Baseline | Optimized | Required GPU/toolkit |
+| --- | --- | --- | --- |
+| `h100_bf16_gemm` | cuBLAS BF16 GEMM | fast.cu Hopper kernel 12 | H100, SM90, CUDA 12+ |
+| `int32_reduction` | CUB sum | Corrected vectorized warp reduction | H100 SM90/CUDA 12+ or B200 SM100/CUDA 12.8+ |
+| `b200_epilogue` | Two 128-bit stores per thread | One 256-bit store per thread | B200, SM100, CUDA 12.9+ |
+| `nvfp4_gemm` | cuBLASLt NVFP4 GEMM | fast.cu r9 NVFP4 GEMM | B300/GB300, SM103, CUDA 13.1+ |
+
+The NVFP4 kernel uses architecture-specific SM103 instructions and is **not a
+B200 kernel**. Changing the compiler target to SM100 does not port its K=96 MMA
+pipeline. The B200 epilogue is a separate experiment applying the article's wide
+store technique to the same FP32-to-FP16 conversion workload in both arms. It
+does not stand in for a complete NVFP4 GEMM.
+
+## Run with the repository harness
+
+Run from `code/` in the normal CUDA-enabled repository environment (PyTorch,
+matching local CUDA toolkit, Ninja, Nsight Compute and Nsight Systems):
+
+```bash
+python -m cli.aisp bench list-targets --chapter labs/fast_cu
+
+# B200: architecture-compatible adaptations
+python -m cli.aisp bench run --targets labs/fast_cu:int32_reduction --profile deep_dive --single-gpu
+python -m cli.aisp bench run --targets labs/fast_cu:b200_epilogue --profile deep_dive --single-gpu
+
+# H100
+python -m cli.aisp bench run --targets labs/fast_cu:h100_bf16_gemm --profile deep_dive --single-gpu
+
+# B300/GB300 with CUDA 13.1+
+python -m cli.aisp bench run --targets labs/fast_cu:nvfp4_gemm --profile deep_dive --single-gpu
+```
+
+Select an allocated GPU using `CUDA_VISIBLE_DEVICES` before launching. Unsupported
+hardware/toolkits produce explicit `SKIPPED:` diagnostics; no precision or
+architecture fallback is substituted. Extension compilation and preparation run
+in setup, outside timed execution. The source tree is checked before compilation,
+and build cache names include all upstream file digests and adapter sources.
+
+Use the repository's strict validity profile and clock-lock/profiler machinery for
+performance evidence. The original upstream Makefile and shell scripts are retained
+as source evidence; their direct `sudo nvidia-smi`/`ncu` helpers are not the lab's
+execution interface. Do not build inside the pinned `upstream/` directory: generated
+files intentionally fail the exact file-set integrity check.
+
+## Correctness and measurement boundaries
+
+- Both arms use the same logical inputs, precision, shape, and full timed output.
+  Preparation, JIT, workspace allocation, and reference checks are separate from
+  the steady-state hot path.
+- The reduction adaptation moves output initialization ahead of the grid on the
+  same stream, preventing the original cross-block reset race. Its bounded input
+  distribution avoids signed 32-bit overflow.
+- The Hopper adapter normalizes the kernel's physical transposed output into a
+  logical matrix view for verification. A successful upstream process exit is not
+  treated as evidence of correctness.
+- The NVFP4 native pair uses a single set of resident buffers. The article's
+  headline instead rotates input buffers and alternates five cooled-down timing
+  rounds. Those are different workload/cache protocols; this lab does not import
+  the article's speedup as an expectation.
+- [Validation](validation.md) records what has actually run and what remains
+  hardware-gated. Source preservation and CPU tests do not prove GPU correctness.
+
+Focused checks, from `code/`:
+
+```bash
+python -m pytest tests/test_fast_cu_*.py -q
+python -m core.scripts.linting.check_benchmarks labs/fast_cu --fail-on-warnings
+```
+
+See [performance_intake.yaml](performance_intake.yaml) and
+[workload_spec.yaml](workload_spec.yaml) for cost hypotheses and fixed workloads.

code/labs/fast_cu/__init__.py

@@ -0,0 +1 @@
+"""Pinned fast.cu CUDA kernels and repository-native benchmark adapters."""