From e36da2efa9fa2889485ac5a1a38467b36cfc926d Mon Sep 17 00:00:00 2001 From: project_6 Date: Wed, 5 Aug 2026 03:35:06 +0000 Subject: [PATCH] [docs+code] lookback delay is a no-op on BI-V100 + V2 compound reduce pattern Two findings from reading CCCL source code: 1. single_pass_scan_operators.cuh: delay() has GridThreshold=500 gate. BI-V100 scan launches ~12 blocks (100K elements / tile_size). 12 < 500, so ALL delay policies collapse to __threadfence_block(). Conclusion: delay_ns, delay_l2w, delay_algorithm are IRRELEVANT on BI-V100. Only threads/items/load/scan algorithms matter. 2. summary_statistics.cu compound reduce pattern maps directly to paged_attention V2's cross-partition reduce. Updated muh_kernel_map.py with the structural mapping and the V2 dispatch bug (use_v1=True hardcoded in paged_attn.py line 99). Source: cccl_upstream/cub/cub/agent/single_pass_scan_operators.cuh cccl_upstream/thrust/examples/summary_statistics.cu --- docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md | 65 ++++++++++++++++++++++++++ muh_kernel_map.py | 33 ++++++++++++- 2 files changed, 96 insertions(+), 2 deletions(-) diff --git a/docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md b/docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md index 7e8b327b..911736a7 100644 --- a/docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md +++ b/docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md @@ -230,3 +230,68 @@ This means: - tuning_reduce.cuh: consider increasing items beyond scale_mem_bound cap for better ILP, especially for small types (fp16, int8) - tuning_scan.cuh: current values are correctly SMEM-bounded, don't increase + +--- + +## CCCL Lookback Delay Protocol (single_pass_scan_operators.cuh) + +> Added: 2026-08-04 + +### delay() has a GridThreshold gate — renders delay_ns IRRELEVANT on BI-V100 + +```cpp +template +void delay() { + if (Delay > 0) { + if (gridDim.x < GridThreshold) // <-- THIS IS THE KEY + __threadfence_block(); // small grid: just fence + else + __nanosleep(Delay); // large grid: actual sleep + } +} +``` + +GridThreshold defaults to 500. BI-V100 scan with 100K fp32 elements: + tile_size = 384 * 22 = 8448 + num_tiles = ceil(100000/8448) = 12 blocks + 12 << 500 → ALL delay calls reduce to __threadfence_block() + +This means: on BI-V100, the entire delay infrastructure (ns, dcid, l2w) +is a no-op. no_delay, fixed_delay(1904), exponential_backon_jitter(1904,830) +ALL execute the same __threadfence_block(). + +### Why our benchmark showed no_delay as "best" + +Not because no_delay is a better strategy, but because ALL strategies +produce identical machine code on a 12-block grid. The ~3% speedup +difference between dcid=0 and dcid=6 in bench_bi100.py is noise. + +### Impact on tuning_scan.cuh + +All scan delay parameters (delay_ns, delay_l2w, delay algorithm) can be +simplified to no_delay for BI-V100. The heuristic scaling (ns×0.5, l2w×0.6) +was both wrong AND irrelevant — the values don't matter because they're +never used as nanosleep arguments. + +The only scan tuning parameters that matter on BI-V100 are: + - threads_per_block (affects SMEM usage and occupancy) + - items_per_thread (affects SMEM usage and ILP) + - load_algorithm (WARP_TRANSPOSE vs DIRECT) + - scan_algorithm (RAKING vs WARP_SCANS) + - load_modifier (DEFAULT vs LDG) + +### summary_statistics.cu → paged_attention V2 compound reduce + +The Welford parallel merge in summary_statistics.cu is structurally +identical to paged_attention V2's cross-partition reduce: + +| summary_statistics | paged_attention V2 | +|---|---| +| summary_stats_data{n,min,max,mean,M2,M3,M4} | partition_result{max_logit, exp_sum, output_partial} | +| unary_op: x → {n=1, mean=x, M2=0, ...} | per-partition attention: Q@K^T → softmax → V·weights | +| binary_op: Welford parallel merge | online softmax merge: rescale by exp(old_max - new_max) | +| thrust::transform_reduce | DeviceReduce pass 2 | + +The compound accumulator size for V2 is sizeof(float)*3 = 12 bytes. +This affects tuning: scale_mem_bound(512, 16, 12) → different items/threads +than a simple float32 reduce. diff --git a/muh_kernel_map.py b/muh_kernel_map.py index dde0f152..52ee7297 100644 --- a/muh_kernel_map.py +++ b/muh_kernel_map.py @@ -187,9 +187,38 @@ VLLM_KERNEL_MAP = { }, "tuning_dimensions": { "NUM_THREADS": {"cccl_field": "threads_per_block"}, - "PARTITION_SIZE": {"cccl_field": "threads_per_block * items_per_thread"}, + "PARTITION_SIZE": {"cccl_field": "threads_per_block * items_per_thread", + "note": "hardcoded 512 in paged_attn.py, should be tunable"}, + }, + "cccl_pattern": "compound reduce: summary_statistics.cu Welford parallel merge pattern", + "cccl_parallel": { + "source": "thrust/examples/summary_statistics.cu", + "mapping": { + "summary_stats_data": "(max_logits, exp_sums, output) per partition", + "summary_stats_unary_op": "per-KV-block attention: Q@K^T → softmax → V weighted sum", + "summary_stats_binary_op": "cross-partition online softmax merge", + "thrust::transform_reduce": "DeviceReduce pass 2 merging partition results", + }, + "insight": "V2 reduce pass is structurally identical to CCCL compound reduce. " + "The accumulator is a 3-field struct (max, exp_sum, output_partial). " + "The binary op is the online softmax merge: " + "new_max = max(A.max, B.max), rescale exp_sums by exp(old_max - new_max), " + "merge weighted outputs. This is exactly the Welford parallel " + "variance pattern with different field semantics. " + "CCCL's AgentReduce handles compound structs natively — " + "the same tuning_reduce.cuh parameters apply, with accum_size = " + "sizeof(float32)*3 = 12 bytes (the compound accumulator).", + }, + "v2_dispatch_bug": { + "file": "paged_attn.py", + "line": 99, + "issue": "use_v1 = True hardcodes V1 for all seq_lens, disabling V2 entirely", + "impact": "For 100K token sequences, V1 makes one CTA iterate ALL KV blocks. " + "V2 would partition into PARTITION_SIZE chunks and reduce across partitions, " + "matching CCCL's two-pass GridEvenShare pattern.", + "fix": "Remove use_v1=True override. Use original heuristic: " + "V2 when max_seq_len > 8192 AND max_num_partitions > 1 AND num_seqs*num_heads <= 512", }, - "cccl_pattern": "reduce pass 1 (per-partition) + reduce pass 2 (cross-partition merge)", }, "context_attention_fwd": {