Merge pull request #80 from ROCm/revert_trload_dispatch

Revert "Avoid to use qr_ks_vs_tr_load_async pipeline for MTile <= 64 …

Author: qianfengz

xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_infer_dispatch.h

@@ -160,8 +160,7 @@ struct batched_infer_mask_bias_dropout_dispatch {
         }
       } else {
         // use qr_async_trload pipeline if seqlen <= switch_seqlen_threshold
-        // for MTile <= 64, qr_ks_vs_whole_k_prefetch gives better performance
-        if constexpr (MaxK == 128 && MTile > 64) {
+        if constexpr (MaxK == 128) {
           using FmhaTraits = ck_tile::TileFmhaTraits<
               false, // kPadSeqLenQ,
               false, // kPadSeqLenK,

xformers/csrc/attention/hip_fmha/ck_tiled_fmha_grouped_infer_dispatch.h

@@ -160,8 +160,7 @@ struct grouped_infer_mask_bias_dropout_dispatch {
         }
       } else {
         // use qr_async_trload pipeline if seqlen <= switch_seqlen_threshold
-        // for MTile <= 64, qr_ks_vs_whole_k_prefetch gives better performance
-        if constexpr (MaxK == 128 && MTile > 64) {
+        if constexpr (MaxK == 128) {
           using FmhaTraits = ck_tile::TileFmhaTraits<
               false, // kPadSeqLenQ,
               false, // kPadSeqLenK,