Skip to content

GitLab

"...composable_kernel.git" did not exist on "e4584d91acc14a22426cbf081c8cc8394c136f6b"
Commit 9809f5d4 authored by danyao12's avatar danyao12
Browse files

remove hardcode

parent de4494a2
Hide whitespace changes
Inline Side-by-side
Showing with 14 additions and 12 deletions
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_bwd_xdl_cshuffle_qloop_light_v1_protro.hpp
View file @ 9809f5d4
...@@ -357,9 +357,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_Light_V1 ...@@ -357,9 +357,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_Light_V1
static constexpr auto I1 = Number<1>{}; static constexpr auto I1 = Number<1>{};
static constexpr auto I2 = Number<2>{}; static constexpr auto I2 = Number<2>{};
static constexpr index_t V_O1 = 8; static constexpr index_t V_O1 = BK1;
static constexpr index_t Y_O1 = 8; static constexpr index_t Y_O1 = AK1;
static constexpr index_t Y_M1 = 2; static constexpr index_t Y_M1 = AK1;
static constexpr auto padder = GemmGemmPadder<GemmSpec, static constexpr auto padder = GemmGemmPadder<GemmSpec,
Number<MPerBlock>, Number<MPerBlock>,
......
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_bwd_xdl_cshuffle_qloop_light_v2_protro.hpp
View file @ 9809f5d4
...@@ -364,8 +364,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_Light_V2 ...@@ -364,8 +364,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_Light_V2
static constexpr auto I1 = Number<1>{}; static constexpr auto I1 = Number<1>{};
static constexpr auto I2 = Number<2>{}; static constexpr auto I2 = Number<2>{};
static constexpr index_t Y_O1 = 8; static constexpr index_t V_O1 = BK1;
static constexpr index_t Y_M1 = 2; static constexpr index_t Y_O1 = AK1;
static constexpr index_t Y_M1 = AK1;
static constexpr auto padder = GemmGemmPadder<GemmSpec, static constexpr auto padder = GemmGemmPadder<GemmSpec,
Number<MPerBlock>, Number<MPerBlock>,
...@@ -556,7 +557,7 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_Light_V2 ...@@ -556,7 +557,7 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_Light_V2
Sequence<padder.PadN, padder.PadO>{}); Sequence<padder.PadN, padder.PadO>{});
// N_O to O0_N_O1; to refactor // N_O to O0_N_O1; to refactor
return Transform::MakeB0GridDescriptor_BK0_N_BK1(v_grid_desc_n_o, Number<BK1>{}); return Transform::MakeB0GridDescriptor_BK0_N_BK1(v_grid_desc_n_o, Number<V_O1>{});
} }
// Z in Gemm0 C position // Z in Gemm0 C position
......
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_bwd_xdl_cshuffle_qloop_v1_protro.hpp
View file @ 9809f5d4
...@@ -304,9 +304,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_V1 ...@@ -304,9 +304,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_V1
static constexpr auto I1 = Number<1>{}; static constexpr auto I1 = Number<1>{};
static constexpr auto I2 = Number<2>{}; static constexpr auto I2 = Number<2>{};
static constexpr index_t V_O1 = 8; static constexpr index_t V_O1 = BK1;
static constexpr index_t Y_O1 = 8; static constexpr index_t Y_O1 = AK1;
static constexpr index_t Y_M1 = 2; static constexpr index_t Y_M1 = AK1;
static constexpr auto padder = GemmGemmPadder<GemmSpec, static constexpr auto padder = GemmGemmPadder<GemmSpec,
Number<MPerBlock>, Number<MPerBlock>,
......
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_bwd_xdl_cshuffle_qloop_v2_protro.hpp
View file @ 9809f5d4
...@@ -311,8 +311,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_V2 ...@@ -311,8 +311,9 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_V2
static constexpr auto I1 = Number<1>{}; static constexpr auto I1 = Number<1>{};
static constexpr auto I2 = Number<2>{}; static constexpr auto I2 = Number<2>{};
static constexpr index_t Y_O1 = 8; static constexpr index_t V_O1 = BK1;
static constexpr index_t Y_M1 = 2; static constexpr index_t Y_O1 = AK1;
static constexpr index_t Y_M1 = AK1;
static constexpr auto padder = GemmGemmPadder<GemmSpec, static constexpr auto padder = GemmGemmPadder<GemmSpec,
Number<MPerBlock>, Number<MPerBlock>,
...@@ -494,7 +495,7 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_V2 ...@@ -494,7 +495,7 @@ struct DeviceBatchedMultiheadAttentionBackward_Qloop_Xdl_CShuffle_V2
Sequence<padder.PadN, padder.PadO>{}); Sequence<padder.PadN, padder.PadO>{});
// N_O to O0_N_O1; to refactor // N_O to O0_N_O1; to refactor
return Transform::MakeB0GridDescriptor_BK0_N_BK1(v_grid_desc_n_o, Number<BK1>{}); return Transform::MakeB0GridDescriptor_BK0_N_BK1(v_grid_desc_n_o, Number<V_O1>{});
} }
// Z in Gemm0 C position // Z in Gemm0 C position
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment