4 Commits

Include renames

Author	SHA1	Message	Date
ViperEkura	16a55bb474	refactor: reorganize CUDA kernels into per-family directories ... - move attention kernels to csrc/kernels/attention/ and rotary to rotary/ - add shared common/mma.cuh (mma_sync, ldmatrix) and device.cuh (sm checks) - split fp8_mm into three-layer fp8/common.h, gemm.cuh, mm.cu - fix fused FP8 GEMM ldmatrix lane indexing to fix OOB shared reads - update extension ops, loader, and kernel tests	2026-08-22 20:40:31 +08:00
ViperEkura	6ac3b51496	refactor: separate extension ops and backends	2026-08-16 21:15:52 +08:00
ViperEkura	69f35c46e0	fix: quantize amax from raw values, not scaled fp8 values ... - amax for delayed scale was the quantized max (always ~448), so scale collapsed to 1 - this made fp8 gradients diverge (cosine 0.05) and training stall - stop w/x transpose-quantize amax from polluting the grad scale	2026-08-14 14:26:16 +08:00
ViperEkura	0378e62e17	refactor: split fp8 into fp8_ops adapter and fp8 policy module ... - fp8_ops is the only module touching the pybind (kernel interface) - fp8.py keeps scaling state, delayed amax and aten::linear dispatch - remove circular imports between old fp8_ops/fp8_state/fp8_dispatch	2026-08-14 12:25:29 +08:00