"...composable_kernel_rocm.git" did not exist on "75ab874e02955279cf45c36a9f1209baa1764a09"
Commit 487a05d6 authored by aska-0096's avatar aska-0096
Browse files

refine blockgemm pipeline version as base struct.

parent 22fe522d
...@@ -9,6 +9,7 @@ ...@@ -9,6 +9,7 @@
#include "ck/ck.hpp" #include "ck/ck.hpp"
#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp" #include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp" #include "ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp" #include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp"
......
...@@ -9,13 +9,6 @@ ...@@ -9,13 +9,6 @@
namespace ck { namespace ck {
enum struct BlockGemmPipelineVersion
{
v1, // Naive
v2, // Mem
v3, // Comp
};
template <BlockGemmPipelineVersion BlkGemmPipelineVer, template <BlockGemmPipelineVersion BlkGemmPipelineVer,
BlockGemmPipelineScheduler BlkGemmPipeSche, BlockGemmPipelineScheduler BlkGemmPipeSche,
index_t BlockSize, index_t BlockSize,
......
...@@ -7,12 +7,6 @@ ...@@ -7,12 +7,6 @@
#include "ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v2.hpp" #include "ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_v2.hpp"
namespace ck { namespace ck {
enum struct BlockGemmPipelineVersion
{
v1, // Single lds buffer
v2, // Double lds buffer
};
template <BlockGemmPipelineVersion BlkGemmPipelineVer, template <BlockGemmPipelineVersion BlkGemmPipelineVer,
BlockGemmPipelineScheduler BlkGemmPipeSche, BlockGemmPipelineScheduler BlkGemmPipeSche,
index_t BlockSize, index_t BlockSize,
......
...@@ -11,15 +11,6 @@ ...@@ -11,15 +11,6 @@
namespace ck { namespace ck {
enum struct BlockGemmPipelineVersion
{
v1, // Naive
v2, // Mem
v3, // Comp
v4, // Comp, double lds buffer
v5, // Comp, double global prefetch register buffer
};
template <BlockGemmPipelineVersion BlkGemmPipelineVer, template <BlockGemmPipelineVersion BlkGemmPipelineVer,
BlockGemmPipelineScheduler BlkGemmPipeSche, BlockGemmPipelineScheduler BlkGemmPipeSche,
index_t BlockSize, index_t BlockSize,
......
...@@ -8,6 +8,19 @@ ...@@ -8,6 +8,19 @@
namespace ck { namespace ck {
enum struct BlockGemmPipelineVersion
{
// For GEMM
v1, // Naive
v2, // Mem
v3, // Comp
v4, // Comp, double lds buffer
v5, // Comp, double global prefetch register buffer
// For GEMM with preshuffled weight
// v1, single lds buffer
// v2, double lds buffer
};
enum struct BlockGemmPipelineScheduler enum struct BlockGemmPipelineScheduler
{ {
Intrawave, Intrawave,
......
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