Skip to content

GitLab

"...composable_kernel_rocm.git" did not exist on "987cc54d90d4c89141127d97d3160a4d13674b3f"
Commit b469ec5f authored by Harisankar Sadasivan's avatar Harisankar Sadasivan
Browse files

fixing clang-format issues

parent fbce4790
Hide whitespace changes
Inline Side-by-side
Showing with 196 additions and 77 deletions
example/01_gemm/README.md 100644 → 100755
View file @ b469ec5f
File mode changed from 100644 to 100755
example/01_gemm/common.hpp
View file @ b469ec5f
...@@ -43,8 +43,8 @@ struct ProblemSizeStreamK final ...@@ -43,8 +43,8 @@ struct ProblemSizeStreamK final
ck::index_t StrideB = 4096; ck::index_t StrideB = 4096;
ck::index_t StrideC = 4096; ck::index_t StrideC = 4096;
ck::index_t Grid_size = -1; //defaults to max occupancy ck::index_t Grid_size = -1; // defaults to max occupancy
ck::index_t Streamk_sel = 1; //defaults to 1-tile SK ck::index_t Streamk_sel = 1; // defaults to 1-tile SK
}; };
struct ProblemSizeSplitK final struct ProblemSizeSplitK final
......
example/01_gemm/run_gemm_example_streamk_v2.inc
View file @ b469ec5f
...@@ -94,13 +94,13 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) ...@@ -94,13 +94,13 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
using namespace ck::literals; using namespace ck::literals;
auto M = problem_size.M; auto M = problem_size.M;
auto N = problem_size.N; auto N = problem_size.N;
auto K = problem_size.K; auto K = problem_size.K;
auto StrideA = problem_size.StrideA; auto StrideA = problem_size.StrideA;
auto StrideB = problem_size.StrideB; auto StrideB = problem_size.StrideB;
auto StrideC = problem_size.StrideC; auto StrideC = problem_size.StrideC;
auto Grid_size = problem_size.Grid_size; auto Grid_size = problem_size.Grid_size;
auto Streamk_sel = problem_size.Streamk_sel; auto Streamk_sel = problem_size.Streamk_sel;
auto f_host_tensor_descriptor = auto f_host_tensor_descriptor =
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_default_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_default_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_default_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmDefault>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmDefault>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_kpadding_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_kpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_kpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmKPadding>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_mnkpadding_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_mnkpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_mnkpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmMNKPadding>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmMNKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_mnpadding_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_mnpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_mnpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmMNPadding>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_comp_instances<GemmMNPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_default_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_default_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_default_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Intrawave, GemmDefault>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Intrawave,
GemmDefault>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_kpadding_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_kpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_kpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Intrawave, GemmKPadding>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Intrawave,
GemmKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_mnkpadding_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_mnkpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v1_mnkpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Intrawave, GemmMNKPadding>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Intrawave,
GemmMNKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_default_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_default_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_default_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Interwave, GemmDefault>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Interwave,
GemmDefault>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_kpadding_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_kpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_kpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Interwave, GemmKPadding>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Interwave,
GemmKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_mnkpadding_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_mnkpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_v2_mnkpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Row, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Row,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Interwave, GemmMNKPadding>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_kn_mn_mem_instances<Interwave,
GemmMNKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_default_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_default_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_default_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Col, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Col,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmDefault>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmDefault>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_kpadding_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_kpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_kpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Col, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Col,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmKPadding>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_mnkpadding_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_mnkpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_mnkpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Col, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Col,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmMNKPadding>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmMNKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_mnpadding_instance.cpp
View file @ b469ec5f
...@@ -9,12 +9,19 @@ namespace device { ...@@ -9,12 +9,19 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_mnpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_mnpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Col, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Col,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmMNPadding>{}); instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_comp_instances<GemmMNPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_default_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_default_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_default_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Col, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Col,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_instances<Intrawave, GemmDefault>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_instances<Intrawave,
GemmDefault>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_kpadding_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_kpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_kpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Col, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Col,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_instances<Intrawave, GemmKPadding>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_instances<Intrawave,
GemmKPadding>{});
} }
} // namespace instance } // namespace instance
......
library/src/tensor_operation_instance/gpu/gemm_universal_streamk/device_gemm_xdl_universal_streamk_f16_f16_f16/device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_mnkpadding_instance.cpp
View file @ b469ec5f
...@@ -9,13 +9,20 @@ namespace device { ...@@ -9,13 +9,20 @@ namespace device {
namespace instance { namespace instance {
void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_mnkpadding_instances( void add_device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_v1_mnkpadding_instances(
std::vector<std::unique_ptr< std::vector<std::unique_ptr<DeviceGemm_Streamk_V2<Row,
DeviceGemm_Streamk_V2<Row, Col, Row, F16, F16, F16, PassThrough, PassThrough, PassThrough>>>& Col,
instances) Row,
F16,
F16,
F16,
PassThrough,
PassThrough,
PassThrough>>>& instances)
{ {
add_device_operation_instances( add_device_operation_instances(
instances, instances,
device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_instances<Intrawave, GemmMNKPadding>{}); device_gemm_xdl_universal_streamk_f16_f16_f16_mk_nk_mn_mem_instances<Intrawave,
GemmMNKPadding>{});
} }
} // namespace instance } // namespace instance
......
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