Commit c0fb1dbc authored by Jing Zhang's avatar Jing Zhang
Browse files

add macro for pipeline_v2

parent 146972f4
......@@ -129,6 +129,9 @@
#define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING 0
#define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING_MAC_CLUSTERS 1
// experimental feature: gridwise_gemm_pipeline_v2
#define CK_EXPERIMENTAL_GRIDWISE_GEMM_PIPELINE_V2 0
// hack: have underlying assumption that need to be satsified, otherwise it's a bug
// hack for forcing register to keep idx_diff_low_const in SGPR. idx_diff_low_const must be
// thread-invariant, otherwise it's a bug
......
......@@ -137,11 +137,11 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
// FIXME: pass GridwiseGemmPipe as a template arguement into GridwiseGemm
using GridwiseGemmPipe =
#if 1
#if CK_EXPERIMENTAL_GRIDWISE_GEMM_PIPELINE_V2
GridwiseGemmPipeline_v2;
#else
remove_cvref_t<decltype(
GridwiseGemmPipeline_v1_Selector<NumGemmKPrefetchStage, LoopSched>())>;
#else
GridwiseGemmPipeline_v2;
#endif
__host__ __device__ static constexpr auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1()
......
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