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 fd027ac8b7..9c371215db 100644 --- a/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py +++ b/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py @@ -349,7 +349,7 @@ struct __attribute__((packed)) fmha_bwd_v3_genl_args }}; struct __attribute__((packed)) fmha_bwd_v3_swa_genl_args -{ +{{ void* ptr_dq; void* ptr_dk; void* ptr_dv; @@ -411,7 +411,7 @@ struct __attribute__((packed)) fmha_bwd_v3_swa_genl_args p1 _p24; unsigned int mask_y; p1 _p25; -}; +}}; struct fmha_bwd_v3_traits {{ @@ -547,6 +547,8 @@ template<> struct FmhaBwdV3Buf struct FmhaBwdV3Buf> {{ static constexpr unsigned char * bwd_v3_buf = bwd_hd64_fp16_causal_a16; }}; template<> struct FmhaBwdV3Buf> {{ static constexpr unsigned char * bwd_v3_buf = bwd_hd64_fp16_causal_a32_pssk; }}; +template<> struct FmhaBwdV3Buf> {{ static constexpr unsigned char * bwd_v3_buf = bwd_v3_bf16_swa_a32_rtne_psskddv; }}; + template struct FmhaBwdV3Ts; // ######################################################|HDim| DataType| MaskType|kIsAtomic32|BF16Cvt|kIsSEQPad|kIsHDPad| template<> struct FmhaBwdV3Ts> {{ static constexpr int ts_qo = 16; static constexpr int ts_kv = 192; }}; @@ -598,6 +600,8 @@ template<> struct FmhaBwdV3Ts struct FmhaBwdV3Ts> {{ static constexpr int ts_qo = 32; static constexpr int ts_kv = 192; }}; template<> struct FmhaBwdV3Ts> {{ static constexpr int ts_qo = 32; static constexpr int ts_kv = 192; }}; +template<> struct FmhaBwdV3Ts> {{ static constexpr int ts_qo = 16; static constexpr int ts_kv = 192; }}; + class fmha_bwd_v3_kernel {{ public: