diff --git a/include/ck_tile/ops/fmha/kernel/fmha_bwd_kernel.hpp b/include/ck_tile/ops/fmha/kernel/fmha_bwd_kernel.hpp index 8750c8b377..5e16fc563b 100644 --- a/include/ck_tile/ops/fmha/kernel/fmha_bwd_kernel.hpp +++ b/include/ck_tile/ops/fmha/kernel/fmha_bwd_kernel.hpp @@ -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; + ck_tile::fmha_bwd_qr_qtr_dor_pipeline::value; static_assert(!kUseQrQtrDorPipeline || !std::is_same_v, "QrQtrDorPipeline needs QGradEpiloguePipeline"); diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_trload_qr_qtr_dor.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_trload_qr_qtr_dor.hpp index 3112070271..789cfb3ea4 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_trload_qr_qtr_dor.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_trload_qr_qtr_dor.hpp @@ -738,6 +738,24 @@ struct BlockFmhaBwdDQDKDVPipelineTrLoadQRQTRDOR } }; -template -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 +// 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 +struct fmha_bwd_qr_qtr_dor_pipeline : std::false_type +{ +}; + +template +struct fmha_bwd_qr_qtr_dor_pipeline> + : std::bool_constant +{ +}; + } // namespace ck_tile