ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-06 07:51:52 +00:00
Code Issues Packages Projects Releases Wiki Activity

Only use v2 pipeline in one gridwise GEMM type

Browse Source
This commit is contained in:
Po-Yen, Chen
2022-07-06 17:04:56 +08:00
parent 77bd29a3c4
commit 8f537a4ccf
2 changed files with 2 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp
View File

@@ -15,8 +15,6 @@
#include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include <tuple>
namespace ck {
template <typename GridwiseGemm,
@@ -129,12 +127,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v2r3
using ThisThreadBlock = ThisThreadBlock<BlockSize>;
static constexpr std::size_t GridwiseGemmPipelineVersion = 2;
using GridwiseGemmPipe = typename std::tuple_element<
GridwiseGemmPipelineVersion,
std::tuple<char, GridwiseGemmPipeline_v1<NumGemmKPrefetchStage>, GridwiseGemmPipeline_v2>>::
type;
using GridwiseGemmPipe = GridwiseGemmPipeline_v1<NumGemmKPrefetchStage>;
__host__ __device__ static constexpr auto GetABlockDescriptor_K0PerBlock_MPerBlock_K1()
{

9
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp
View File

@@ -16,8 +16,6 @@
#include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include <tuple>
namespace ck {
template <typename GridwiseGemm,
@@ -143,12 +141,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v3r2
using ThisThreadBlock = ThisThreadBlock<BlockSize>;
static constexpr std::size_t GridwiseGemmPipelineVersion = 2;
using GridwiseGemmPipe = typename std::tuple_element<
GridwiseGemmPipelineVersion,
std::tuple<char, GridwiseGemmPipeline_v1<NumGemmKPrefetchStage>, GridwiseGemmPipeline_v2>>::
type;
using GridwiseGemmPipe = GridwiseGemmPipeline_v1<NumGemmKPrefetchStage>;
__host__ __device__ static constexpr auto GetABlockDescriptor_K0PerBlock_MPerBlock_K1()
{