From 8f537a4ccfb8719605ce62fb7ed5764930d3264e Mon Sep 17 00:00:00 2001 From: "Po-Yen, Chen" Date: Wed, 6 Jul 2022 17:04:56 +0800 Subject: [PATCH] Only use v2 pipeline in one gridwise GEMM type --- .../gpu/grid/gridwise_gemm_xdlops_v2r3.hpp | 9 +-------- .../gpu/grid/gridwise_gemm_xdlops_v3r2.hpp | 9 +-------- 2 files changed, 2 insertions(+), 16 deletions(-) diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp index 2b34f2655c..abaf45054d 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp @@ -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 - namespace ck { template ; - static constexpr std::size_t GridwiseGemmPipelineVersion = 2; - - using GridwiseGemmPipe = typename std::tuple_element< - GridwiseGemmPipelineVersion, - std::tuple, GridwiseGemmPipeline_v2>>:: - type; + using GridwiseGemmPipe = GridwiseGemmPipeline_v1; __host__ __device__ static constexpr auto GetABlockDescriptor_K0PerBlock_MPerBlock_K1() { diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp index 8c9a715bdc..5350ebd997 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp @@ -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 - namespace ck { template ; - static constexpr std::size_t GridwiseGemmPipelineVersion = 2; - - using GridwiseGemmPipe = typename std::tuple_element< - GridwiseGemmPipelineVersion, - std::tuple, GridwiseGemmPipeline_v2>>:: - type; + using GridwiseGemmPipe = GridwiseGemmPipeline_v1; __host__ __device__ static constexpr auto GetABlockDescriptor_K0PerBlock_MPerBlock_K1() {