From fab2f10a554974998e8a979d7992c02784bfc848 Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Mon, 12 Aug 2019 15:48:35 -0500 Subject: [PATCH] clean up --- ...plicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.hpp | 9 ++++----- .../include/tensor_description/tensor_coordinate.hpp | 2 -- .../threadwise_generic_tensor_slice_copy.hpp | 4 ++-- ...ice_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp | 6 ++++-- 4 files changed, 10 insertions(+), 11 deletions(-) diff --git a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.hpp b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.hpp index 8c172111f3..a8b330458d 100644 --- a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.hpp +++ b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.hpp @@ -49,10 +49,9 @@ template struct GridwiseConvolutionImplicitGemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer { - __device__ void __launch_bounds__(BlockSize, 2) - Run(const Float* const __restrict__ p_in_global, - const Float* const __restrict__ p_wei_global, - Float* const __restrict__ p_out_global) const + __device__ void Run(const Float* const __restrict__ p_in_global, + const Float* const __restrict__ p_wei_global, + Float* const __restrict__ p_out_global) const { // this is a mess // TODO: find more elegent way of specifying (or calculating) performance parameters @@ -268,7 +267,7 @@ struct GridwiseConvolutionImplicitGemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer // c_thread_mtx definition: this is a mess // TODO:: more elegent way of defining c_thread_mtx constexpr auto c_k0k2_n1n2_thread_mtx_desc = make_ConstantMatrixDescriptor_packed( - Number{}, Number{}); + Number{}, Number{}); const auto blockwise_gemm = BlockwiseGemmBlockABlockBThreadCTransANormalBNormalC_v2< BlockSize, diff --git a/composable_kernel/include/tensor_description/tensor_coordinate.hpp b/composable_kernel/include/tensor_description/tensor_coordinate.hpp index 7b76eac8bc..eee8c27502 100644 --- a/composable_kernel/include/tensor_description/tensor_coordinate.hpp +++ b/composable_kernel/include/tensor_description/tensor_coordinate.hpp @@ -234,8 +234,6 @@ struct MergedTensorCoordinate { static_assert(is_same{} && T::GetSize() == nDim, "wrong!"); - index_t normal_offset_diff = 0; - static_for<0, nDim, 1>{}([&](auto idim) { if(step_sizes[idim] != 0) { diff --git a/composable_kernel/include/tensor_operation/threadwise_generic_tensor_slice_copy.hpp b/composable_kernel/include/tensor_operation/threadwise_generic_tensor_slice_copy.hpp index 57dec923f1..859d5fc164 100644 --- a/composable_kernel/include/tensor_operation/threadwise_generic_tensor_slice_copy.hpp +++ b/composable_kernel/include/tensor_operation/threadwise_generic_tensor_slice_copy.hpp @@ -198,7 +198,7 @@ struct ThreadwiseGenericTensorSliceCopy_v1r1 dst_vector_access_dim, dst_access_id[dst_vector_access_dim] * dst_data_per_access); - vector_t vector_data; + vector_t vector_data{}; // pack vector from buffer static_for<0, DstDataPerAccess, 1>{}([&](auto i) { @@ -224,7 +224,7 @@ struct ThreadwiseGenericTensorSliceCopy_v1r1 dst_data_begin_id(dst_vector_access_dim) = dst_access_id[dst_vector_access_dim] * dst_data_per_access; - vector_t vector_data; + vector_t vector_data{}; // pack vector from buffer for(index_t i = 0; i < DstDataPerAccess; ++i) diff --git a/driver/include/device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp b/driver/include/device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp index 3b37d08132..96ec3ff3bc 100644 --- a/driver/include/device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp +++ b/driver/include/device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp @@ -91,8 +91,6 @@ void device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(InDesc, constexpr index_t WeiBlockCopySrcDataPerRead_E = 4; constexpr index_t WeiBlockCopyDstDataPerWrite_K = 1; - - constexpr index_t OutThreadCopyDataPerAccess_W = 1; #elif 1 // each thread hold 64 data constexpr index_t BlockSize = 256; @@ -101,6 +99,8 @@ void device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(InDesc, constexpr index_t KPerBlock = 128; constexpr index_t EPerBlock = 8; + constexpr index_t GemmNRepeat = 2; + constexpr index_t GemmMPerThreadSubC = 4; constexpr index_t GemmNPerThreadSubC = 4; constexpr index_t GemmMLevel0Cluster = 4; @@ -136,6 +136,8 @@ void device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(InDesc, constexpr index_t KPerBlock = 64; constexpr index_t EPerBlock = 8; + constexpr index_t GemmNRepeat = 2; + constexpr index_t GemmMPerThreadSubC = 2; constexpr index_t GemmNPerThreadSubC = 4; constexpr index_t GemmMLevel0Cluster = 4;