mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-04-20 06:49:15 +00:00
Remove unsupported use of c++20 concept. (#2719)
Downstream libraries aren't migrated to c++20 yet, so replace a use of c++20 concept with equivalent SFINAE logic. The template checks for both the existence and the truthiness of the static member variable.
This commit is contained in:
@@ -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");
|
||||
|
||||
|
||||
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user