ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-18 20:09:25 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge commit 'c71d7ddd7473b1c952f961e29b09f4a61f0a87d5' into develop

Browse Source
This commit is contained in:
assistant-librarian[bot]
2025-08-25 05:12:56 +00:00
parent ac06a4ed66
commit 2b1a426db5
2 changed files with 21 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/ck_tile/ops/fmha/kernel/fmha_bwd_kernel.hpp
View File

@@ -40,7 +40,7 @@ struct FmhaBwdDQDKDVKernel
static constexpr ck_tile::index_t kBlockSize = FmhaPipeline::kBlockSize;
static constexpr ck_tile::index_t kBlockPerCu = FmhaPipeline::kBlockPerCu;
static constexpr bool kUseQrQtrDorPipeline =
ck_tile::fmha_bwd_qr_qtr_dor_pipeline_c<FmhaPipeline>;
ck_tile::fmha_bwd_qr_qtr_dor_pipeline<FmhaPipeline>::value;
static_assert(!kUseQrQtrDorPipeline || !std::is_same_v<QGradEpiloguePipeline_, void>,
"QrQtrDorPipeline needs QGradEpiloguePipeline");

22
include/ck_tile/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_trload_qr_qtr_dor.hpp
View File

@@ -738,6 +738,24 @@ struct BlockFmhaBwdDQDKDVPipelineTrLoadQRQTRDOR
}
};
template <class T>
concept fmha_bwd_qr_qtr_dor_pipeline_c = T::is_qr_qtr_dor_pipeline;
// We don't support C++20 concepts yet, so we use SFINAE check the existence and truthiness
// of is_qr_qtr_dor_pipeline static member instead of using concepts directly.
//
// The template struct's value field is equivalent to the following commented concept definition.
//
// template <class T>
// concept fmha_bwd_qr_qtr_dor_pipeline_c = T::is_qr_qtr_dor_pipeline;
// SFINAE test for existence and truthiness of static member is_qr_qtr_dor_pipeline.
template <typename, typename = void>
struct fmha_bwd_qr_qtr_dor_pipeline : std::false_type
{
};
template <typename T>
struct fmha_bwd_qr_qtr_dor_pipeline<T, std::void_t<decltype(T::is_qr_qtr_dor_pipeline)>>
: std::bool_constant<T::is_qr_qtr_dor_pipeline>
{
};
} // namespace ck_tile