/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_streamk.hpp Source File#
device_gemm_xdl_streamk.hpp
Go to the documentation of this file.
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: kernel_launch.hpp:14
float launch_and_time_kernel_with_preprocess(const StreamConfig &stream_config, PreProcessFunc preprocess, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
Definition: kernel_launch.hpp:91
Definition: ck.hpp:268
Definition: stream_config.hpp:10
Definition: block_to_ctile_map.hpp:1022
Definition: gridwise_gemm_xdlops_streamk.hpp:140
Definition: gridwise_gemm_xdlops_streamk.hpp:115
__host__ static constexpr __device__ bool CheckValidity(const Argument &karg)
Definition: gridwise_gemm_xdlops_streamk.hpp:315
__host__ static constexpr __device__ index_t GetSharedMemoryNumberOfByte()
Definition: gridwise_gemm_xdlops_streamk.hpp:289
static std::string GetTypeString()
Definition: gridwise_gemm_xdlops_streamk.hpp:1163
FloatAcc_ FloatAcc
Definition: gridwise_gemm_xdlops_streamk.hpp:132
Definition: integral_constant.hpp:20
Definition: device_base.hpp:197
void * p_workspace_
Definition: device_base.hpp:204
Definition: device_base.hpp:208
Definition: device_gemm_streamk.hpp:25
Definition: device_gemm_xdl_streamk.hpp:131
void Print(const Argument_ &karg)
Definition: device_gemm_xdl_streamk.hpp:133
INVOKER_RUN3_IMPL float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition: device_gemm_xdl_streamk.hpp:223
float RunImp(const typename GridwiseGemm::Argument &karg, const StreamConfig &stream_config=StreamConfig{})
Definition: device_gemm_xdl_streamk.hpp:139
Definition: device_gemm_xdl_streamk.hpp:70
static constexpr bool IsValidCompilationParameter()
Definition: device_gemm_xdl_streamk.hpp:263
static constexpr auto I3
Definition: device_gemm_xdl_streamk.hpp:78
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, uint32_t NumSKBlocks=0xffffffff)
Definition: device_gemm_xdl_streamk.hpp:299
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition: device_gemm_xdl_streamk.hpp:294
static auto MakeInvoker()
Definition: device_gemm_xdl_streamk.hpp:368
std::string GetTypeString() const override
Definition: device_gemm_xdl_streamk.hpp:448
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition: device_gemm_xdl_streamk.hpp:442
static constexpr GET_NXDL_PER_WAVE_IMPL auto NXdlPerWave64
Definition: device_gemm_xdl_streamk.hpp:72
typename GridwiseGemm64::Argument Argument
Definition: device_gemm_xdl_streamk.hpp:127
static constexpr auto I2
Definition: device_gemm_xdl_streamk.hpp:77
static bool IsSupportedArgument(const Argument &karg)
Definition: device_gemm_xdl_streamk.hpp:269
static constexpr auto I0
Definition: device_gemm_xdl_streamk.hpp:75
void SetWorkSpacePointer(BaseArgument *pArg, void *p_workspace, const StreamConfig &=StreamConfig{}) const override
Definition: device_gemm_xdl_streamk.hpp:254
static constexpr auto NXdlPerWave32
Definition: device_gemm_xdl_streamk.hpp:73
size_t GetWorkSpaceSize(const BaseArgument *pArg) const override
Definition: device_gemm_xdl_streamk.hpp:230
static constexpr auto I1
Definition: device_gemm_xdl_streamk.hpp:76
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, index_t NumSKBlocks=0) override
Definition: device_gemm_xdl_streamk.hpp:371