3 Commits

Include renames

Author	SHA1	Message	Date
ViperEkura	d5067af064	refactor: harden param PODs and CUTLASS-style fp8 layout tags ... - NSDMI null/-1 defaults for AttentionParams/FP8Params pointer+flag members: partially packed structs can no longer hold garbage non-null pointers that gate optional paths (root cause class of the paged test bug); still aggregates, still trivially copyable - move per-lane ldmatrix wrappers (ldsm_x2/x4) from fp8/gemm.cuh to common/mma.cuh as ldmatrix_x2_lane/x4_lane, next to the single-address variants - DEVICE_FORCEINLINE macro in common/mma.cuh (matches layout_policies.cuh, internal linkage) - frag_addr now delegates to tile_at: the swizzle math has one source - operand layouts as CUTLASS-style RowMajor/ColMajor tags threaded from launch_fp8_gemm through the kernel to load_operand_tile; B's operand view via transpose_layout_t; call sites read <Fmt, false, RowMajor, ColMajor> instead of <Fmt, false, false, true>	2026-08-24 15:27:52 +08:00
ViperEkura	31ca357c61	refactor: namespace csrc kernels and extract common helpers ... - attention family -> astrai::attention; fp8 family -> astrai::fp8 - new common/reduce.cuh (warp/group reductions, atomic_max_float) - new common/cp_async.cuh (predicated cp_async_16, commit/wait group) - move MAX_SPLITS into attention/common.h; delete warp_utils.cuh - .cu bindings and pure C tests open family namespaces via using	2026-08-24 14:49:55 +08:00
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