block_fmha_fwd_splitkv_pipeline_qr_ks_vs_default_policy.hpp Source File

block_fmha_fwd_splitkv_pipeline_qr_ks_vs_default_policy.hpp Source File#

Composable Kernel: block_fmha_fwd_splitkv_pipeline_qr_ks_vs_default_policy.hpp Source File

// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
 
#pragma once
 
#include "ck_tile/core.hpp"
#include "ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy.hpp"
 
namespace ck_tile {
 
// This pipeline is qkv all located in LDS
struct BlockFmhaFwdSplitKVPipelineQRKSVSDefaultPolicy
    : BlockFmhaPipelineQXKSVSCustomPolicy</* QLoadOnce = */ true,
                                          /* AsyncCopy = */ false,
                                          /* NumPrefetchK = */ 1,
                                          /* NumPrefetchV = */ 1>
{
    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