dylanyunlong/project_6
1
0
Fork 0
Code Issues Pull Requests Actions Projects Releases Wiki Activity

[docs] CCCL transform architecture — vectorized vs prefetch, bytes_in_flight scope

Browse Source 
From reading cccl_upstream/cub/cub/device/dispatch/tuning/tuning_transform.cuh:

1. BI-V100 can only use prefetch and vectorized algorithms.
   ldgsts (SM80+ cp.async) and ublkcp (SM90+ bulk copy) are NVIDIA-only.

2. bytes_in_flight only affects the PREFETCH path. For vllm's
   contiguous fp16 element-wise ops (RMSNorm/SiLU/RoPE), the
   VECTORIZED path is selected, where items_per_thread is fixed
   at compile time, not derived from bytes_in_flight.

3. CCCL's cc_to_min_bytes_in_flight: B200=64KB, H100=48KB, A100=16KB,
   V100=12KB. Our 64KB matches B200 level (56 GB/s/SM ≈ B200).

4. Bench result alg=1 confirms vectorized path is used on BI-V100.
   The vectorized default {256, 8, 4} matches the benchmark winner.

Source: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_transform.cuh
This commit is contained in:
project_6
2026-08-05 03:36:00 +00:00
parent 5d6f159906
commit c17e517e9e
1 changed files with 57 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

57
docs/CCCL_REDUCE_ARCHITECTURE_NOTES.md
View File

@@ -295,3 +295,60 @@ identical to paged_attention V2's cross-partition reduce:
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)