|
1 | 1 | # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. |
2 | 2 | # SPDX-License-Identifier: Apache-2.0 |
3 | 3 |
|
4 | | -"""Blackwell FMHA-prefill port from CUTLASS Python DSL to CUDA Lang. |
| 4 | +"""Blackwell FMHA-prefill CUDA Lang tutorial. |
5 | 5 |
|
6 | 6 | This is a native fused implementation, with no SIMT, PyTorch-SDPA, library, |
7 | | -source-kernel, or precompiled-PTX fallback. It supports the source tutorial's |
8 | | -fixed and variable sequence storage, D=32/64/128, FP16/BF16/E4M3 inputs and |
9 | | -outputs, MHA/GQA, mask variants, LSE, sink logits, scaling controls, and |
10 | | -correction/epilogue options. |
| 7 | +or precompiled-PTX fallback. It supports fixed and variable sequence storage, |
| 8 | +D=32/64/128, FP16/BF16/E4M3 inputs and outputs, MHA/GQA, mask variants, LSE, |
| 9 | +sink logits, scaling controls, and correction/epilogue options. |
11 | 10 |
|
12 | | -The device kernel retains the source's structural contract: |
| 11 | +The device kernel uses: |
13 | 12 |
|
14 | 13 | * one 512-thread CTA and one CTA per ``(256 Q rows, Q head, batch)`` tile; |
15 | 14 | * softmax warps 0-7, correction warps 8-11, MMA warp 12, TMA warp 13, |
16 | 15 | epilogue warp 14, and CLC scheduler warp 15; |
17 | 16 | * CTA-group-1 tcgen05 QK (128x128x16) and PV (128x64x16); |
18 | 17 | * two Q stages, three unified K/V stages, two output stages, and all 512 |
19 | | - TMEM columns with the source S/P/O aliases; |
| 18 | + TMEM columns with S/P/O aliases; |
20 | 19 | * online softmax, packed-FP16 P in TMEM, FP32 O correction in TMEM, TMA |
21 | 20 | global/shared transfers, and a one-slot 480-consumer CLC mailbox. |
22 | 21 |
|
23 | | -The source-compatible host configuration, chunked reference, correctness |
24 | | -matrix, cold-L2/CUDA-graph benchmark, and CLI share this public entry point. |
| 22 | +The host configuration, chunked reference, correctness matrix, |
| 23 | +cold-L2/CUDA-graph benchmark, and CLI share this public entry point. |
25 | 24 | Invalid combinations raise clear diagnostics and are never silently redirected |
26 | 25 | to another attention implementation. |
27 | 26 | """ |
@@ -376,9 +375,8 @@ def _store_output_values( |
376 | 375 | word_count = 16 |
377 | 376 | if output_kind == ELEMENT_E4M3: |
378 | 377 | word_count = 8 |
379 | | - # One PTX shared-memory vector store carries four 32-bit registers. The |
380 | | - # swizzle is constant within each 16-byte group, so form the same v4 |
381 | | - # transfer units as CUTLASS's store_swizzled lowering. |
| 378 | + # One PTX shared-memory vector store carries four 32-bit registers. The |
| 379 | + # swizzle is constant within each 16-byte group, so form v4 transfer units. |
382 | 380 | for group in cl.static_iter(range(word_count // 4)): |
383 | 381 | byte_offset = ( |
384 | 382 | slice_index * MMA_M * slice_bytes |
@@ -421,8 +419,8 @@ def _store_output_global( |
421 | 419 | alignment=1, |
422 | 420 | ) |
423 | 421 | else: |
424 | | - # CUTLASS keeps the converted FP16/BF16 lanes scalar on the guarded |
425 | | - # tail path, which lowers to one b16 store per valid element. |
| 422 | + # Keep converted FP16/BF16 lanes scalar on the guarded tail path, which |
| 423 | + # lowers to one b16 store per valid element. |
426 | 424 | dst_u16 = cl.bitcast( |
427 | 425 | dst, |
428 | 426 | cl.pointer_dtype(cl.uint16, cl.MemorySpace.GLOBAL), |
@@ -489,12 +487,12 @@ def _key_trip_context( |
489 | 487 | window_left, |
490 | 488 | window_right, |
491 | 489 | ): |
492 | | - """Return the source super-tile K interval and its Q0 invalid tail. |
| 490 | + """Return the super-tile K interval and its Q0 invalid tail. |
493 | 491 |
|
494 | | - CUTLASS prunes whole 128-column K/V tiles using the union of the two |
495 | | - 128-row Q halves. Q0 can therefore have one fully-invalid trailing tile |
496 | | - that Q1 still needs. That tile remains in the barrier protocol but does |
497 | | - not issue QK0/PV0 math. |
| 492 | + Whole 128-column K/V tiles are pruned using the union of the two 128-row Q |
| 493 | + halves. Q0 can therefore have one fully-invalid trailing tile that Q1 still |
| 494 | + needs. That tile remains in the barrier protocol but does not issue QK0/PV0 |
| 495 | + math. |
498 | 496 | """ |
499 | 497 | query_begin = seq_tile * CTA_M |
500 | 498 | offset = cl.int32(0) |
@@ -792,8 +790,8 @@ def _fmha_prefill_kernel( |
792 | 790 | sched_valid = cl.shared_array(1, cl.int32, alignment=4) |
793 | 791 |
|
794 | 792 | tid = cl.thread_index(0) |
795 | | - # Match CUTLASS's make_warp_uniform(warp_idx): broadcasting lane 0's |
796 | | - # value lets the backend treat every role predicate as warp-uniform. |
| 793 | + # Broadcasting lane 0's warp index lets the backend treat every role |
| 794 | + # predicate as warp-uniform. |
797 | 795 | warp = cl.shfl_sync(tid // WARP_SIZE, 0) |
798 | 796 | lane = tid % WARP_SIZE |
799 | 797 | if variable_length: |
@@ -855,9 +853,9 @@ def _fmha_prefill_kernel( |
855 | 853 | l2_promotion=cl.TensorMapL2Promotion.NONE, |
856 | 854 | ) |
857 | 855 |
|
858 | | - # CUTLASS initializes Blackwell pipeline stages cooperatively with warp 0. |
859 | | - # Dynamic lane indexing keeps one static PTX site per pipeline half while |
860 | | - # still initializing every physical barrier exactly once. |
| 856 | + # Initialize pipeline stages cooperatively with warp 0. Dynamic lane |
| 857 | + # indexing keeps one static PTX site per pipeline half while initializing |
| 858 | + # every physical barrier exactly once. |
861 | 859 | if warp == SOFTMAX0_WARPS[0]: |
862 | 860 | if lane < Q_STAGES: |
863 | 861 | cl.mbarrier_initialize(q_full.get_element_pointer(lane), 1) |
@@ -1313,8 +1311,8 @@ def _fmha_prefill_kernel( |
1313 | 1311 | kv_index = 0 |
1314 | 1312 | kv_full_phase = 0 |
1315 | 1313 | while valid: |
1316 | | - # CUTLASS reuses one elected MMA-warp lane for all UMMA math in a |
1317 | | - # work tile. Pipeline commits elect independently in their helper. |
| 1314 | + # Reuse one elected MMA-warp lane for all UMMA math in a work tile. |
| 1315 | + # Pipeline commits elect independently in their helper. |
1318 | 1316 | mma_elect_one = cl.elect_sync() |
1319 | 1317 | query_base, key_base, task_seq_q, task_seq_k = _sequence_context( |
1320 | 1318 | cumulative_q, |
@@ -1364,7 +1362,7 @@ def _fmha_prefill_kernel( |
1364 | 1362 |
|
1365 | 1363 | # QK0(0), QK1(0): both halves share K0 before its ring slot is |
1366 | 1364 | # released. The acquire mbarrier wait orders the TMA async writes |
1367 | | - # before the tcgen MMA, matching CUTLASS's TMA-to-UMMA pipeline. |
| 1365 | + # before the tcgen MMA in the TMA-to-UMMA pipeline. |
1368 | 1366 | k_stage = kv_index |
1369 | 1367 | _wait_mbarrier(kv_full.get_element_pointer(k_stage), kv_full_phase) |
1370 | 1368 | k_desc = _qk_descriptor( |
@@ -1842,9 +1840,9 @@ def _fmha_prefill_kernel( |
1842 | 1840 | for chunk in cl.static_iter(range(4)) |
1843 | 1841 | ) |
1844 | 1842 |
|
1845 | | - # Match CUTLASS's four 32-value vector reductions: |
1846 | | - # MAXIMUMF propagates NaNs within each chunk, while the |
1847 | | - # outer maxnum combines retain the source's semantics. |
| 1843 | + # Use four 32-value vector reductions. MAXIMUMF propagates |
| 1844 | + # NaNs within each chunk, while the outer maxnum combines |
| 1845 | + # retain the required semantics. |
1848 | 1846 | new_row_max = cl.float32(-float("inf")) |
1849 | 1847 | for chunk in cl.static_iter(score_chunks): |
1850 | 1848 | chunk_max = chunk[:32].reduce( |
@@ -1878,18 +1876,16 @@ def _fmha_prefill_kernel( |
1878 | 1876 | stats_vec, |
1879 | 1877 | ) |
1880 | 1878 | # Acquire the first P-ready stage while the statistics |
1881 | | - # store is in flight. The second acquire and STTM wait |
1882 | | - # are split into the initial FMA lookahead below, matching |
1883 | | - # CUTLASS's machine schedule. |
| 1879 | + # store is in flight. The second acquire and STTM wait are |
| 1880 | + # split into the initial FMA lookahead below. |
1884 | 1881 | _wait_mbarrier( |
1885 | 1882 | p_empty.get_element_pointer((qid, 0)), |
1886 | 1883 | 1 ^ (pv_event & 1), |
1887 | 1884 | ) |
1888 | 1885 |
|
1889 | 1886 | block_sums = cl.Vector(0.0, 0.0, dtype=cl.float32) |
1890 | 1887 | minus_safe_max_scale = -safe_max * scale_softmax_log2 |
1891 | | - # Keep the DKG constexpr-loop topology while using flattened |
1892 | | - # local arrays in place of its mutable Python lists. |
| 1888 | + # Use constexpr loops with flattened local arrays. |
1893 | 1889 | if input_kind == ELEMENT_E4M3: |
1894 | 1890 | p_window_elems = 4 |
1895 | 1891 | fma_pipe_windows = 3 |
@@ -3211,10 +3207,10 @@ def benchmark( |
3211 | 3207 | return 0.0 |
3212 | 3208 |
|
3213 | 3209 | if use_cuda_graphs: |
3214 | | - # Match the CUTLASS testing helper's semantics: capture the complete |
3215 | | - # warmup and profiling loops on a non-default stream, replay each graph |
3216 | | - # once, and divide the profiled graph time by its launch count. Each |
3217 | | - # workspace has static graph arguments, so cold-L2 rotation is retained. |
| 3210 | + # Capture the complete warmup and profiling loops on a non-default |
| 3211 | + # stream, replay each graph once, and divide the profiled graph time by |
| 3212 | + # its launch count. Each workspace has static graph arguments, so |
| 3213 | + # cold-L2 rotation is retained. |
3218 | 3214 | workspace_index = 0 |
3219 | 3215 | if warmup_iterations: |
3220 | 3216 | warmup_graph = torch.cuda.CUDAGraph() |
|
0 commit comments