Skip to content

GitLab

"script/profile_resnet50.sh" did not exist on "b491ebf38480bc0d6cb329ba6825dee610c59097"
Commit 232864b4 authored by Po Yen Chen's avatar Po Yen Chen
Browse files

Complete splitkv async default policy

parent 8f9f4ae5
Show whitespace changes
Inline Side-by-side
Showing with 15 additions and 6 deletions
include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs_async_default_policy.hpp
View file @ 232864b4
...@@ -9,11 +9,20 @@ ...@@ -9,11 +9,20 @@
namespace ck_tile { namespace ck_tile {
// This pipeline is q in register; kv all located in LDS // This pipeline is q in register; kv all located in LDS
using BlockFmhaFwdSplitKVPipelineQRKSVSAsyncDefaultPolicy = struct BlockFmhaFwdSplitKVPipelineQRKSVSAsyncDefaultPolicy
BlockFmhaPipelineQXKSVSCustomPolicy</* QLoadOnce = */ true, : BlockFmhaPipelineQXKSVSCustomPolicy</* QLoadOnce = */ true,
/* AsyncCopyK = */ true, /* AsyncCopyK = */ true,
/* AsyncCopyV = */ false, /* AsyncCopyV = */ false,
/* NumPrefetchK = */ 3, /* NumPrefetchK = */ 3,
/* NumPrefetchV = */ 3>; /* NumPrefetchV = */ 3>
{
template <typename Problem>
CK_TILE_HOST_DEVICE static constexpr auto GetAlignmentOacc()
{
using OaccDataType = remove_cvref_t<typename Problem::OaccDataType>;
return static_cast<index_t>(16 / sizeof(OaccDataType));
}
};
} // namespace ck_tile } // namespace ck_tile
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment