biondizzle fae61d3ef7 Add c10/cuda/CUDAStream.h include for getCurrentCUDAStream

2026-06-04 04:13:40 +00:00

..

__init__.py

Cleanup Step 2: Archive Lineage P code, fix broken imports

2026-06-02 19:27:07 +00:00

_hash_router.py

Router: full kernel stack — hash, topk, activation+topk, dense decode/prefill

2026-05-21 21:54:05 +00:00

activation_topk.cu

Router: full kernel stack — hash, topk, activation+topk, dense decode/prefill

2026-05-21 21:54:05 +00:00

amax_gsa.cu

Fix CUDA kernel compilation: use c10::cuda::getCurrentCUDAStream

2026-06-01 20:49:55 +00:00

append_swa.cu

KV Cache: schema, allocator, pools, manager, append_swa kernel

2026-05-22 00:08:38 +00:00

blackwell_swizzle.cu

Add c10/cuda/CUDAStream.h include for getCurrentCUDAStream

2026-06-04 04:13:40 +00:00

compressor_reduce.cu

Fix compressor: do not add positional bias to KV content

2026-06-03 15:52:00 +00:00

deinterleave_quantize.cu

Restructure: cutedsl/ -> dsv4/ with proper layering

2026-05-21 17:30:44 +00:00

dequant_nvfp4.cu

KV-1/KV-2: Fused compress+NVFP4 quantize kernels + dequant

2026-06-02 09:37:53 +00:00

flush_write.cu

Flush compressor: schema fix, prepare_forward, flush_write kernels, state rotation

2026-05-22 00:25:47 +00:00

fp8_attention_io.cu

B1: mixed FP8/BF16 decode FMHA integration

2026-06-02 22:53:14 +00:00

fused_amax_quantize.cu

CRITICAL FIX: fused_amax_quantize cross-CTA race condition

2026-06-01 21:26:51 +00:00

fused_deinterleave_amax_quantize.cu

P0 COMPLETE: Eliminate ALL .item() CPU-GPU syncs from NVFP4 activation path

2026-06-01 21:05:03 +00:00

fused_mhc_rmsnorm_quantize.cu

P5: Fused mHC pre_block + RMSNorm + NVFP4 quantize kernel

2026-06-02 16:39:42 +00:00

fused_rmsnorm_quantize.cu

P4: Fix fused RMSNorm kernel — match quantize_nvfp4.cu encoding

2026-06-02 16:28:44 +00:00

gather_kv.cu

fix: extern declarations for gather_swa functions in gather_kv.cu

2026-05-30 21:14:15 +00:00

gather_swa.cu

fix: single PYBIND11_MODULE for combined gather .so

2026-05-30 21:13:24 +00:00

hash_router.cu

Router: full kernel stack — hash, topk, activation+topk, dense decode/prefill

2026-05-21 21:54:05 +00:00

indexer_fp8_score_topk.cu

B2 indexer: adopt TMEM warp-to-row mapping fix

2026-06-03 01:42:38 +00:00

indexer_score_topk.cu

FIX: Deadlock in indexer_score_topk kernel — __syncthreads inside strided loop

2026-06-02 18:11:56 +00:00

kv_quantize.cu

Fix FP8 quantize return type (2-tuple not 3)

2026-06-02 10:02:01 +00:00

loader.py

Fix stale lock file in CUDA loader — prevents infinite spin on crash recovery

2026-06-02 21:34:58 +00:00

mhc_sinkhorn.cu

Fix mHC Sinkhorn kernel: remove VLA, remove Python fallback

2026-06-02 10:44:53 +00:00

quantize_nvfp4.cu

fix quantize_nvfp4 kernel: use proven single-thread-per-CTA pattern from deinterleave_quantize.cu

2026-05-25 16:21:44 +00:00

rope_cuda.cu

P3: CUDA RoPE kernel — single launch per call (vs 5-6 PyTorch ops)

2026-06-02 09:05:22 +00:00

sampler.cu

P7: compressor early return + decode buffering (skip GEMMs when n_complete=0); sampler SMEM fix (LK=24 fits 48KB default); topk on float not bf16

2026-06-01 22:29:56 +00:00

sparse_topk_metadata.cu

Restructure: cutedsl/ -> dsv4/ with proper layering

2026-05-21 17:30:44 +00:00

topk_select.cu

Router: full kernel stack — hash, topk, activation+topk, dense decode/prefill

2026-05-21 21:54:05 +00:00