diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_appendkv_pipeline_default_policy.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_appendkv_pipeline_default_policy.hpp index 76066befeb..2e0592ed14 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_appendkv_pipeline_default_policy.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_appendkv_pipeline_default_policy.hpp @@ -58,7 +58,7 @@ struct BlockFmhaFwdAppendKVPipelineDefaultPolicy } template - CK_TILE_DEVICE static constexpr auto GetQNumElemsPerRead() + CK_TILE_HOST_DEVICE static constexpr auto GetQNumElemsPerRead() { using DataType = typename Problem::QDataType; @@ -121,7 +121,7 @@ struct BlockFmhaFwdAppendKVPipelineDefaultPolicy } template - CK_TILE_DEVICE static constexpr auto GetKnewNumElemsPerRead() + CK_TILE_HOST_DEVICE static constexpr auto GetKnewNumElemsPerRead() { using DataType = typename Problem::KDataType; @@ -190,7 +190,7 @@ struct BlockFmhaFwdAppendKVPipelineDefaultPolicy } template - CK_TILE_DEVICE static constexpr auto MakeVnewDramTileDistribution() + CK_TILE_HOST_DEVICE static constexpr auto MakeVnewDramTileDistribution() { using VLayout = remove_cvref_t; using VDataType = remove_cvref_t; @@ -237,7 +237,7 @@ struct BlockFmhaFwdAppendKVPipelineDefaultPolicy } template - CK_TILE_DEVICE static constexpr auto GetRotaryCosSinTileSize() + CK_TILE_HOST_DEVICE static constexpr auto GetRotaryCosSinTileSize() { constexpr index_t height = (IsRotaryCosSinForQ ? Problem::kM0 : Problem::kN0); @@ -245,14 +245,14 @@ struct BlockFmhaFwdAppendKVPipelineDefaultPolicy { return make_tuple(number{}, number{}); } - else // Problem::RotaryEnum == RotaryEmbeddingEnum::INTERLEAVED + else { return make_tuple(number{}, number{}); } } template - CK_TILE_DEVICE static constexpr auto MakeRotaryCosSinTileDistribution() + CK_TILE_HOST_DEVICE static constexpr auto MakeRotaryCosSinTileDistribution() { using DataType = std::conditional_t