device_gemm_xdl_cshuffle_v2.hpp Source File#
device_gemm_xdl_cshuffle_v2.hpp
Go to the documentation of this file.
23// Note: inter-wave loop scheduler is rolled out to c-shuffle version first. Becuase non c-shuffle
float launch_and_time_kernel(const StreamConfig &stream_config, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
Definition host_utility/kernel_launch.hpp:14
Definition convolution_backward_data_specialization.hpp:8
std::string getGemmSpecializationString(const GemmSpecialization &s)
Definition gemm_specialization.hpp:32
GemmSpecialization
Definition gemm_specialization.hpp:11
@ MKPadding
Definition gemm_specialization.hpp:18
@ KPadding
Definition gemm_specialization.hpp:16
@ MNKPadding
Definition gemm_specialization.hpp:20
@ NKPadding
Definition gemm_specialization.hpp:19
Definition convolution_backward_data_specialization.hpp:7
Definition ck.hpp:268
__global__ void kernel_gemm_xdl_cshuffle_v2(typename GridwiseGemm::Argument karg)
Definition gridwise_gemm_xdl_cshuffle_v2.hpp:26
constexpr LoopScheduler make_default_loop_scheduler()
Definition loop_scheduler.hpp:20
Definition ck/stream_config.hpp:10
Definition gridwise_gemm_xdl_cshuffle_v2.hpp:126
Definition device_base.hpp:197
BaseInvoker()=default
Definition device_gemm_xdl_cshuffle_v2.hpp:148
INVOKER_RUN3_IMPL float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition device_gemm_xdl_cshuffle_v2.hpp:188
float RunImp(const typename GridwiseGemm::Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition device_gemm_xdl_cshuffle_v2.hpp:150
Definition device_gemm_xdl_cshuffle_v2.hpp:80
static auto MakeArgument(const ADataType *p_a, const BDataType *p_b, CDataType *p_c, index_t M, index_t N, index_t K, index_t StrideA, index_t StrideB, index_t StrideC, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation)
Definition device_gemm_xdl_cshuffle_v2.hpp:239
GridwiseGemmBase< math::max(NXdlPerWave64, 1)> GridwiseGemm64
Definition device_gemm_xdl_cshuffle_v2.hpp:141
GridwiseGemm_xdl_cshuffle_v2< ALayout, BLayout, CLayout, ADataType, BDataType, GemmAccDataType, CShuffleDataType, CDataType, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation, GemmSpec, InMemoryDataOperationEnum::Set, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1, BK1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave_, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, false, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, false, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CShuffleBlockTransferScalarPerVector_NPerBlock, LoopSched, PipelineVer, ComputeTypeA, ComputeTypeB > GridwiseGemmBase
Definition device_gemm_xdl_cshuffle_v2.hpp:92
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition device_gemm_xdl_cshuffle_v2.hpp:234
typename GridwiseGemm64::Argument Argument
Definition device_gemm_xdl_cshuffle_v2.hpp:144
static GET_NXDL_PER_WAVE_IMPL constexpr auto NXdlPerWave64
Definition device_gemm_xdl_cshuffle_v2.hpp:83
static constexpr auto I0
Definition device_gemm_xdl_cshuffle_v2.hpp:86
std::string GetTypeString() const override
Definition device_gemm_xdl_cshuffle_v2.hpp:289
std::unique_ptr< BaseArgument > MakeArgumentPointer(const void *p_a, const void *p_b, void *p_c, index_t M, index_t N, index_t K, index_t StrideA, index_t StrideB, index_t StrideC, AElementwiseOperation, BElementwiseOperation, CElementwiseOperation) override
Definition device_gemm_xdl_cshuffle_v2.hpp:258
static constexpr auto I1
Definition device_gemm_xdl_cshuffle_v2.hpp:87
static constexpr bool IsValidCompilationParameter()
Definition device_gemm_xdl_cshuffle_v2.hpp:195
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition device_gemm_xdl_cshuffle_v2.hpp:283
static auto MakeInvoker()
Definition device_gemm_xdl_cshuffle_v2.hpp:255
static bool IsSupportedArgument(const Argument &arg)
Definition device_gemm_xdl_cshuffle_v2.hpp:201
DeviceGemm_Xdl_CShuffleV2 DeviceOp
Definition device_gemm_xdl_cshuffle_v2.hpp:81
static constexpr auto I2
Definition device_gemm_xdl_cshuffle_v2.hpp:88
GridwiseGemmBase< NXdlPerWave32 > GridwiseGemm32
Definition device_gemm_xdl_cshuffle_v2.hpp:142
static constexpr auto NXdlPerWave32
Definition device_gemm_xdl_cshuffle_v2.hpp:84
Definition device_gemm.hpp:22