From 162d3dac4b303ebece490b5f6f70055044e63c91 Mon Sep 17 00:00:00 2001 From: danyao12 Date: Thu, 17 Apr 2025 18:41:26 +0800 Subject: [PATCH] remove dq convert --- example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py b/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py index 291d63518e..bc16344a68 100755 --- a/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py +++ b/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py @@ -791,7 +791,7 @@ float fmha_bwd_v3_(const ck_tile::stream_config& s, fmha_bwd_args a) if(s.log_level_ > 0) std::cout << ", " << fmha_bwd_dot_do_o_get_name_() << ", " << FmhaBwdV3Name::bwd_v3_name << std::flush; fmha_bwd_v3_args args; - args.ptr_dq = a.dq_ptr; + args.ptr_dq = a.dq_acc_ptr; args.ptr_dk = a.dk_ptr; args.ptr_dv = a.dv_ptr; args.ptr_q = a.q_ptr; @@ -1167,9 +1167,9 @@ float fmha_bwd<2>(fmha_bwd_traits t, fmha_bwd_args a, const ck_tile::stream_conf (a.batch_stride_q >= a.stride_q) && (a.batch_stride_do >= a.stride_do) && ((a.batch_stride_dk / a.batch_stride_k) == (a.nhead_q / a.nhead_k)) && ((a.batch_stride_dv / a.batch_stride_v) == (a.nhead_q / a.nhead_k))){{ using dot_do_o_trait_ = fmha_bwd_dot_do_o_traits_<128, FmhaBwdFp16, false, false, false>; using dq_dk_dv_v3_traits_ = fmha_bwd_dq_dk_dv_v3_traits_<128, FmhaBwdFp16, false, true, 0, false, false>; - using convert_dq_trait_ = fmha_bwd_convert_dq_traits_<128, FmhaBwdFp16, false, false, false, false>; + // using convert_dq_trait_ = fmha_bwd_convert_dq_traits_<128, FmhaBwdFp16, false, false, false, false>; // const std::string bwd_v3_name = "bwd_v3_fp16_a32"; - r = fmha_bwd_v3_(s, a); + r = fmha_bwd_v3_(s, a); return r; }} else if((a.seqlen_q % 64 == 0) && (a.hdim_q == 128)){{