Critical fix based on commit 41ecb8c's discovery:
enginex-vllm-bi100-qwen36 has NO .cu source files. All 9 csrc/*.cu
injection targets in VLLM_INJECTION_POINTS are dead — patches generated
by gen_patch.py have zero effect on the running system.
Old (DEAD):
reduce → csrc/attention/attention_kernels.cu (does not exist)
topk → csrc/sampling/sampling_kernels.cu (does not exist)
scan → csrc/attention/paged_attention_v1.cu (does not exist)
... etc
New (REAL):
prefill → prefix_prefill.py BLOCK/NUM_WARPS (Triton JIT tl.constexpr)
flash_attn → triton_flash_attention.py BLOCK_M/BLOCK_N (Triton autotune)
moe → fused_moe.py BLOCK_SIZE_M (only param ixformer reads)
runtime → _custom_ops.py SMEM (48KB fix)
scheduler → computility-run.yaml num-scheduler-steps
Dead targets preserved as comments for documentation.
Also read: cub/device/dispatch/kernels/kernel_scan.cuh
- DeviceScanInitKernel initializes tile_state for lookback
- __launch_bounds__(threads, 1): max 1 CTA/SM for scan (full SMEM)
- Lookahead requires CUDACC >= 12.8 (not available on BI-V100)
Source: cccl_upstream/cub/cub/device/dispatch/kernels/kernel_scan.cuh
ce42a8579d