From 0ba2221d9c72472b5c08391b65738291131e2a26 Mon Sep 17 00:00:00 2001
From: dylanyunlon <dylanyunlon@users.noreply.github.com>
Date: Fri, 7 Aug 2026 03:19:39 +0000
Subject: [PATCH] =?UTF-8?q?docs:=20CCCL=20=E2=86=92=20vllm=20kernel=20patt?=
 =?UTF-8?q?ern=20mapping=20for=20competition?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Maps 5 CCCL patterns to actual vllm kernels:
1. Multi-field reduction → paged_attention (83% weight)
2. Prefix scan + transform → softmax in prefix_prefill
3. Transform → SiLU/GeLU/RMSNorm (14% weight)
4. TopK → sampling (via precompiled .so)
5. Flash attention → all patterns combined in triton kernel

Identifies 30 competition-critical files from 5205 in cccl_upstream.
---
 CCCL_PATTERN_MAP.md | 104 ++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 104 insertions(+)
 create mode 100644 CCCL_PATTERN_MAP.md

diff --git a/CCCL_PATTERN_MAP.md b/CCCL_PATTERN_MAP.md
new file mode 100644
index 00000000..d6681176
--- /dev/null
+++ b/CCCL_PATTERN_MAP.md
@@ -0,0 +1,104 @@
+# CCCL → vllm Kernel Pattern Mapping
+## BI-V100 Competition Reference
+
+### Pattern 1: Multi-field Reduction (paged_attention)
+
+**CCCL source**: `thrust/examples/bounding_box.cu`, `summary_statistics.cu`
+**vllm kernel**: `paged_attn.py` → ixformer paged_attention_v1/v2
+
+```
+CCCL:  transform_reduce(begin, end, unary_op, init, binary_op)
+vllm:  for each KV block: score = Q·K, max_score = reduce_max, exp_sum = reduce_sum
+```
+
+**Tuning surface**:
+- `_PARTITION_SIZE`: controls how many KV tokens per CTA in V2 mode
+- V1/V2 dispatch threshold: `total_tiles vs 2 × sm_count`
+- BI-V100: 16 SMs → V2 beneficial when seq_len > 1024 (2 waves of 16 CTAs × 512 partition)
+
+**CCCL parameter**: `ReducePassPolicy{threads=512, items=24, vec=2, WARP_REDUCTIONS, LDG}`
+
+### Pattern 2: Prefix Scan + Transform (softmax)
+
+**CCCL source**: `thrust/examples/simple_moving_average.cu`, `cub/benchmarks/bench/scan/exclusive/sum.cu`
+**vllm kernel**: `prefix_prefill.py` context_attention_fwd_kernel
+
+```
+CCCL:  inclusive_scan(begin, end, output, plus<float>)
+vllm:  for each BLOCK_N chunk: qk = Q·K, m_new = max(m_old, max(qk)),
+       l_new = l_old * exp(m_old - m_new) + sum(exp(qk - m_new))
+```
+
+**Tuning surface**:
+- `BLOCK_M`: Q tile rows (32 or 64 for BI-V100)
+- `BLOCK_N`: K/V sweep width (32 or 64)
+- `NUM_WARPS`: 4 (16 SMs don't benefit from 8 warps per CTA)
+- `num_stages`: 1 (no cp.async) or 2 (software pipeline)
+
+**CCCL parameter**: `ScanLookbackPolicy{threads=384, items=22, WARP_TRANSPOSE, DEFAULT, WARP_SCANS, {backon_jitter_window, 952, 415}}`
+
+### Pattern 3: Transform (activation functions)
+
+**CCCL source**: `cub/benchmarks/bench/transform/babelstream.cu`
+**vllm kernel**: Triton SiLU, GeLU, RMSNorm kernels (via `_custom_ops.py`)
+
+```
+CCCL:  transform(begin, end, output, silu_op)    // x * sigmoid(x)
+vllm:  @triton.jit def silu_kernel(x): tl.sigmoid(x) * x
+```
+
+**Tuning surface**:
+- `bytes_in_flight`: 64KB on BI-V100 (56 GB/s per-SM × 1100ns latency)
+- Triton `num_stages=2` maps to BIF=64KB (2× prefetch window)
+- `SMEM = 49152` (fixed by _custom_ops.py)
+
+**CCCL parameter**: `TransformPrefetchPolicy{threads=256, bif=64KB, prefetch_stride=128}`
+
+### Pattern 4: TopK (sampling)
+
+**CCCL source**: `cub/benchmarks/bench/topk/keys.cu`
+**vllm kernel**: sampling_kernels (precompiled .so)
+
+```
+CCCL:  DeviceTopk::TopK(keys, k, output)
+vllm:  ixformer topk_sampling → radix_sort + select partial
+```
+
+**Tuning surface** (via .so, limited):
+- `bits_per_pass`: 11 for float32 (32 bits / 3 passes)
+- Thread count: 512 (baked into .so)
+
+### Pattern 5: Triton Flash Attention (all patterns combined)
+
+**CCCL source**: All of the above + `cub/agent/agent_scan.cuh` union SMEM model
+**vllm kernel**: `triton_flash_attention.py`
+
+```
+Q_resident × K_streaming × softmax_online → Output
+= transform_reduce (Q·K) + scan (softmax) + transform (V matmul)
+```
+
+**Tuning surface**: 17 existing + 19 new autotune configs from gen_config.py
+**Key configs for BI-V100**:
+```python
+# Best for long context (seq_len > 4096):
+Config(BLOCK_M=64, BLOCK_N=64, num_warps=4, num_stages=2)  # 40KB SMEM, 1 CTA/SM
+
+# Best for short context (seq_len < 1024):
+Config(BLOCK_M=32, BLOCK_N=32, num_warps=2, num_stages=2)  # 32KB SMEM, 2 CTAs/SM
+```
+
+---
+
+### CCCL Asset Utilization Summary
+
+| CCCL Asset | Files | Used for BI-V100 | Competition Impact |
+|-----------|-------|-------------------|-------------------|
+| Tuning headers (26) | 18094 lines | 3568 lines (20%) | P0: reduce/scan/transform |
+| CUB benchmarks (80) | reduce/scan/topk/transform | benchmark framework | P0: parameter search |
+| Thrust examples (52) | summary_stats/bounding_box/norm | pattern mapping | P1: architecture understanding |
+| CUB tests (243) | correctness verification | 0% (need BI-V100) | P2: correctness |
+| libcudacxx (1463) | type traits, atomics | implicit (via CUB) | Infra |
+
+**Total usable CCCL assets**: 5205 files in cccl_upstream
+**Competition-critical subset**: ~30 files (5 tuning headers + 10 benchmarks + 15 examples)