🎉FA2/HGEMM SMEM Swizzle
What's Changed
- [FA2] split-q + tiling-qk D=512 performance🎉 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/178
- [FA2] split-q + tiling-qk D=512 performance🎉 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/179
- [FA2] split-q + tiling-qk D=512 performance🎉 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/180
- [Doc] Refactor README.md to improve readability✔️ by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/181
- [Doc] Refactor README.md for better readability✔️ by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/182
- [FA2] flash-attn-mma 3080/L20/4090 bench✔️ by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/183
- [FA2] flash-attn-mma 3080/L20/4090 bench✔️ by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/184
- [FA2] fa2/hgemm manually smem swizzle🎉 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/185
flash_attn_mma_stages_split_q_tiling_qk_swizzle_kernel
void flash_attn_mma_stages_split_q_tiling_qk_swizzle_kernel<512, 16, 8, 16, 8, 1, 8, 1, 1, 16, 1, 64, 2, 0, 0, 8>(__half *, __half *, __half *, __half *, int, int) (8, 48, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 8.9
Section: Command line profiler metrics
------------------------------------------------------------------ ----------- ------------
Metric Name Metric Unit Metric Value
------------------------------------------------------------------ ----------- ------------
sm__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldsm.avg 0
sm__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldsm.max 0
sm__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldsm.min 0
sm__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldsm.sum 0
------------------------------------------------------------------ ----------- ------------Full Changelog: DefTruth/CUDA-Learn-Notes@v2.6.11...v2.6.12
Contributors
DefTruth