mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-12 01:10:17 +00:00
41cdd3801a
* gemm+activation * move C pointwise operation into threadwise copy * add pointwise operation to A/B matrix * update ckProfiler * adding bias add * adding bias add * adding bias add * added bias add; worked around compiler issues * clean up * clean up * Update README.md * Update README.md * Update README.md * clean up * add conv_xdl example * adding conv_xdl_bias_relu_add example * add conv+bias+relu+add, but has register spill issue * tweak * tweak * refactor * Update README.md update readme for example/2_gemm_xdl_bias_relu_add * clean up * Update README.md update readme for example/3_conv_xdl * Update README.md
43 lines
1.3 KiB
C++
43 lines
1.3 KiB
C++
#ifndef DEVICE_GEMM_HPP
|
|
#define DEVICE_GEMM_HPP
|
|
|
|
#include <iostream>
|
|
#include "device_base.hpp"
|
|
|
|
namespace ck {
|
|
namespace tensor_operation {
|
|
namespace device {
|
|
|
|
template <typename AElementwiseOperation,
|
|
typename BElementwiseOperation,
|
|
typename CElementwiseOperation>
|
|
struct DeviceGemm : public BaseOperator
|
|
{
|
|
virtual std::unique_ptr<BaseArgument>
|
|
MakeArgumentPointer(const void* p_a,
|
|
const void* p_b,
|
|
void* p_c,
|
|
ck::index_t M,
|
|
ck::index_t N,
|
|
ck::index_t K,
|
|
ck::index_t StrideA,
|
|
ck::index_t StrideB,
|
|
ck::index_t StrideC,
|
|
AElementwiseOperation a_element_op,
|
|
BElementwiseOperation b_element_op,
|
|
CElementwiseOperation c_element_op) = 0;
|
|
|
|
virtual std::unique_ptr<BaseInvoker> MakeInvokerPointer() = 0;
|
|
};
|
|
|
|
template <typename AElementwiseOperation,
|
|
typename BElementwiseOperation,
|
|
typename CElementwiseOperation>
|
|
using DeviceGemmPtr = std::unique_ptr<
|
|
DeviceGemm<AElementwiseOperation, BElementwiseOperation, CElementwiseOperation>>;
|
|
|
|
} // namespace device
|
|
} // namespace tensor_operation
|
|
} // namespace ck
|
|
#endif
|