This is the complete list of members for gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >, including all inherited members.

bgm_perf_tuning_knob typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
call(sycl::nd_item< 3 > &item, arguments_t *args)	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	inlinestatic
compute_policy_out typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
compute_policy_out_b_trnp_a typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
compute_policy_QKT typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
dtype_acc typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
dtype_bin typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
dtype_bot typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
dtype_sfx typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_128x128 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_128x256 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_128x64 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_128x64_trnp_a typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_128x64_trnp_af typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_16x2048 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_256x64_trnp_a typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_256x64_trnp_af typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_32x1024 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_64x384 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_arguments_64x512 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_mem_layout_a	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
gemm_mem_layout_out_b	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
gemm_mem_layout_QKT_b	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
gemm_mem_layout_trnp_a	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
gemm_mem_space_a	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
gemm_mem_space_b	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
gemm_mem_space_trnp_a	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
gemm_op_128x128_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_128x256_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_128x64_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_128x64_trnp_a_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_128x64_trnp_af_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_16x2048_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_256x64_trnp_a_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_256x64_trnp_af_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_32x1024_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_64x384_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
gemm_op_64x512_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
global_kslicing	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
k_stride	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
matAcc_128x128_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_128x256_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_128x64_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_128x64_trnp_a_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_128x64_trnp_af_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_16x2048_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_256x64_trnp_a_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_256x64_trnp_af_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_32x1024_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_64x384_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matAcc_64x512_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x128_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x128_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x128_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x256_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x256_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x256_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_trnp_a_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_trnp_a_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_trnp_a_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_trnp_af_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_trnp_af_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_128x64_trnp_af_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_16x2048_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_16x2048_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_16x2048_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_256x64_trnp_a_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_256x64_trnp_a_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_256x64_trnp_a_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_256x64_trnp_af_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_256x64_trnp_af_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_256x64_trnp_af_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_32x1024_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_32x1024_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_32x1024_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_64x384_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_64x384_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_64x384_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_64x512_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_64x512_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matC_64x512_tile_desc_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_128x128_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_128x128_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_128x256_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_128x256_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_16x2048_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_16x2048_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_32x1024_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_32x1024_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_64x384_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_64x384_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_64x512_payload_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
matW_64x512_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
mem_desc_a_out typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
mem_desc_a_out_b_trnp_a typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
mem_desc_a_QKT typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
mem_desc_b_out typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
mem_desc_b_out_b_trnp_a typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
mem_desc_b_QKT typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
mem_layout_a	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
mem_layout_c	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
mem_layout_out_b	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
mem_layout_QKT_b	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
mem_layout_trnp_a	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
mem_space_a	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
mem_space_b	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
mem_space_c	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
periodic_sync_interval	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
pre_processing_128x128 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_128x256 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_128x64 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_128x64_af typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_16x2048 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_256x64 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_256x64_af typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_32x1024 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_64x384 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
pre_processing_64x512 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
prefetch_distance	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
sfx_type_size	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
ThreadNum	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >	static
tile_attr_128x128 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
tile_attr_128x256 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
tile_attr_128x64 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
tile_attr_16x2048 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
tile_attr_256x64 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
tile_attr_32x1024 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
tile_attr_64x384 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
tile_attr_64x512 typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >
work_group_t typedef	gpu::xetla::kernel::xetla_mha_attn_reg_bwd_t< dtype_bwd_bin_, dtype_bwd_bot_, dtype_bwd_sfx_, dtype_bwd_acc_, HWThreadNum, Dopt_RandGenflag, Mkin_flag, Max_SeqLen >