under test, not sure no errors

2026-09-02 07:03:56 +00:00
commit 43c43b491c
4211 changed files with 1013777 additions and 0 deletions
--- a/docs/CCCL_BENCHMARK_REFERENCE.md
+++ b/docs/CCCL_BENCHMARK_REFERENCE.md
@@ -0,0 +1,136 @@
+# CCCL benchmark reference data — extracted from 27 tuning headers
+
+> Auto-extracted from `cccl_upstream/cub/cub/device/dispatch/tuning/*.cuh`
+> 199 benchmark annotations across 27 files, 286 template specializations
+> This is the data NVIDIA spent millions of GPU-hours generating on A100/H100/B200.
+> muh needs equivalent data for BI-V100.
+
+## Summary
+
+| Algorithm | CCCL lines | muh lines | muh/CCCL | Benchmarks | Specializations | Priority |
+|-----------|-----------|-----------|----------|------------|-----------------|----------|
+| radix_sort | 2382 | 223 | 9% | 70 | 0 | P0 — sampling hot path |
+| select_if | 2730 | 460 | 17% | 0* | 37 | P0 — top-p filtering |
+| scan_by_key | 2009 | 146 | 7% | 30 | 58 | P1 — softmax denominator |
+| reduce_by_key | 1736 | 172 | 10% | 32 | 63 | P1 — score aggregation |
+| unique_by_key | 1540 | 167 | 11% | 29 | 58 | P1 — KV cache dedup |
+| scan | 1526 | 371 | 24% | 16 | 12 | P0 — prefix sum hot path |
+| three_way_partition | 789 | 100 | 13% | 0 | 0 | P1 — token classification |
+| rle_non_trivial_runs | 692 | 69 | 10% | 8 | 15 | P1 — attention mask |
+| segmented_sort | 641 | 190 | 30% | 0 | 0 | P1 — per-seq token ranking |
+| rle_encode | 627 | 64 | 10% | 4 | 15 | P1 — mask compression |
+| transform | 550 | 186 | 34% | 0 | 0 | P1 — RMSNorm/SiLU/RoPE |
+| reduce | 479 | 298 | 62% | 6 | 2 | P0 — attention score reduce |
+| histogram | 364 | 77 | 21% | 4 | 3 | P1 — repetition penalty |
+| topk | 122 | 114 | 93% | 0 | 0 | P0 — sampling core |
+
+*select_if has 37 muh benchmark annotations from the 3-dimension restore fix
+
+## Decode hot path — benchmark reference values
+
+### reduce (Output TPS × 16.796 = 83% weight)
+
+CCCL SM100 benchmarks (the target we need to match or beat on BI-V100):
+
+```
+# float32, offset=4, accum=4:
+ipt_16.tpb_512.ipv_2  1.061295  1.000000  1.065478  1.167139   geo=1.072
+
+# float64, offset=4, accum=8:
+ipt_16.tpb_640.ipv_1  1.017834  1.000000  1.015835  1.057092   geo=1.023
+
+# int64, offset=4, accum=8:
+ipt_15.tpb_512.ipv_2  1.019887  1.000000  1.017636  1.058036   geo=1.024
+
+# int64, offset=8, accum=8:
+ipt_15.tpb_512.ipv_1  1.019414  1.000000  1.017218  1.057143   geo=1.023
+
+# Deterministic float32 (SM90):
+ipt_13.tpb_224        1.107188  1.009709  1.097114  1.316820   geo=1.127
+
+# Deterministic float64 (SM86):
+ipt_11.tpb_128        1.232089  1.002124  1.245336  1.582279   geo=1.250
+```
+
+Current muh BI-V100 values (theoretical, NOT benchmarked):
+- float32: items=24, threads=512, vec=2  (CCCL SM100: items=16, threads=512, vec=2)
+- float64: items=16, threads=384, vec=2  (CCCL SM100: items=16, threads=640, vec=1)
+- det float32: items=32, threads=384     (CCCL SM90: items=13, threads=224)
+- det float64: items=16, threads=384     (CCCL SM86: items=11, threads=128)
+
+**Critical gap**: muh items are 1.5-2.5× CCCL SM100 values. Rationale was "16 SMs need
+larger tiles to compensate for fewer CTAs." This MUST be validated on real hardware.
+If register pressure causes occupancy drop, the 2.5× items advantage evaporates.
+
+### scan (softmax denominator, TTFT impact)
+
+CCCL SM100 benchmarks with full delay tuning:
+
+```
+# int8, offset=4:
+ipt_18.tpb_512.ns_768.dcid_7.l2w_820.trp_1.ld_0   1.189  1.006  1.173  1.305   geo=1.163
+
+# int16, offset=4:
+ipt_13.tpb_512.ns_1384.dcid_7.l2w_720.trp_1.ld_0  1.128  1.003  1.120  1.308   geo=1.135
+
+# float32, offset=4:
+ipt_22.tpb_384.ns_1904.dcid_6.l2w_830.trp_1.ld_0  1.148  0.997  1.140  1.463   geo=1.182
+
+# float32, offset=8:
+ipt_19.tpb_416.ns_956.dcid_7.l2w_550.trp_1.ld_1   1.146  0.994  1.137  1.456   geo=1.178
+
+# float64, offset=4:
+ipt_23.tpb_416.ns_772.dcid_5.l2w_710.trp_1.ld_0   1.089  1.016  1.086  1.265   geo=1.111
+
+# float64, offset=8:
+ipt_22.tpb_320.ns_328.dcid_2.l2w_965.trp_1.ld_0   1.080  1.000  1.076  1.249   geo=1.100
+
+# SM90 int128:
+tpb_576.ipt_21.ns_860.l2w_630                      (no speedup data in comment)
+```
+
+Key tuning dimensions absent from muh:
+- `dcid` (delay constructor ID): 8 variants (0-7), each a different backoff strategy
+- `l2w` (L2 write latency in ns): BI-V100 L2=6MB vs SM100 L2=50MB — needs recalibration
+- `ns` (delay in nanoseconds): range 64-2044ns across all scan benchmarks
+- `trp` (transpose): 0=DIRECT, 1=WARP_TRANSPOSE
+- `ld` (load modifier): 0=LOAD_DEFAULT, 1=LOAD_CA/LOAD_LDG
+
+### radix_sort (70 benchmarks — most data-rich algorithm)
+
+Top-performing SM100 configurations:
+
+```
+# Large key (8B), offset=4:
+ipt_14.tpb_320  1.256  1.000  1.228  1.487   geo=1.231
+
+# Small key (1B), offset=4:
+ipt_20.tpb_512  1.013  0.968  1.016  1.048   geo=1.011
+
+# Medium key (4B), offset=4:
+ipt_21.tpb_512  1.003  0.995  1.004  1.019   geo=1.005
+```
+
+Qwen3.6 sampling: vocab_size=152064, logits are float32 (4B keys).
+Bits per pass: sizeof(float32)=4 → bits_per_pass=11 → ⌈32/11⌉=3 passes.
+Each pass: 2^11=2048 histogram bins × sizeof(int)=4 = 8KB SMEM for histogram.
+Total sort SMEM ≈ 8KB + threads×items×sizeof(float32) staging.
+
+## Delay algorithms reference (for lookback-based algorithms)
+
+| dcid | Algorithm | Description |
+|------|-----------|-------------|
+| 0 | no_delay | No delay between lookback iterations |
+| 1 | fixed_delay | Fixed ns delay |
+| 2 | exponential_backoff | Double delay each retry |
+| 3 | exponential_backoff_jitter | Backoff + random jitter |
+| 4 | exponential_backoff_jitter_window | Backoff + jitter + window |
+| 5 | exponential_backon_jitter_window | Increase delay (backon) + jitter + window |
+| 6 | exponential_backon_jitter | Increase delay + jitter |
+| 7 | exponential_backon | Increase delay monotonically |
+
+BI-V100 implications:
+- L2 cache 6MB (SM100: 50MB) → tile_state fits in L2 for fewer concurrent CTAs
+- 16 SMs → max 32 concurrent tiles → lower contention → shorter delays likely optimal
+- bandwidth_per_SM=56GB/s (SM100: 54GB/s) → similar per-SM behavior
+- Recommended starting point: dcid=7 (exponential_backon) with ns×0.5, l2w×0.6 scaling
--- a/docs/CCCL_ENGINEX_ARCHITECTURE_ALIGNMENT.md
+++ b/docs/CCCL_ENGINEX_ARCHITECTURE_ALIGNMENT.md
@@ -0,0 +1,80 @@
+# CCCL ↔ EngineX Architecture Alignment
+
+## Executive Summary
+
+EngineX ships precompiled `.so` kernels — **zero `.cu` source files** are available.
+The optimization surface is Python runtime params + Triton JIT kernels.
+
+CCCL's value is NOT parameter values. It's the **architectural patterns** that
+tell us which parameters matter, what their constraints are, and why.
+
+## Three-Layer Architecture Mapping
+
+### CCCL Layer → EngineX Layer → What We Control
+
+| CCCL | EngineX | Controllable? |
+|------|---------|--------------|
+| `dispatch_reduce.cuh` (GridEvenShare work distribution) | `paged_attn.py` (V1/V2 dispatch, _PARTITION_SIZE) | **Yes** — Python runtime |
+| `kernel_reduce.cuh` (kernel entry, atomic vs 2-phase) | `_C_flashattention.so` (paged_attention_v1/v2) | **No** — precompiled |
+| `agent_reduce.cuh` (tile consumption, vectorized load) | internal to `.so` | **No** — precompiled |
+| `tuning_reduce.cuh` (policy_selector) | `_custom_ops.py` (SMEM=49152) | **Partially** — SMEM limit |
+| `dispatch_transform.cuh` (spread_out_items_per_thread) | `rmsnorm_kernels.py` (BLOCK_SIZE heuristic) | **Yes** — Triton autotune |
+| `kernel_scan.cuh` (lookback/lookahead) | `prefix_prefill.py` (BLOCK_M/N, num_warps) | **Yes** — Triton config |
+| `dispatch_scan.cuh` (tile init + scan kernel) | `triton_splitk.py` (split-K attention) | **Yes** — Triton config |
+
+### Key CCCL Patterns We Apply
+
+1. **GridEvenShare** (`grid_even_share.cuh`):
+   - `max_blocks = sm_occupancy × sm_count × subscription_factor`
+   - BI-V100: 1 × 16 × 5 = 80 max CTAs
+   - Applied to: `paged_attn.py` _BI100_TARGET_TILES, V1/V2 threshold
+
+2. **Compound Reduce** (`summary_statistics.cu`):
+   - Accumulator = struct{m, l, o} (max, sum_exp, weighted_output)
+   - unary_op: score_tile → partial softmax stats
+   - binary_op: online softmax merge with correction factor
+   - Applied to: `_forward_decode_pytorch` online softmax loop
+
+3. **Two-Phase Reduce** (`kernel_reduce.cuh`):
+   - Phase 1: each CTA reduces a partition → `d_block_reductions[blockIdx.x]`
+   - Phase 2: single CTA reduces all block results
+   - Applied to: paged_attention_v2 partition → merge
+
+4. **spread_out_items_per_thread** (`dispatch_transform.cuh`):
+   - Reduce items/thread when there aren't enough items to fill all SMs
+   - `items = min(max, ceil_div(N, sm_count × threads × occupancy))`
+   - Applied to: Triton kernel BLOCK_SIZE selection
+
+5. **Lookback Delay** (`tuning_scan.cuh`):
+   - 16 SMs → ~32 concurrent CTAs → tile_state fits in 6MB L2
+   - Inter-CTA contention near zero → no_delay optimal
+   - Applied to: scan-based operations (softmax denominator)
+
+## BI-V100 Hardware Profile (Confirmed)
+
+| Property | Value | Impact |
+|----------|-------|--------|
+| SM count | 16 | 3.1x fewer CTAs than spec (50) → larger tiles per CTA |
+| SMEM | 48KB | Same as NVIDIA → CCCL SMEM constraints apply directly |
+| HBM BW | 900 GB/s | BW/SM = 56 GB/s ≈ B200 level → bytes_in_flight = 64KB |
+| L2 cache | 6MB | 8.3x smaller than SM100 → faster coherence, no_delay wins |
+| Warp size | 32 | Same as NVIDIA → CCCL warp-level primitives work |
+
+## Files Inventory
+
+### Precompiled (CANNOT modify)
+- `_C_flashattention.so` — paged_attention_v1, paged_attention_v2, reshape_and_cache
+- `_C.so` — xformers attention backends
+- `libtriton.so` — Triton compiler/runtime
+
+### Triton JIT (CAN modify)
+- `pkgs/triton/ops/flash_attention.py` — Flash Attention (head_dim ≤ 128 only)
+- `pkgs/xformers/ops/fmha/triton_splitk.py` — Split-K attention (V2 pattern)
+- `pkgs/xformers/ops/triton/rmsnorm_kernels.py` — RMSNorm
+- `pkgs/xformers/ops/triton/rope_padded_kernels.py` — RoPE
+
+### Python runtime (CAN modify)
+- `paged_attn.py` — V1/V2 dispatch, _PARTITION_SIZE, decode fallback
+- `prefix_prefill.py` — Prefill attention BLOCK_M/N/NUM_WARPS
+- `vllm/_custom_ops.py` — SMEM=49152 (already fixed from 32768)
+- `computility-run.yaml` — Server launch params
--- a/docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md
+++ b/docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md
@@ -0,0 +1,412 @@
+# CCCL Reduce Architecture Notes
+
+> Source: `dispatch_reduce.cuh`, `kernel_reduce.cuh`, `agent_reduce.cuh`, `tuning_reduce.cuh`, `util_arch.cuh`
+> Read: 2026-08-04 by Claude from CCCL upstream in project_6/cccl_upstream/
+
+## Key Architecture
+
+### Two-pass dispatch (dispatch_reduce.cuh)
+
+```
+num_items <= single_tile.threads * single_tile.items
+  → SingleTile: one CTA, one kernel launch
+  → DeviceReduceSingleTileKernel(d_in, d_out, num_items, ...)
+
+num_items > single_tile threshold
+  → Pass 1: DeviceReduceKernel — N CTAs each reduce their share → d_block_reductions[N]
+  → Pass 2: DeviceReduceSingleTileKernel — 1 CTA reduces d_block_reductions[N] → d_out
+```
+
+Grid size for Pass 1: `max_blocks = sm_occupancy * sm_count * subscription_factor(5)`
+For BI-V100: `2 * 16 * 5 = 160 blocks` max.
+Each block processes `ceil(num_items / 160)` elements.
+
+### Tile consumption (agent_reduce.cuh)
+
+**Critical: tile data is in registers, NOT SMEM.**
+
+```cpp
+AccumT items[ITEMS_PER_THREAD];  // <-- register array, per-thread
+// ... load from global memory ...
+thread_aggregate = ThreadReduce(items, reduction_op);  // per-thread reduction
+
+// Only SMEM used:
+BlockReduce(temp_storage.reduce).Reduce(thread_aggregate, reduction_op);
+```
+
+`TempStorage` = `BlockReduce::TempStorage` ≈ threads * sizeof(AccumT) bytes.
+NOT threads * items * sizeof(AccumT).
+
+### Vectorized loads
+
+```cpp
+ATTEMPT_VECTORIZATION = (vec_size > 1) && (ITEMS_PER_THREAD % vec_size == 0)
+    && is_pointer<InputIteratorT>
+    && (is_primitive<InputT> || is_trivially_relocatable<InputT>)
+    && sizeof(InputT) <= 8;
+```
+
+For fp32 scores: vec_size=2 → loads 8 bytes (2 floats) per instruction.
+For fp16 KV cache: vec_size=4 → loads 8 bytes (4 halfs) per instruction.
+
+### scale_mem_bound vs scale_reg_bound (util_arch.cuh)
+
+Two scaling functions with different constraints:
+
+**scale_mem_bound** (memory-bound algorithms: reduce, transform):
+- items = clamp(nominal * 4 / type_size, 1, nominal * 2)  ← allows 2x expansion
+- threads = min(nominal, round_up(48KB / (type_size * items), 32))
+
+**scale_reg_bound** (register-bound algorithms: scan with complex state):
+- items = max(1, nominal * 4 / max(4, type_size))  ← no expansion past nominal
+- threads = min(nominal, ceil_div(48KB / (type_size * items), 32) * 32)
+
+Key difference: scale_reg_bound uses `max(4, type_size)` preventing items from exceeding nominal for small types, and uses `ceil_div` instead of `round_up` for thread count. Both use 48KB as the cap, but this limits REGISTER PRESSURE (spill to local memory), not actual SMEM usage.
+
+## Impact on muh tuning
+
+### Our SMEM model was wrong for reduce
+
+`test_smem_safety.py` and `check_smem()` in `muh_kernel_map.py` compute
+`tile_bytes = threads * items * type_size` and check against 49152.
+
+This is the scale_mem_bound cap, NOT the actual SMEM usage. The actual SMEM
+for reduce is approximately `threads * max(sizeof(AccumT), 4)` bytes — about
+2-8 KB, not 32-49 KB.
+
+CCCL's SM100 float64 tuning uses `threads=640, items=16` → scale_mem_bound
+"tile" = 640*16*8 = 81920 > 49152. But this doesn't overflow SMEM — it only
+means scale_mem_bound will cap threads down. The actual kernel SMEM usage
+with threads=640 is only ~5120 bytes.
+
+### Our float64/int64 tuning may be too conservative
+
+We use threads=384 items=16 for float64, capped by scale_mem_bound. CCCL
+uses threads=640 items=16 on SM100. The question is whether BI-V100's register
+file (255 regs/thread) can hold 16 float64 items without spilling.
+
+16 * 8 = 128 bytes = 32 registers per thread for tile data alone.
+With overhead (thread_aggregate, loop variables, etc.), ~40 registers/thread.
+255 max registers → no spill risk. threads=640 may be safe on BI-V100.
+
+**TODO**: Benchmark threads=640 items=16 for float64 on BI-V100.
+
+### paged_attn.py forces V1
+
+Line 99: `use_v1 = True` overrides V1/V2 heuristic. V2 is completely disabled.
+For 100K token sequences, V1 makes one CTA iterate over all KV blocks — bad
+for latency. V2 would partition the work and reduce across partitions, which
+is exactly CCCL's two-pass pattern.
+
+**TODO**: Re-enable V2 for max_seq_len > 8192. Use muh's partition_size tuning.
+
+### _PARTITION_SIZE = 512 is hardcoded
+
+Not controlled by muh. Should be tunable: larger partition = fewer blocks =
+less overhead but more work per block. Optimal value depends on SM count.
+For 16 SMs: partition_size=1024 may be better (fewer partitions to reduce).
+
+---
+
+## CCCL Scan Architecture (dispatch_scan.cuh)
+
+> Added: 2026-08-04
+
+### Two algorithm paths
+
+**Lookback** (all GPUs including BI-V100):
+- Each CTA processes one tile, uses `ScanTileState` in global memory for inter-CTA communication
+- Lookback delay policy controls how aggressively CTAs poll predecessors
+- SMEM: static only (`__shared__`), passed as `0` dynamic SMEM
+- BI-V100 optimal: `no_delay` (dcid=0) because 16 SMs → ~32 CTAs → tile_status fits in 6MB L2
+
+**Lookahead** (SM100+ only, PTX ISA >= 860):
+- Pipeline-based with `__pipeline_memcpy_async` and bulk copy
+- Uses dynamic SMEM with auto-selected `num_stages`
+- **Not available on BI-V100** — requires NVIDIA PTX ISA 860+ instructions
+- All lookahead structs in our tuning_scan.cuh can remain empty shells
+
+### ScanTileState allocation
+
+Scan requires `d_temp_storage` for tile status descriptors:
+```
+tile_size = threads * items
+num_tiles = ceil(num_items / tile_size)
+temp_bytes = tile_state.AllocationSize(num_tiles)
+```
+
+For BI-V100 with 100K tokens and tile_size=384*22=8448:
+num_tiles = ceil(100000/8448) = 12 tiles → negligible temp storage.
+
+### Grid size for scan
+
+Lookback scan launches `num_tiles` blocks (one per tile), NOT `sm_count * subscription_factor`.
+This is different from reduce, which uses `GridEvenShare`.
+For scan, every CTA processes exactly one tile and communicates with neighbors.
+
+With 12 tiles on 16 SMs: all tiles fit in one wave, zero lookback contention.
+This is why `no_delay` works on BI-V100 — the entire scan completes in a single wave.
+
+### Lookahead num_stages optimization (SM100 only)
+
+CCCL dynamically selects pipeline depth:
+```cpp
+max_stages = ceil(num_items / (sm_count * tile_size)) + 1
+while (smem_for_stages(num_stages+1) <= max_dynamic_smem) num_stages++
+```
+
+For BI-V100 this is irrelevant (no pipeline support), but the formula shows
+NVIDIA's strategy: match pipeline depth to problem size / SM count ratio.
+
+---
+
+## CCCL Scan Agent Architecture (agent_scan.cuh)
+
+> Added: 2026-08-04
+
+### Critical difference from reduce: scan DOES use SMEM for tile data
+
+```cpp
+union _TempStorage {
+    BlockLoadT::TempStorage load;       // SMEM for WARP_TRANSPOSE load
+    BlockStoreT::TempStorage store;     // SMEM for WARP_TRANSPOSE store
+    struct {
+        TilePrefixCallbackOpT::TempStorage prefix;  // lookback state
+        BlockScanT::TempStorage scan;                // block scan
+    } scan_storage;
+};
+```
+
+This is a **union** — load, store, and scan share the same SMEM, used
+in phases separated by `__syncthreads()`. Actual SMEM = max of three.
+
+For `BLOCK_LOAD_WARP_TRANSPOSE`:
+  load_smem ≈ threads * items * sizeof(InputT)
+
+For `BlockScan`:
+  scan_smem ≈ threads * sizeof(AccumT) + prefix_callback
+
+The dominant term is load/store: threads * items * type_size.
+
+**Conclusion: our SMEM constraint `threads * items * type_size ≤ 48KB`
+is CORRECT for scan but WRONG (overly conservative) for reduce.**
+
+### Tile processing flow
+
+```
+1. BlockLoad(SMEM).Load(d_in + offset, items[ITEMS_PER_THREAD])
+2. __syncthreads()
+3. BlockScan(SMEM).Scan(items, ..., prefix_op)  // lookback here
+4. __syncthreads()
+5. BlockStore(SMEM).Store(d_out + offset, items)
+```
+
+Each CTA processes exactly one tile (tile_idx = start_tile + blockIdx.x).
+Inter-CTA communication happens in step 3 via TilePrefixCallbackOp,
+which reads predecessor tile states from global memory (the lookback).
+
+### Lookback protocol (TilePrefixCallbackOp)
+
+For tile k, the callback:
+1. Sets own tile state to PARTIAL with local aggregate
+2. Looks back at tiles k-1, k-2, ... until finding an INCLUSIVE prefix
+3. Combines found prefix with local aggregate → own INCLUSIVE prefix
+4. Sets own tile state to INCLUSIVE
+
+The LookbackDelayPolicy controls how aggressively step 2 polls:
+- no_delay: spin immediately (best when few CTAs, e.g., BI-V100 16 SMs)
+- exponential_backon: exponentially increase delay between polls
+  (best when many CTAs compete for L2 coherence, e.g., SM100 148 SMs)
+
+### Impact on muh tuning
+
+For reduce: items_per_thread can be larger because SMEM only stores
+~threads*4 bytes for BlockReduce. The 48KB cap prevents register spill.
+
+For scan: items_per_thread is genuinely SMEM-limited because
+BlockLoad/BlockStore use threads*items*type_size bytes of SMEM.
+
+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 <int Delay, unsigned int GridThreshold = 500>
+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.
+
+---
+
+## CCCL Transform Architecture (tuning_transform.cuh)
+
+> Added: 2026-08-04
+
+### 4 algorithms, only 2 available on BI-V100
+
+| Algorithm | Requirement | BI-V100 |
+|---|---|---|
+| prefetch | universal | ✓ available |
+| vectorized | contiguous + trivially_relocatable + power-of-2 size | ✓ available |
+| ldgsts | SM80+ cp.async (NVIDIA-specific PTX) | ✗ |
+| ublkcp | SM90+ bulk copy (NVIDIA-specific PTX) | ✗ |
+
+### cc_to_min_bytes_in_flight — the correct value for BI-V100
+
+CCCL's hardcoded mapping:
+  B200  (SM100, 54 GB/s/SM): 64KB
+  H200  (SM90,  25 GB/s/SM): 48KB
+  A100  (SM80,  19 GB/s/SM): 16KB
+  V100  (SM70,  11 GB/s/SM): 12KB
+
+BI-V100 (56 GB/s/SM) is closest to B200. Our 64KB is aggressive but
+bench_bi100 confirms bif=8 (64KB) dominates bif=0 (32KB). So 64KB stands.
+
+However: bytes_in_flight only affects the PREFETCH algorithm path.
+For vllm's RMSNorm/SiLU/RoPE (contiguous fp16 arrays), the VECTORIZED
+path is selected instead, where bytes_in_flight is ignored and
+items_per_thread is set directly.
+
+### vectorized policy selection for BI-V100
+
+CCCL's tuned_vectorized_policy for fallback (cc < 8.0):
+  TransformVectorizedPolicy{256, 8, 4}  // 256 threads, 8 items, vec=4
+
+For RMSNorm with fp16 (store_size=2):
+  items_per_thread=8, vec_size=4 → 8 elements/thread, 4 per vector load
+  tile = 256 * 8 = 2048 elements per CTA
+  With 16 SMs × 2 occupancy = 32 CTAs → 65536 elements/wave
+
+Qwen3.6 hidden_size=3584 → RMSNorm processes 3584 elements.
+3584 / 2048 = 2 tiles → fits in one wave on BI-V100. Good.
+
+### Impact on our tuning_transform.cuh
+
+Our bi100_bytes_in_flight=64KB is correct for prefetch but irrelevant
+for vectorized. We should also set vectorized policy parameters directly:
+  threads=256, items=8, vec=4 (CCCL default for older arch)
+  OR test threads=128 with items=16 (A100 triad tuning) for higher ILP.
+
+Benchmark result (bif=8, alg=1, pref=2, tpb=256, unrl=1, vsp2=1):
+  alg=1 = vectorized (confirmed — prefetch would be alg=0)
+  tpb=256 = matches CCCL default
+  vsp2=1 → vec_size parameter (powers of 2, so vsp2=1 means vec_size=2)
+  speedups: 1.203199 1.058919 1.019168 (fp16, 1M/16M/64M)
+
+---
+
+## GridEvenShare Work Distribution (grid_even_share.cuh)
+
+> Added: 2026-08-04
+
+### Two strategies: RAKE vs STRIP_MINE
+
+**RAKE** (scan uses this): consecutive tiles per block
+  block k gets tiles [k*avg .. k*avg + avg-1]
+  block_stride = TILE_ITEMS (contiguous, no gaps)
+
+**STRIP_MINE** (reduce uses this): interleaved tiles
+  block k gets tiles k, k+grid_size, k+2*grid_size, ...
+  block_stride = grid_size * TILE_ITEMS (strided)
+
+### Concrete numbers for BI-V100 attention score reduce
+
+tile_items = 512 * 24 = 12288 (bi100_plus_float32_o4)
+max_grid = 2 * 16 * 5 = 160 (occupancy * SMs * subscription)
+
+| seq_len | total_tiles | grid_size | tiles/block | waves |
+|---------|-------------|-----------|-------------|-------|
+| 1K      | 1           | 1         | 1           | 1     |
+| 8K      | 1           | 1         | 1           | 1     |
+| 32K     | 3           | 3         | 1           | 1     |
+| 100K    | 9           | 9         | 1           | 1     |
+| 1M      | 82          | 82        | 1           | 3     |
+
+Even at 100K tokens, only 9 CTAs are needed → everything fits in one
+wave on 16 SMs. This means:
+
+1. Reduce tuning (items/threads) matters less than expected — there 
+   are so few tiles that the per-tile overhead dominates, not throughput.
+
+2. The V1/V2 choice in paged_attn.py matters MORE — V1 doesn't use
+   GridEvenShare at all, it's a single CTA iterating sequentially.
+   V2's partition-based approach enables parallel reduction.
+
+3. For short sequences (≤8K, 1 tile), SingleTile path triggers:
+   just 1 CTA, 1 kernel launch, no temp storage.
+
+### The "big shares" distribution
+
+GridEvenShare handles uneven tile counts:
+  avg_tiles_per_block = total_tiles / grid_size
+  big_shares = total_tiles % grid_size (blocks that get +1 tile)
+
+For 100K tokens with 9 tiles and 9 blocks: avg=1, big_shares=0.
+All blocks equal. No imbalance.
+
+For 1M tokens with 82 tiles and 82 blocks: avg=1, big_shares=0.
+Still perfectly balanced at 1 tile/block.
+
+Only when max_grid_size limits grid_size do we get imbalance:
+e.g., 200 tiles with max_grid=160 → avg=1, big_shares=40 (40 blocks
+get 2 tiles, 120 blocks get 1 tile).
--- a/docs/CCCL_TO_TRITON_METHODOLOGY.md
+++ b/docs/CCCL_TO_TRITON_METHODOLOGY.md
@@ -0,0 +1,130 @@
+# CCCL → Triton 方法论迁移
+
+> 核心观点: CCCL 的 policy_selector 和 %RANGE% benchmark 框架是 NVIDIA 几十年 GPU 性能优化的结晶。
+> 竞赛中几万人都在用同一份 EngineX 代码调参数。我们的差异化来自 CCCL 的方法论——不是复制参数，是复制思维方式。
+
+---
+
+## 一、CCCL 的 tuning 方法论
+
+NVIDIA 在 CCCL 中的参数搜索基础设施:
+
+```
+%RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1      ← 每线程处理的元素数
+%RANGE% TUNE_THREADS_PER_BLOCK tpb 128:1024:32 ← 每 CTA 的线程数
+%RANGE% TUNE_ITEMS_PER_VEC_LOAD_POW2 ipv 1:2:1 ← 向量化加载宽度
+```
+
+这些 %RANGE% 注释由 CCCL 的 benchmark runner 读取，生成笛卡尔积，每个组合跑 4 个 problem size，输出:
+```
+ipt_22.tpb_384.ns_1904.dcid_6.l2w_830.trp_1.ld_0 1.148 0.997 1.140 1.463
+```
+
+选几何均值最高的组合写入 policy_selector。
+
+**关键**: 不是人类凭经验猜参数，是系统化的笛卡尔积搜索 + 实测数据驱动。
+
+---
+
+## 二、EngineX Triton 的参数对应
+
+### 2.1 prefix_prefill.py (Context Attention — 影响 Input TPS 14%)
+
+CCCL scan 的 %RANGE%:
+```
+ipt 7:24:1      → Triton: BLOCK_M ∈ {32, 64, 128, 256}
+tpb 128:1024:32 → Triton: num_warps ∈ {2, 4, 8, 16} (warps × 32 = threads)
+ns 0:2048:4     → BI-V100 不适用 (Triton 没有 delay policy)
+trp 0:1:1       → BI-V100 不适用 (Triton 自动选择 memory layout)
+ld 0:1:1        → BI-V100 不适用 (Triton 自动选择 cache modifier)
+```
+
+| CCCL 参数 | Triton 参数 | 当前值 | 搜索范围 |
+|-----------|-----------|--------|---------|
+| items_per_thread | BLOCK_M (和 BLOCK_N) | 128 or 64 | {32, 64, 128} |
+| threads_per_block | num_warps × 32 | 8×32=256 | {2,4,8}×32 |
+| N/A | num_stages | 1 | {1, 2, 3} |
+
+### 2.2 triton_flash_attention.py (Decode Attention — 影响 Output TPS 83%)
+
+CCCL reduce 的 %RANGE%:
+```
+ipt 7:24:1  → BLOCK_M ∈ {16, 32, 64, 128, 256}
+tpb 128:1024:32 → num_warps ∈ {4, 8}
+ipv 1:2:1   → PRE_LOAD_V ∈ {True, False}
+```
+
+| 当前 Triton config | CCCL 对应 | BI-V100 评估 |
+|-------------------|----------|-------------|
+| BLOCK_M=256, BLOCK_N=64, warps=8 | ipt=高, tpb=高 | ⚠️ SMEM 可能不够 |
+| BLOCK_M=128, BLOCK_N=128, warps=4 | ipt=中, tpb=低 | ✓ 可能最优 |
+| BLOCK_M=128, BLOCK_N=64, warps=4 | ipt=中, tpb=低 | ✓ 安全 |
+| BLOCK_M=64, BLOCK_N=64, warps=8 | ipt=低, tpb=高 | ✓ 安全 |
+| BLOCK_M=32, BLOCK_N=32, warps=8 | ipt=极低, tpb=高 | ✓ 保守 |
+| BLOCK_M=16, BLOCK_N=16, warps=4 | ipt=极低, tpb=低 | ✓ 最保守 |
+
+### 2.3 fused_moe.py (MoE GEMM — Qwen3.6 的核心瓶颈)
+
+CCCL 没有直接的 MoE tuning，但 transform 和 reduce 的参数搜索逻辑适用:
+
+| Triton 参数 | 当前值 (batch≤8) | 搜索范围 | CCCL 类比 |
+|-----------|----------------|---------|---------|
+| BLOCK_SIZE_M | 32 | {16, 32, 64} | threads_per_block 的 M 维度 |
+| BLOCK_SIZE_N | 64 | {32, 64, 128} | items 的 N 维度 |
+| BLOCK_SIZE_K | 32 | {32, 64, 128} | vec_size 的 K 维度 |
+| GROUP_SIZE_M | 8 | {1, 4, 8} | CTA swizzle pattern |
+
+---
+
+## 三、执行计划: 从 CCCL benchmark runner 到 Triton autotune
+
+### Step 1: 在 Phanthy Cloud 确认硬件参数 (阻塞一切)
+```python
+import torch
+props = torch.cuda.get_device_properties(0)
+print(f"SMEM: {props.max_shared_memory_per_block}")  # 32KB? 48KB?
+print(f"SMs: {props.multi_processor_count}")          # 16? 50?
+print(f"Warp size: {props.warp_size}")                 # 32?
+```
+
+### Step 2: prefix_prefill BLOCK/NUM_WARPS 网格搜索
+```python
+# 等价于 CCCL: %RANGE% TUNE_ITEMS ipt 32:128:32 × %RANGE% TUNE_THREADS tpb 64:256:32
+for BLOCK in [32, 64, 128]:
+    for NUM_WARPS in [2, 4, 8]:
+        if BLOCK * 128 * 2 <= SMEM_LIMIT:  # SMEM check (CCCL scale_mem_bound 等价)
+            measure_input_tps(BLOCK, NUM_WARPS)
+```
+
+### Step 3: fused_moe BLOCK_SIZE 网格搜索
+```python
+for M in [16, 32, 64]:
+    for N in [32, 64, 128]:
+        for K in [32, 64, 128]:
+            if M * K * 2 + K * N * 2 <= SMEM_LIMIT:  # A_tile + B_tile
+                measure_moe_latency(M, N, K)
+```
+
+### Step 4: triton_flash_attention 过滤不安全 configs
+```python
+# 从 CCCL 的 scale_mem_bound 逻辑: tile_bytes = BLOCK_M * head_dim * 2 (fp16)
+safe_configs = [c for c in autotune_configs 
+                if c.BLOCK_M * 128 * 2 <= SMEM_LIMIT]  # head_dim=128 for Qwen3.6
+# 添加 BI-V100 特化 config
+safe_configs.append(triton.Config(
+    {'BLOCK_M': 64, 'BLOCK_N': 32, 'waves_per_eu': 2, 'PRE_LOAD_V': False},
+    num_stages=1, num_warps=4
+))
+```
+
+---
+
+## 四、为什么这比其他参赛者的方法强
+
+| 方法 | 其他参赛者 | 我们 |
+|------|----------|------|
+| 参数来源 | 猜 / 从 NVIDIA 博客抄 / 凭经验 | CCCL 27 个 tuning header 的 160+ 条 benchmark 注释 |
+| 搜索策略 | 手动试几个值 | CCCL %RANGE% 笛卡尔积系统搜索 |
+| SMEM 约束 | 运行时 crash 才发现 | CCCL scale_mem_bound 编译期检查 |
+| 硬件适配 | 用 NVIDIA 默认值 | muh 27 个 BI-V100 policy_selector |
+| MoE 调优 | 用 EngineX 默认 config | 从 CCCL partition + select_if 逻辑指导 MoE tile 选择 |
--- a/docs/LAYERWISE_SPLIT_KV_CACHE.md
+++ b/docs/LAYERWISE_SPLIT_KV_CACHE.md
@@ -0,0 +1,101 @@
+# Layerwise Split KV Cache Sharding
+
+**Commit:** 494f293b5629 · **PR:** #2260 · **Upstream:** xLLM  
+**Adaptation:** Iluvatar BI-V100 (PCIe topology)  
+**LOC:** +455 −10 across 19 files
+
+## Problem
+
+For models with heterogeneous layer structures (e.g., DeepSeek-V3 with dense
+attention layers interleaved with MoE layers), the KV cache is uniformly
+sharded across all tensor-parallel ranks.  Each rank stores KV for all layers,
+even though different layers may have vastly different head counts.
+
+On Iluvatar BI-V100 (32 GB HBM per card), this wastes memory on ranks that
+serve layers with fewer KV heads and prevents optimal utilisation of each
+card's HBM.
+
+## Solution
+
+Introduce **layerwise split KV cache sharding**: a new KV cache layout
+strategy where each layer's KV cache can be sharded independently across a
+configurable subset of TP ranks.
+
+### Key Components
+
+| # | Component | Files | Description |
+|---|-----------|-------|-------------|
+| 1 | `LayerwiseSplitLayout` | `core/framework/kv_cache/layerwise_split_layout.h` | Per-layer KV shard mappings.  Dense layers spread across all TP ranks; MoE layers concentrate on fewer ranks. |
+| 2 | Layerwise allocation | `core/framework/kv_cache/kv_cache_layerwise.{h,cpp}` | `allocate_kv_caches_layerwise()` — allocates per-layer shard sizes from layout. Handles the ILU/MLU transposed cache layout `[n_blocks, n_heads, block_size, head_dim]`. |
+| 3 | Memory estimation | `core/framework/kv_cache/kv_cache_estimation_layerwise.{h,cpp}` | Reports peak/average per-rank memory; computes savings vs uniform. |
+| 4 | ILU topology mapping | `core/framework/parallel_state/mapping_ilu.{h,cpp}` | PCIe-aware assignment: MoE layers placed on ranks sharing a PCIe switch to maximise intra-group bandwidth. |
+| 5 | Engine integration | `core/distributed_runtime/layerwise_split_{engine_ext,master}.{h,cpp}` | Master computes layout at startup; engines propagate to workers. |
+| 6 | Worker init | `core/runtime/worker_layerwise_init.{h,cpp}` | Workers receive and apply per-layer KV shard assignments. |
+| 7 | Config flag | `core/config/parallel_config_layerwise.{h,cpp}` | `--enable_layerwise_split` gflag (default: false). |
+
+### Iluvatar BI-V100 Hardware Context (verified via ixsmi + debug_warpsize.py)
+
+- **4× BI-V100**, Bus-Id `4B:00.0` – `4E:00.0`, NUMA node 1
+- **Warp size: 64** (NOT 32 — verified via CUDA kernel `warpSize` builtin)
+- **32768 MiB HBM** per card, 1500 MHz SM clock, 1200 MHz mem clock
+- **Flat PIX topology** — all pairs connected via single PCIe bridge (equal BW)
+- IX-ML 3.2.3, Driver 3.2.1, CUDA 10.2 (CoreX)
+- CoreX SDK at `/usr/local/corex/`
+- NCCL for collective communication (same process group as CUDA)
+- KV cache tensor layout (ILU): `[n_blocks, n_heads, block_size, head_dim]` (axis 1 = heads)
+- All verified constants centralized in `core/config/ilu_hw_constants.h`
+
+### Usage
+
+```bash
+# Enable layerwise split KV cache
+./xllm_server --model deepseek-v3 --enable_layerwise_split=true
+
+# Disable (default — uniform sharding, no regression)
+./xllm_server --model deepseek-v3 --enable_layerwise_split=false
+```
+
+## Test Plan
+
+| ID | Level | Description | Criteria |
+|----|-------|-------------|----------|
+| TC-01 | L1 | Allocation correctness | Per-rank allocation matches layout; unassigned layers get zero KV; total equals sum |
+| TC-02 | L1 | Memory estimation accuracy | Layerwise peak ≤ uniform; estimation error ≤ 5% |
+| TC-03 | L1 | ILU PCIe topology mapping | All layers assigned; MoE layers on same-switch ranks; no oversubscription |
+| TC-04 | L1 | Engine layout propagation | All 8 workers receive consistent layout; full layer coverage |
+| TC-05 | L1 | Worker KV shard application | KV populated for assigned layers; zero for unassigned; ASAN clean |
+| TC-06 | L2 | Fallback when disabled | Uniform allocation identical to pre-feature behaviour |
+| TC-07 | L2 | Speculative engine | No crash; KV correctly partitioned per model |
+
+## File Summary
+
+```
+core/
+├── CMakeLists.txt
+├── config/
+│   ├── ilu_hw_constants.h
+│   ├── parallel_config_layerwise.cpp
+│   └── parallel_config_layerwise.h
+├── distributed_runtime/
+│   ├── layerwise_split_engine_ext.cpp
+│   ├── layerwise_split_engine_ext.h
+│   ├── layerwise_split_master.cpp
+│   └── layerwise_split_master.h
+├── framework/
+│   ├── kv_cache/
+│   │   ├── kv_cache_estimation_layerwise.cpp
+│   │   ├── kv_cache_estimation_layerwise.h
+│   │   ├── kv_cache_layerwise.cpp
+│   │   ├── kv_cache_layerwise.h
+│   │   └── layerwise_split_layout.h
+│   └── parallel_state/
+│       ├── mapping_ilu.cpp
+│       └── mapping_ilu.h
+└── runtime/
+    ├── worker_layerwise_init.cpp
+    └── worker_layerwise_init.h
+docs/
+└── LAYERWISE_SPLIT_KV_CACHE.md
+tests/core/
+└── test_layerwise_split_kv_cache.cpp
+```
--- a/docs/MOE_EXECUTION_ANALYSIS.md
+++ b/docs/MOE_EXECUTION_ANALYSIS.md
@@ -0,0 +1,60 @@
+# MoE Execution Path Analysis
+
+> Source: vllm/model_executor/layers/fused_moe/fused_moe.py + vllm/_custom_ops.py
+> Read: 2026-08-04
+
+## The Real Bottleneck
+
+Qwen3.6-35B-A3B has 64 MoE layers, each with:
+- 256 experts, top-8 routing
+- Gate up projection (w1): hidden_dim → intermediate_dim
+- SiLU activation
+- Down projection (w2): intermediate_dim → hidden_dim
+- Weighted sum of 8 expert outputs
+
+### Per-decode-step kernel launches
+
+| Operation | Count | Implementation |
+|---|---|---|
+| fused_moe_kernel (w1) | 64 | ixf_F.vllm_invoke_fused_moe_kernel |
+| silu_and_mul | 64 | ixf_F.silu_and_mul |
+| fused_moe_kernel (w2) | 64 | ixf_F.vllm_invoke_fused_moe_kernel |
+| topk_softmax | 64 | ixf_F.vllm_moe_topk_softmax |
+| moe_align_block_size | 64 | ixf_F.vllm_moe_align_block_size |
+| torch.sum (expert merge) | 64 | PyTorch |
+| paged_attention_v1 | 1 | ixf_F.vllm_single_query_cached_kv_attention |
+| rms_norm | 128 | ixf_F.rms_norm |
+| fused_add_rms_norm | 64 | ixf_F.fused_add_rms_norm |
+| rotary_embedding | 64 | ixf_F.vllm_rotary_embedding_neox |
+| **Total** | **~640+** | |
+
+640+ kernel launches per decode step. At target Output TPS ≥ 395,
+that's 395 × 640 = 253,000 kernel launches per second.
+
+### Memory allocation per step
+
+```python
+# Inside fused_experts, called 64 times per step:
+intermediate_cache1 = torch.empty((M, topk, N))        # 64 × alloc
+intermediate_cache2 = torch.empty((M * topk, N // 2))   # 64 × alloc
+intermediate_cache3 = torch.empty((M, topk, w2_shape[1]))  # 64 × alloc
+```
+
+192 torch.empty calls per decode step = 192 CUDA mallocs.
+At 395 TPS = 75,840 mallocs/second.
+
+### What we can actually change
+
+1. **BLOCK_SIZE_M** (passed to ixformer): 16 for decode (numel=8, M=1×topk=8)
+   - Already optimized: 16 for ≤16 tokens, 32 for ≤64, 64 for ≤1024
+   - ixformer may or may not respect N/K/GROUP values
+
+2. **Intermediate cache pre-allocation**: move torch.empty outside the layer loop
+   - Allocate once, reuse across 64 layers
+   - Saves 192 CUDA mallocs per decode step
+
+3. **torch.sum → ixformer?**: the expert merge `torch.sum(dim=1)` is PyTorch,
+   could potentially be fused into the second fused_moe_kernel call
+
+4. **Chunk size**: VLLM_FUSED_MOE_CHUNK_SIZE controls batching.
+   For decode M=1, chunking adds overhead for no benefit.
--- a/docs/PORTING_ASSESSMENT.md
+++ b/docs/PORTING_ASSESSMENT.md
@@ -0,0 +1,269 @@
+# BI-V100 移植评估：全仓库编译目标清单
+
+## 架构差异
+
+| | NVIDIA V100 | Iluvatar BI-V100 |
+|---|---|---|
+| 架构标识 | `sm_70` | `ivcore10` |
+| 编译器 | `nvcc` / `clang --cuda-gpu-arch=sm_70` | `corex clang/16 --cuda-gpu-arch=ivcore10` |
+| 运行时编译 | `nvrtc` + `nvjitlink` | **不支持** |
+| Driver API | `cuLibraryLoadData` / `cuLibraryGetKernel` | **不支持** |
+| Tensor Core | HMMA (SM70) | **不支持** |
+| Warp size | 32 | 32 (确认) |
+| SMEM | 96KB (configurable) | 48KB |
+| L2 Cache | 6MB | 不同 |
+| SMs | 80 | 16 |
+| CUB block-level | ✅ header-only | ✅ 可通过 corex clang 编译 |
+| CUB device-level | ✅ via nvrtc JIT | ❌ 需要 AOT 替代方案 |
+
+## 1. NVIDIA/CCCL (10,083 files)
+
+### 1.1 c/parallel SHARED LIBRARY — cccl.c.parallel.so
+
+**状态: ❌ 不能直接移植**
+
+12 个算法全部依赖 NVRTC JIT 编译。每个 .cu 通过 `nvrtc_translation_unit` 生成源码，`-arch=sm_XX` 编译，`cuLibraryLoadData` 加载。
+
+| 算法 | 源文件 | 行数 | NVRTC 依赖 | 移植方案 |
+|---|---|---|---|---|
+| reduce | reduce.cu | 783 | nvrtc × 30 | AOT: 直接调用 cub::DeviceReduce with corex |
+| scan | scan.cu | 943 | nvrtc × 25 | AOT: cub::DeviceScan |
+| radix_sort | radix_sort.cu | 947 | nvrtc × 24 | AOT: cub::DeviceRadixSort |
+| merge_sort | merge_sort.cu | 763 | nvrtc × 25 | AOT: cub::DeviceMergeSort |
+| transform | transform.cu | 1014 | nvrtc × 38 | AOT: cub::DeviceTransform |
+| select_if | three_way_partition.cu | 697 | nvrtc × 29 | AOT: cub::DeviceSelect |
+| histogram | histogram.cu | 858 | nvrtc × 18 | AOT: cub::DeviceHistogram |
+| segmented_reduce | segmented_reduce.cu | 655 | nvrtc × 26 | AOT: cub::DeviceSegmentedReduce |
+| segmented_sort | segmented_sort.cu | 1306 | nvrtc × 40 | AOT: cub::DeviceSegmentedSort |
+| binary_search | binary_search.cu | 547 | nvrtc × 8 | AOT: cub::DeviceBinarySearch |
+| unique_by_key | unique_by_key.cu | 768 | nvrtc × 19 | AOT: cub::DeviceUniqueByKey |
+| for | for.cu | 426 | nvrtc × 15 | AOT: cub::DeviceFor |
+
+**移植策略**: 不搬 c/parallel，而是直接用 CUB header-only API 写 AOT .cu 文件，用 corex clang 编译成 .so。每个算法 = 一组固定类型特化。
+
+### 1.2 c/parallel.v2 SHARED LIBRARY
+
+**状态: ❌ 不能直接移植 (依赖 hostjit/libnvcc)**
+
+v2 用嵌入式 clang 做 JIT，不用 nvrtc。理论上可以用 corex clang 替换 libnvcc 的 clang，但改造量大。
+
+### 1.3 CUB block/warp/thread 原语 (header-only)
+
+**状态: ✅ 可直接使用**
+
+| 类别 | 文件数 | 说明 |
+|---|---|---|
+| block primitives | 25 .cuh | BlockReduce, BlockScan, BlockSort, BlockLoad, BlockStore 等 |
+| warp primitives | 17 .cuh | WarpReduce, WarpScan, WarpSort 等 |
+| thread primitives | 8 .cuh | ThreadReduce, ThreadScan, ThreadSort 等 |
+| agent implementations | 26 .cuh | 每个 device algorithm 的 kernel 实现 |
+| dispatch kernels | 17 .cuh | kernel launch 模板 |
+| tuning policies | 27 .cuh | SM-specific 参数选择 (需适配 ivcore10) |
+
+**移植策略**: `#include <cub/block/block_reduce.cuh>` 直接在 corex .cu 中使用。tuning policy 需要为 ivcore10 写新的参数表。
+
+### 1.4 CUB/Thrust benchmarks + examples
+
+| 类别 | 数量 | 移植状态 |
+|---|---|---|
+| CUB benchmarks | 82 | 需适配 ivcore10 编译 |
+| CUB examples | 18 | 需适配 ivcore10 编译 |
+| Thrust examples | 60 | 需适配 ivcore10 编译 |
+| Thrust benchmarks | 75 | 需适配 ivcore10 编译 |
+| cudax examples | 68 | 依赖 cudax runtime，暂不移植 |
+| libcudacxx benchmarks | 62 | 需适配 ivcore10 编译 |
+
+---
+
+## 2. NVIDIA/CUTLASS (7,787 files)
+
+### 2.1 核心 GEMM 库 (header-only)
+
+**状态: ⚠️ 部分可移植**
+
+| SM 架构 | 文件数 | BI-V100 兼容 |
+|---|---|---|
+| SM70 (Volta SIMT) | ~20 | ✅ 需验证 ivcore10 兼容性 |
+| SM75 (Turing) | ~30 | ⚠️ 部分 (SIMT mode) |
+| SM80 (Ampere Tensor) | ~200 | ❌ 需要 HMMA |
+| SM90 (Hopper) | ~300 | ❌ |
+| SM100/120 (Blackwell) | ~200 | ❌ |
+
+### 2.2 Grouped GEMM (MoE 核心)
+
+| Example | 文件 | SM 要求 | 移植状态 |
+|---|---|---|---|
+| 24_gemm_grouped | gemm_grouped.cu | SM70+ SIMT | ✅ 可移植 |
+| 57_hopper_grouped_gemm | — | SM90 | ❌ |
+| 64_ada_fp8_gemm_grouped | — | SM89 | ❌ |
+| 92_blackwell_moe_gemm | — | SM100 | ❌ |
+
+**移植策略**: example 24 (SIMT grouped GEMM) 是唯一能在 BI-V100 跑的。搬过来，接口适配到 xllm group_gemm。
+
+### 2.3 编译目标汇总
+
+| 类别 | 数量 |
+|---|---|
+| Example executables | 164 .cu |
+| Test executables | 862 .cu |
+| Include headers | 785 |
+| SM70 兼容子集 | ~20 examples + ~50 tests |
+
+---
+
+## 3. Dao-AILab/flash-attention (606 .cu files)
+
+### 3.1 flash_attn_2_cuda.so
+
+**状态: ❌ 不能直接移植 (SM80+ Tensor Core)**
+
+所有 kernel 使用 `cute::MMA_Atom<SM80_16x8x16_F16F16F16F16_TN>` — 依赖 Ampere Tensor Core。
+
+| Kernel 类别 | .cu 数量 | SM 要求 |
+|---|---|---|
+| SM80 fwd | 48 | ❌ Tensor Core |
+| SM80 bwd | 24 | ❌ Tensor Core |
+| SM80 fwd_split | 48 | ❌ Tensor Core |
+| SM80 fwd_split_align | 42 | ❌ Tensor Core |
+| Hopper (SM90+) | 453 | ❌ |
+
+### 3.2 可用的算法模板
+
+| 文件 | 行数 | 价值 |
+|---|---|---|
+| flash_fwd_kernel.h | 1301 | attention 算法流程 (Q×K softmax V) |
+| softmax.h | 189 | online softmax 实现 |
+| kernel_traits.h | 344 | SMEM/register 分配策略 |
+| mask.h | 214 | causal mask 实现 |
+| rotary.h | 153 | RoPE in-kernel 实现 |
+
+**移植策略**: 不搬 .cu kernel（依赖 Tensor Core），搬算法模板头文件，基于 CUB block primitives 重写 SIMT attention kernel for ivcore10。或者直接用 ixformer base image 的 `ixinfer_flash_attn_unpad_with_block_tables`（已编译好）。
+
+### 3.3 Layer Norm kernels
+
+| 类别 | .cu 数量 | SM 要求 |
+|---|---|---|
+| ln_fwd | 14 (256~8192 width) | ✅ 纯 SIMT |
+| ln_bwd | 14 | ✅ 纯 SIMT |
+| ln_parallel_fwd | 14 | ✅ 纯 SIMT |
+| ln_parallel_bwd | 14 | ✅ 纯 SIMT |
+
+**移植策略**: Layer norm kernel 是纯 SIMT，不依赖 Tensor Core。可直接用 corex clang 编译。hidden_size=5120 对应 ln_fwd_5120.cu。
+
+---
+
+## 4. jd-opensource/xllm (全平台推理引擎)
+
+### 4.1 ILU (BI-V100) 专用代码
+
+**状态: ✅ 已在项目中 (upstream_ref + ex_engine)**
+
+| 文件 | 行数 | 作用 | 状态 |
+|---|---|---|---|
+| ilu/activation.cpp | 32 | silu_and_mul → ixformer::infer | ✅ 已搬 |
+| ilu/norm.cpp | 50 | rms_norm → ixformer::infer | ✅ 已搬 |
+| ilu/rope.cpp | 31 | rotary_embedding → ixformer::infer | ✅ 已搬 |
+| ilu/attention.cpp | 162 | prefill + decode → ixformer::infer | ✅ 已搬 |
+| ilu/fused_moe.cpp | 99 | topk + expand + combine → ixformer::infer | ✅ 已搬 |
+| ilu/group_gemm.cpp | 39 | group_gemm → ixformer::infer | ✅ 已搬 |
+| ilu/matmul.cpp | 73 | linear → ixformer::infer | ✅ 已搬 |
+| ilu/ixformer.h | 147 | 完整 ixformer::infer API 声明 | ✅ 已搬 |
+| ilu/ilu_ops_api.h | 153 | xllm kernel 层 API | ✅ 已搬 |
+| ilu/utils.h | 62 | 工具函数 | ✅ 已搬 |
+| layers/ilu/fused_moe.cpp | 806 | 完整 MoE 7步 pipeline | ✅ 已搬 |
+| layers/ilu/attention.cpp | 189 | attention layer 封装 | ✅ 已搬 |
+
+### 4.2 CUDA kernels (SM-agnostic)
+
+| 文件 | 行数 | SM 限制 | 状态 |
+|---|---|---|---|
+| activation.cu | 188 | 无 | ✅ 已搬 |
+| norm.cu | 600 | 需 cub::BlockReduce | ✅ 已搬 |
+| rope.cu | 258 | 无 | ✅ 已搬 |
+| block_copy.cu | 209 | 无 | ✅ 已搬 |
+| reshape_paged_cache.cu | 101 | 无 | ✅ 已搬 |
+| moe/moe_topk_softmax_kernels.cuh | 867 | 无 | ✅ 已搬 |
+| moe/moe_compute_index.cu | 155 | 无 | ✅ 已搬 |
+| moe/moe_combine.cu | 105 | 无 | ✅ 已搬 |
+| moe/moe_fused_topk.cu | 59 | 无 | ✅ 已搬 |
+
+### 4.3 CUDA kernels (SM80+ only)
+
+| 文件 | 行数 | SM 限制 | 移植方案 |
+|---|---|---|---|
+| fused_qknorm_rope.cu | 473 | SM80 (`__CUDA_ARCH__ >= 800`) | 拆出 SIMT 部分 |
+| fp8_quant_utils.cuh | 239 | SM89 (`__CUDA_ARCH__ >= 890`) | 不适用 |
+| cutlass_w8a8/*.cu | ~400 | SM90/100/120 | 不适用 |
+
+### 4.4 其他平台代码 (参考用)
+
+| 平台 | kernel 文件数 | layer 文件数 | 说明 |
+|---|---|---|---|
+| DCU (AMD ROCm) | 14 | 12 | GDN 完整实现可参考 |
+| MLU (Cambricon) | 21 | 35 | GDN + MoE 最完整 |
+| MUSA (Moore Threads) | 14 | 12 | GDN kernel 最近代 |
+| NPU (Ascend) | 30+ | 30+ | tilelang GDN 可参考 |
+
+---
+
+## 5. fla-org/flash-linear-attention (349 Triton kernels)
+
+### 5.1 GatedDeltaNet 专用 kernels
+
+**状态: ⚠️ 需验证 Triton 在 BI-V100 上是否工作**
+
+| 文件 | @triton.jit | 行数 | 说明 |
+|---|---|---|---|
+| chunk_fwd.py | 2 | 428 | GDN 前向 chunk (核心) |
+| fused_recurrent.py | 2 | 478 | GDN decode (单步) |
+| wy_fast.py | 4 | 351 | WY representation |
+| gate.py | 6 | 344 | gate cumsum |
+
+### 5.2 通用 Triton 算子
+
+| 目录 | kernel 数 | 说明 |
+|---|---|---|
+| common/ | 36 | chunk_h, chunk_o, fused_recurrent (所有 linear attention 共享) |
+| utils/ | 44 | cumsum, softmax, matmul, solve_tril |
+| gated_delta_rule/ | 14 | GDN 专用 |
+| gdn2/ | 12 | GDN v2 (新版) |
+| kda/ | 24 | Key-dependent attention |
+| delta_rule/ | 12 | 原始 delta rule |
+| gla/ | 18 | Gated Linear Attention |
+
+### 5.3 Backend 分发
+
+| Backend | SM 要求 | 说明 |
+|---|---|---|
+| FlashQLA | SM90+ | ❌ 不适用 BI-V100 |
+| Triton (default) | 任意 GPU | ⚠️ 需验证 corex Triton |
+| triton_ascend | Ascend NPU | ❌ 不适用 |
+
+---
+
+## 移植优先级
+
+### P0 — 直接可编译 (corex clang ivcore10)
+
+1. **xllm CUDA kernels** (9 files, 2542 lines) — 已搬，需在真机编译测试
+2. **CUB block/warp headers** — 已在 cccl_upstream/，可直接 #include
+3. **ix_moe_bridge.so + ix_attn_bridge.so** — pybind11 桥接 ixformer::infer
+
+### P1 — 需适配后可编 (改 SM 架构 + tuning 参数)
+
+4. **FlashAttention layer_norm kernels** (56 .cu) — 纯 SIMT，改编译 flag
+5. **CUTLASS SM70 SIMT GEMM** (example 24 grouped_gemm) — MoE group_gemm 替代方案
+6. **CUB tuning policies** (27 .cuh) — 为 ivcore10 写参数表 (SMEM=48KB, SM=16)
+
+### P2 — 需要重写 (算法可用，硬件指令不兼容)
+
+7. **FlashAttention fwd kernel** — 基于算法模板用 CUB BlockReduce 重写 SIMT 版
+8. **CCCL c/parallel AOT 版** — 绕过 NVRTC，直接用 CUB device API + corex 编译
+9. **FLA Triton GDN kernels** — 需验证 Triton on corex 可行性
+
+### P3 — 不移植
+
+10. FlashAttention SM80+ Tensor Core kernels
+11. CUTLASS SM80/90/100/120 kernels
+12. CCCL nvrtc/nvjitlink 依赖代码
+13. xllm fp8/cutlass_w8a8 quantization kernels
--- a/docs/SPECIALIZATION_ANALYSIS.md
+++ b/docs/SPECIALIZATION_ANALYSIS.md
@@ -0,0 +1,112 @@
+# muh vs CCCL SM100: Type Specialization Parity Analysis
+
+Generated: 2026-07-31
+
+## Summary
+
+| Algorithm | CCCL SM100 branches | muh BI-V100 branches | Status |
+|-----------|--------------------:|---------------------:|--------|
+| reduce | 4+2 det = 6 | 4+2 det+1 default = 7 | ✓ PARITY+ |
+| scan (lookback) | 7 | 7 (after 35ef79c5) | ✓ PARITY |
+| scan (lookahead) | 6 | 6 | ✓ PARITY |
+| topk | 1 (dynamic by key_size) | 1 (dynamic by key_size) | ✓ PARITY |
+| transform | 1 (dynamic by elem_size) | 1 (dynamic by elem_size) | ✓ PARITY |
+| batch_memcpy | 1 (uniform) | 1 (uniform) | ✓ PARITY |
+| for | 1 (uniform) | 1 (uniform) | ✓ PARITY |
+
+## Detailed Breakdown
+
+### reduce (tuning_reduce.cuh)
+
+CCCL SM100 specializes by `(accum_type × offset_size)`:
+- `int64 + o4`: ipt=15, tpb=512, ipv=2
+- `int64 + o8`: ipt=15, tpb=512, ipv=1
+- `float32 + o4`: ipt=16, tpb=512, ipv=2
+- `float64 + o4`: ipt=16, tpb=640, ipv=1
+
+muh BI-V100 maps these with SMEM-derived corrections:
+- `bi100_float32_plus_o4`: tpb=512, ipt=16, ipv=2 (direct match)
+- `bi100_float64_plus_o4`: tpb=512, ipt=12, ipv=1 (SM100 ipt=16 → SMEM overflow at 8B, reduced)
+- `bi100_int64_plus_o4`: tpb=384, ipt=16, ipv=2 (SM100 tpb=512 → SMEM overflow, reduced threads)
+- `bi100_int64_plus_o8`: tpb=384, ipt=16, ipv=1 (same, vec=1 for 8B offset)
+- `bi100_det_float32`: tpb=224, ipt=13 (deterministic path, RAKING)
+- `bi100_det_float64`: tpb=128, ipt=11 (deterministic path, RAKING)
+- `bi100_default`: tpb=256, ipt=16, ipv=4 (fallback)
+
+### scan (tuning_scan.cuh)
+
+CCCL SM100 lookback specializes by `(input_value_size × offset_size)`:
+```
+offset=4: 1B→(512,18) 2B→(512,13) 4B→(384,22) 8B→(416,23)
+offset=8: 1B→(384,14) [2B=skip] 4B→(416,19) 8B→(320,22)
+```
+
+muh BI-V100 after commit 35ef79c5:
+```
+offset=4: 1B→(512,18) 2B→(512,13) 4B→(384,22) 8B→(416,14*)
+offset=8: 1B→(384,14) 4B→(416,19) 8B→(320,19*)
+```
+*items reduced to fit 49152B SMEM
+
+All delay parameters halved (ns×0.5, l2w×0.6) to account for
+BI-V100 L2=6MB vs SM100 L2=50MB.
+
+### SMEM Constraint Validation
+
+Every muh bi100_* struct satisfies: `nominal_tile = tpb × ipt × 4 ≤ 49152`
+
+| Struct | tpb | ipt | nominal_tile | Status |
+|--------|----:|----:|-------------:|--------|
+| bi100_lookback_1B_o4 | 512 | 18 | 36864 | ✓ |
+| bi100_lookback_2B_o4 | 512 | 13 | 26624 | ✓ |
+| bi100_lookback_4B_o4 | 384 | 22 | 33792 | ✓ |
+| bi100_lookback_4B_o8 | 416 | 19 | 31616 | ✓ |
+| bi100_lookback_8B_o4 | 416 | 14 | 23296 | ✓ |
+| bi100_lookback_8B_o8 | 320 | 19 | 24320 | ✓ |
+| bi100_lookback_1B_o8 | 384 | 14 | 21504 | ✓ |
+| bi100_float32_plus_o4 | 512 | 16 | 32768 | ✓ |
+| bi100_float64_plus_o4 | 512 | 12 | 24576 | ✓ |
+| bi100_int64_plus_o4 | 384 | 16 | 24576 | ✓ |
+| bi100_int64_plus_o8 | 384 | 16 | 24576 | ✓ |
+
+## Non-Hot-Path Algorithms (20 missing)
+
+These 20 CCCL algorithms have muh/schema/*.yaml but no tuning header.
+They are NOT on the vllm inference hot path for Qwen3.6 decode.
+If any competition test case triggers them, they will use CCCL defaults
+which may cause SMEM overflow on BI-V100 for large types.
+
+Priority to add (by SMEM overflow risk):
+1. `radix_sort` (89KB tuning, 161 type dispatches) — HIGH risk
+2. `reduce_by_key` (72KB, 134 dispatches) — HIGH risk
+3. `select_if` (107KB, 2729 lines) — MEDIUM risk
+4. `scan_by_key` (88KB) — MEDIUM risk
+5. `unique_by_key` (61KB) — LOW risk
+6. Others: LOW risk (small tile sizes, unlikely SMEM overflow)
+
+## SMEM Overflow Detection (from muh/dispatch.py)
+
+Running `python3 muh_kernel_map.py` against all 6 tuning headers
+detected 5 lookahead structs with incorrect SMEM estimates:
+
+| Struct | SMEM calc | Limit | Status |
+|--------|----------:|------:|--------|
+| bi100_lookahead_1B | 162,816 | 49,152 | ✗ OVERFLOW |
+| bi100_lookahead_2B | 97,280 | 49,152 | ✗ OVERFLOW |
+| bi100_lookahead_4B | 80,896 | 49,152 | ✗ OVERFLOW |
+| bi100_lookahead_4B_float | 89,088 | 49,152 | ✗ OVERFLOW |
+| bi100_lookahead_8B | 89,088 | 49,152 | ✗ OVERFLOW |
+
+**Root cause**: Lookahead SMEM ≠ `threads × items × elem_bytes`.
+The lookahead pipeline uses multi-stage buffering where SMEM =
+`(reduce_squad + scan_store_squad) × items × accum_size × stages`.
+The simple tile formula overestimates by including lookahead items
+that live in registers, not SMEM.
+
+**Impact**: These are currently non-functional on BI-V100 anyway
+(lookahead requires SM90+ warpspeed pipeline support). The dispatch
+correctly falls back to lookback algorithm. But the values in the
+structs are misleading — they should either be corrected or removed.
+
+**Action**: Issue #27 (scan benchmark) TC-04 covers this:
+"lookahead 可行性评估 — 测试 ScanAlgorithm::lookahead 是否能在 BI-V100 上编译运行"
--- a/docs/paged_attention_kernel_architecture.md
+++ b/docs/paged_attention_kernel_architecture.md
@@ -0,0 +1,306 @@
+# Paged Attention Kernel Architecture for BI-V100
+
+## Derived from CCCL Algorithm Patterns
+
+This document designs a complete paged attention kernel from first principles,
+using CCCL's algorithm implementations as the algorithmic foundation.
+Every module maps to a proven CCCL pattern.
+
+---
+
+## 1. Problem Definition
+
+Paged attention computes, for each query token in a decode step:
+
+    output[h, d] = softmax(Q[h] · K[t]^T / √d) · V[t]
+
+where K and V are stored in a **paged block table** (non-contiguous physical memory).
+
+**Qwen3.6 parameters:**
+- head_dim (d) = 256
+- num_heads (H) = 24
+- num_kv_heads (kv_h) = 4, GQA ratio = 6
+- seq_len (T) = up to 100,000
+- block_size = 16 tokens per physical block
+- SMEM per block = 48KB
+
+**The challenge:** K/V are scattered across physical blocks.
+A naive implementation does 6,250 random memory accesses for 100K tokens.
+
+---
+
+## 2. Algorithm Decomposition (Three Levels from CCCL)
+
+### Level 1: Warp Reduce (from `warp_reduce_shfl.cuh`)
+
+**CCCL pattern:** `shfl.sync.down` butterfly reduction in log2(32) = 5 steps.
+Each step: `output = reduction_op(input, ShuffleDown(input, 1 << step))`.
+
+**In attention:** Within one warp (32 threads), each thread holds QK^T scores
+for a subset of KV tokens. Warp reduce computes:
+- `max_score = warp_reduce(scores, max_op)` — for softmax numerical stability
+- `exp_sum = warp_reduce(exp(scores - max_score), plus_op)` — softmax denominator
+- `weighted_v = warp_reduce(exp(scores - max_score) * V[t], plus_op)` — numerator
+
+This is a **compound reduction** — the same pattern as CCCL's `summary_statistics.cu`
+where (count, mean, M2) are reduced together with a custom binary op.
+
+**Our compound type:**
+```
+struct attention_partial {
+    float max_score;      // running max of QK^T
+    float exp_sum;        // sum of exp(score - max_score)
+    float weighted_v[D];  // sum of exp(score - max_score) * V
+};
+```
+
+**Binary op** (from `summary_statistics.cu`):
+```
+attention_partial combine(attention_partial a, attention_partial b) {
+    float new_max = max(a.max_score, b.max_score);
+    float scale_a = exp(a.max_score - new_max);
+    float scale_b = exp(b.max_score - new_max);
+    return {
+        new_max,
+        scale_a * a.exp_sum + scale_b * b.exp_sum,
+        scale_a * a.weighted_v + scale_b * b.weighted_v  // element-wise
+    };
+}
+```
+
+This is exactly the online softmax from Flash Attention.
+It's also exactly CCCL's binary reduction op pattern.
+
+### Level 2: Block Reduce (from `block_reduce_warp_reductions.cuh`)
+
+**CCCL pattern:** Each warp produces a `warp_aggregate`. Lane 0 of each warp
+writes it to `SMEM warp_aggregates[warp_id]`. Then thread 0 serially reduces
+across warps:
+```
+for (warp_idx = 1; warp_idx < warps; ++warp_idx)
+    aggregate = reduction_op(aggregate, warp_aggregates[warp_idx]);
+```
+
+**In attention:** One thread block processes one partition of the KV sequence
+(e.g., PARTITION_SIZE = 512 tokens). Multiple warps within the block each handle
+a chunk of these 512 tokens.
+
+- Warp 0: tokens 0..63 (BLOCK_N=64 at a time, or 32 for head_dim=256)
+- Warp 1: tokens 64..127
+- ...
+- Warp W-1: tokens (W-1)*64..511
+
+Each warp produces an `attention_partial`. Block reduce merges them:
+```
+__shared__ attention_partial warp_partials[NUM_WARPS];
+warp_partials[warp_id] = my_warp_result;
+__syncthreads();
+if (threadIdx.x == 0) {
+    attention_partial block_result = warp_partials[0];
+    for (int w = 1; w < NUM_WARPS; w++)
+        block_result = combine(block_result, warp_partials[w]);
+    // Write block_result to global: tmp_output, exp_sums, max_logits
+}
+```
+
+**SMEM layout for attention_partial at head_dim=256:**
+- max_score: 4 bytes
+- exp_sum: 4 bytes  
+- weighted_v[256]: 256 × 4 = 1024 bytes
+- Total per warp: 1032 bytes
+- For 4 warps: 4128 bytes (fits easily in 48KB)
+
+### Level 3: Cross-Partition Coordination (from `agent_scan.cuh` + decoupled lookback)
+
+**CCCL pattern:** `TilePrefixCallbackOp` implements decoupled lookback.
+Each tile block:
+1. Computes its local aggregate
+2. Publishes local aggregate to global `tile_state` (PARTIAL status)
+3. Warp 0 looks back through predecessor tiles:
+   - If predecessor has INCLUSIVE status → directly use its prefix
+   - If predecessor has PARTIAL status → accumulate and keep looking back
+4. Once prefix is resolved, update own status to INCLUSIVE
+
+**In attention (V2):** Each partition block has its `attention_partial`.
+The cross-partition reduction is simpler than scan because attention
+partitions are **commutative** — we don't need prefix sums, just a
+global reduce.
+
+But the coordination pattern is the same:
+1. Each partition block writes its (max_logit, exp_sum, partial_output) to
+   global memory: `tmp_output[seq, head, partition, :]`
+2. A separate reduction kernel (or the last partition block) reads all
+   partitions and does the final combine.
+
+**Simplification over CCCL's lookback:** Since attention partitions are
+independent (no prefix dependency), we don't need the lookback polling loop.
+Each partition can run fully independently. The reduction is a simple
+parallel reduce over `num_partitions` compound values.
+
+For 100K tokens / 512 partition_size = ~200 partitions.
+200 `attention_partial` values × (4 + 4 + 256×4) = 200 × 1032 = ~200KB.
+One block can reduce all 200 in registers + SMEM.
+
+---
+
+## 3. Paged K/V Gather (from `block_load.cuh` + `cache_modified_input_iterator.cuh`)
+
+**CCCL pattern:** `BlockLoadWarpTranspose` loads contiguous global memory
+into a striped register layout that enables coalesced access. Each thread
+loads `ITEMS_PER_THREAD` elements, and the warp transposes them so each
+thread gets its tile of the data.
+
+**In paged attention:** K/V are not contiguous — they're indexed through
+`block_tables[seq, logical_block] → physical_block`.
+- Key cache: `[num_blocks, kv_heads, head_dim/x, block_size, x]`
+  where x = 16/sizeof(dtype) is the packing factor
+- Value cache: `[num_blocks, kv_heads, head_dim, block_size]`
+
+The gather pattern (from `prefix_prefill.py`, which works on BI-V100):
+```
+# For BLOCK_N tokens starting at position start_n:
+token_ids = start_n + tl.arange(0, BLOCK_N)
+logical_blocks = token_ids // block_size
+within_block = token_ids % block_size
+physical_blocks = tl.load(block_tables + seq * stride + logical_blocks * stride)
+
+# K gather: compute 2D offset array [HEAD_DIM, BLOCK_N]
+off_k = (physical_blocks[None, :] * stride_kc_b +
+         kv_head * stride_kc_h +
+         (offs_d[:, None] // x) * stride_kc_dx +
+         within_block[None, :] * stride_kc_bs +
+         (offs_d[:, None] % x) * stride_kc_x)
+k = tl.load(key_cache + off_k, mask=valid_mask)
+```
+
+This is an **indirect gather** — the physical block ID comes from a table lookup.
+CCCL's `CacheModifiedInputIterator` handles the cache hint part, but the
+indirect indexing is our addition.
+
+**Memory access pattern:**
+- block_tables lookup: 1 global read per BLOCK_N tokens (amortized)
+- K gather: BLOCK_N × HEAD_DIM / x global reads (scattered by physical block)
+- V gather: BLOCK_N × HEAD_DIM global reads (similar scatter)
+
+For BLOCK_N=32, HEAD_DIM=256, x=8: 32 × 32 = 1024 reads for K per iteration.
+At 16 bytes per read (128-bit): 16KB per K load.
+V is similar. Total per iteration: ~32KB — fits in L2 (6MB on BI-V100).
+
+---
+
+## 4. GQA (Grouped Query Attention) Handling
+
+**The insight:** 6 query heads share 1 KV head. Loading KV once and
+computing 6 sets of QK^T scores is 6x more compute-efficient than
+loading KV 6 times.
+
+**CCCL analogy:** This is like `BlockReduce` where we have 6 different
+reduction operations on the same input data. CCCL doesn't have this exact
+pattern, but the principle is: share data loads, parallelize computation.
+
+**Implementation:**
+- Each thread block handles one `(seq, kv_head, partition)` triple
+- Within the block, 6 query heads are processed simultaneously
+- Q vectors: 6 × HEAD_DIM = 6 × 256 = 1536 values in registers (per thread
+  this is 1536/32 = 48 registers — feasible)
+- K/V: loaded once for the kv_head, broadcast across all 6 query heads
+- Scores: 6 × BLOCK_N values per iteration
+- Weighted V: 6 × HEAD_DIM per thread's accumulator
+
+This reduces K/V cache reads by 6x (the GQA ratio).
+
+Grid: `(num_seqs, num_kv_heads, num_partitions)` = `(1, 4, 200)` = 800 blocks
+instead of `(1, 24, 200)` = 4800 blocks.
+
+Each block does 6x more compute but reads KV only once.
+
+---
+
+## 5. SMEM Budget
+
+For one block processing BLOCK_N=32 KV tokens across 6 query heads:
+
+| Item | Size | Notes |
+|------|------|-------|
+| K tile [HEAD_DIM, BLOCK_N] | 32×256×2 = 16KB | fp16, loaded from paged cache |
+| V tile [BLOCK_N, HEAD_DIM] | 32×256×2 = 16KB | fp16, loaded from paged cache |
+| Warp partials [4 warps × attention_partial] | 4×(4+4+256×4) = 4.1KB | For block-level reduce |
+| Q vectors [6 × HEAD_DIM] | 6×256×4 = 6KB | In registers ideally, SMEM if spills |
+| **Total** | **42.1KB** | **≤ 48KB ✓** |
+
+Tight but feasible. If Q stays in registers (likely with 4 warps × 32 threads
+= 128 threads, each handling 6×256/128 = 12 Q values), total SMEM is 36.1KB.
+
+---
+
+## 6. Kernel Launch Configuration
+
+**Phase 1: Partitioned Attention**
+- Grid: `(num_seqs, num_kv_heads, num_partitions)`
+- Block: `(NUM_WARPS × 32)` = 128 threads (4 warps)
+- Each block processes:
+  - PARTITION_SIZE = 512 KV tokens
+  - 6 query heads (GQA broadcast)
+  - Produces 6 × (max_logit, exp_sum, partial_output[256])
+
+**Phase 2: Cross-Partition Reduction**
+- Grid: `(num_seqs, num_kv_heads)` 
+- Block: 128 threads
+- Each block reduces ~200 partitions × 6 query heads
+- Uses `combine()` op (same as CCCL `BlockReduce` but with `attention_partial`)
+
+**Phase 1 iterations per block:**
+- PARTITION_SIZE / BLOCK_N = 512 / 32 = 16 iterations
+- Each iteration: load K[32, 256] + V[32, 256], compute 6×32 scores, update 6 accumulators
+
+---
+
+## 7. Implementation Mapping
+
+| Module | CCCL Source | Our Implementation |
+|--------|------------|-------------------|
+| Warp-level QK^T + softmax | `warp_reduce_shfl.cuh` | Triton: `tl.sum()` within warp-sized groups |
+| Block-level partition reduce | `block_reduce_warp_reductions.cuh` | Triton: shared memory + `tl.reduce()` |
+| Cross-partition combine | `agent_scan.cuh` (simplified, no lookback) | Separate reduction kernel |
+| Paged K/V gather | `block_load.cuh` + indirect indexing | `prefix_prefill.py` pattern adapted |
+| Online softmax | `summary_statistics.cu` binary op | `combine(attention_partial, attention_partial)` |
+| GQA broadcast | (no exact CCCL analog) | Multiple Q per KV load |
+
+---
+
+## 8. Why This Design Beats Python V2
+
+Current Python V2 (3 bmm launches + Python overhead):
+- gather all KV → permute → contiguous → bmm → reshape → softmax → bmm → reduce
+- **Python-CUDA boundary crossed 10+ times per decode step**
+- **Full KV tensor materialized in GPU memory** (200MB-2.4GB depending on GQA)
+
+This kernel (2 GPU launches, zero Python-CUDA crossings during compute):
+- Phase 1: single kernel, K/V loaded tile-by-tile from paged cache (never materialized)
+- Phase 2: single kernel, reduces 200 partitions in SMEM
+- **KV cache stays in paged format** — no gather/permute/contiguous overhead
+- **GQA broadcast within kernel** — KV loaded once for 6 heads
+
+Expected improvement over Python V2: **10-100x** (eliminating Python overhead
+and memory allocation dominates at decode batch_size=1).
+
+Expected improvement over no V2 (V1 only for seq ≤ 8192): **enables long-context
+decode** which V1 cannot do due to SMEM overflow at 48KB.
+
+---
+
+## 9. Implementation Priority
+
+1. **Triton implementation** — if Triton works on BI-V100 with BLOCK=32, head_dim=256:
+   Use the `prefix_prefill.py` paged gather pattern, add the compound reduction.
+   This is the fastest path to a working kernel.
+
+2. **Compiled CUDA kernel** — if `/usr/local/corex/` has a compiler (ixcc):
+   Write the kernel in CUDA using the CCCL patterns directly.
+   `warp_reduce_shfl` → `__shfl_down_sync` PTX
+   `block_reduce` → SMEM warp_aggregates pattern
+   Compile with `torch.utils.cpp_extension.load()` at Docker build time.
+
+3. **Python V2** (current) — fallback if neither Triton nor CUDA works:
+   Already written, tested, has GQA broadcast optimization.
+   This is the floor, not the ceiling.
--- a/docs/server_recon/machine_profile.md
+++ b/docs/server_recon/machine_profile.md
@@ -0,0 +1,42 @@
+# Competition Server Profile
+**Captured**: 2026-08-01
+
+## Hardware
+- **GPU**: 4× Iluvatar BI-V100 32GB HBM each (128GB total)
+  - Clock: SM 1500MHz / Mem 1200MHz
+  - Driver: 3.2.1, COREX 10.2
+  - Power: 250W TDP per card
+- **CPU**: Intel Xeon Gold 6530
+- **RAM**: 503GB DDR
+- **Disk**: 3.5TB overlay, 100GB JuiceFS (public-storage)
+
+## Software
+- **OS**: Ubuntu 20.04.6 LTS, kernel 5.15.0-119
+- **COREX**: 3.2.3 at `/usr/local/corex`
+- **torch**: 2.1.0+corex.3.2.3
+- **vllm**: 0.6.3+corex.3.2.3
+- **transformers**: 4.51.3
+
+## Model
+- **Path**: `/root/public-storage/models/Qwen/Qwen3.6-35B-A3B/`
+- **Name**: Qwen3.6-35B-A3B (MoE, 35B total, 3B active)
+- **Note**: 4 cards × 32GB = 128GB total, model fits
+
+## Key Paths
+- `/root/llm-infer/` — benchmark scripts, README
+- `/root/public-storage/models/Qwen/` — model weights
+- `/root/apps/llm-modelzoo/benchmark/vllm/` — benchmark tools
+- `/share/fshare/common/models/` — shared model storage (NFS)
+
+## Benchmark Tools
+- `benchmark_server_v0.5.0.py` — automated server benchmark
+  - Sweeps: max-num-seqs=[128,256] × num-prompts=[1,128] × input=[128,1024] × output=[128,1024]
+- `benchmark_server_v0.5.0.sh` — launches vllm server + benchmark client
+  - Sets `NCCL_FORCESYNC_DISABLE=1`
+  - Auto-cleanup of vllm processes
+- `benchmark_serving_tokens.py` — online serving benchmark client
+
+## Scoring Formula
+`Output TPS × 16.796 + Input TPS × 2.799 + Cache TPS × 0.56`
+- Threshold: ≥ 8000 weighted score
+- Output TPS weight: 83% of total score
--- a/docs/server_recon/qwen36_bootstrap_issue.md
+++ b/docs/server_recon/qwen36_bootstrap_issue.md
@@ -0,0 +1,45 @@
+# Qwen3.6-35B-A3B Bootstrap Issue
+
+## Problem
+vllm 0.6.3+corex.3.2.3 does not recognize `qwen3_5_moe` model type.
+
+```
+ValueError: The checkpoint you are trying to load has model type `qwen3_5_moe` 
+but Transformers does not recognize this architecture.
+```
+
+## Root Cause
+- Model `config.json` specifies `"model_type": "qwen3_5_moe"` and `"architectures": ["Qwen3_5MoeForCausalLM"]`
+- Server transformers version: 4.51.3 (needs ≥ 4.57.1)
+- Server vllm version: 0.6.3+corex.3.2.3
+
+## Model Architecture (from config.json)
+- **Type**: Qwen3_5MoeForCausalLM (MoE with linear attention)
+- **Total params**: ~35B
+- **Active params per token**: ~3B (8 of 256 experts)
+- **Hidden size**: 2048
+- **Layers**: 40 (30 linear_attention + 10 full_attention, every 4th is full)
+- **Experts**: 256 total, 8 per token
+- **Expert intermediate**: 512
+- **Shared expert intermediate**: 512
+- **Head dim**: 256
+- **KV heads**: 2 (GQA ratio 8:1)
+- **Max position**: 262144
+- **Vocab**: 248320
+- **Precision**: bfloat16
+- **Linear attention**: conv kernel dim=4, 16 key heads (dim128), 32 value heads (dim128)
+- **MTP**: 1 hidden layer (multi-token prediction)
+- **Vision**: yes (patch16, depth27, hidden1152)
+
+## Key Architecture Features
+1. **Hybrid attention**: 3 linear_attention + 1 full_attention pattern (30+10=40 layers)
+2. **MoE**: 256 experts, top-8 routing = very sparse
+3. **Linear attention with conv**: NOT standard transformer — uses conv kernel dim=4
+4. **Multi-token prediction (MTP)**: 1 extra hidden layer for speculative prediction
+5. **Multimodal**: has vision encoder (but competition likely tests text only)
+
+## Solution Paths
+1. **EngineX route**: Check if the competition's enginex-vllm package already supports this model
+   - The repo has `enginex-vllm-bi100-qwen36-main.zip` (96MB) — THIS is likely the answer
+2. **Upgrade transformers**: `pip install transformers>=4.57.1` (may break corex compatibility)
+3. **Custom model registration**: Register Qwen3_5MoeForCausalLM in vllm's model registry