Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[FA2] tiling-qk + fully smem swizzle✔️ #188

Merged
merged 10 commits into from
Dec 27, 2024
Merged

[FA2] tiling-qk + fully smem swizzle✔️ #188

merged 10 commits into from
Dec 27, 2024

Conversation

DefTruth
Copy link
Owner

  • no padding, 0 ldsm bank_conflicts
void flash_attn_mma_stages_split_q_tiling_qk_fully_swizzle_kernel<64, 16, 8, 16, 4, 1, 4, 1, 1, 8, 1, 8, 1, 0, 0, 0>(__half *, __half *, __half *, __half *, int, int) (16, 48, 1)x(128, 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
    ------------------------------------------------------------------ ----------- ------------
  • no padding, 0 ld & st bank_conflicts
  void flash_attn_mma_stages_split_q_tiling_qk_fully_swizzle_kernel<64, 16, 8, 16, 4, 1, 4, 1, 1, 8, 1, 8, 2, 0, 0, 0>(__half *, __half *, __half *, __half *, int, int) (16, 48, 1)x(128, 1, 1), Context 1, Stream 7, Device 0, CC 8.9
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.avg                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.max                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.min                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
    -------------------------------------------------------- ----------- ------------

  void flash_attn_mma_stages_split_q_tiling_qk_fully_swizzle_kernel<64, 16, 8, 16, 4, 1, 4, 1, 1, 8, 1, 8, 2, 0, 0, 0>(__half *, __half *, __half *, __half *, int, int) (16, 48, 1)x(128, 1, 1), Context 1, Stream 7, Device 0, CC 8.9
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.avg                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.max                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.min                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------

@DefTruth DefTruth merged commit 4165873 into main Dec 27, 2024
@DefTruth DefTruth deleted the fa-mma-swizzle branch December 27, 2024 03:55
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

1 participant