From 3276a5e9b94dc9ffccacf469e57b1f8502dabccb Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Fri, 5 Jul 2019 15:33:58 -0500 Subject: [PATCH] update build --- ...mm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp | 20 ++--- .../include/utility/config_amd.hpp.in | 3 +- .../include/utility/config_nvidia.hpp.in | 3 +- driver/include/conv_common.hpp | 14 +++- driver/include/device.hpp | 6 +- ...lution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp | 6 +- ...lution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp | 4 +- driver/include/tensor.hpp | 47 +++++++++++- driver/src/driver.cpp | 75 +++++-------------- 9 files changed, 90 insertions(+), 88 deletions(-) diff --git a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp index bb9b6cbd07..2745d9dd3c 100644 --- a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp +++ b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp @@ -209,15 +209,6 @@ struct GridwiseConvolutionImplicitGemm_v3_nchw_cyxk_nkhw_lds_double_buffer GemmDataPerReadA, GemmDataPerReadB>{}; - // choose GEMM implementation here - const auto run_blockwise_gemm = [&](auto... Xs) { -#if 1 - return blockwise_gemm.Run(Xs...); -#else - return blockwise_gemm.Run_amd_asm(Xs...); -#endif - }; - // LDS allocation for input and weight: be careful of alignment constexpr index_t max_align = math::lcm(InBlockCopyDstDataPerWrite_N2, WeiBlockCopyDataPerAccess_K, @@ -225,9 +216,10 @@ struct GridwiseConvolutionImplicitGemm_v3_nchw_cyxk_nkhw_lds_double_buffer GemmDataPerReadB); constexpr index_t in_block_space = - in_c_n1_b_n2_block_mem_desc.GetElementSpace(Number{}); + math::integer_least_multiple(in_c_n1_b_n2_block_mem_desc.GetElementSpace(), max_align); - constexpr index_t wei_block_space = wei_c_k_block_desc.GetElementSpace(Number{}); + constexpr index_t wei_block_space = + math::integer_least_multiple(wei_c_k_block_desc.GetElementSpace(), max_align); __shared__ Float p_in_block_double[2 * in_block_space]; __shared__ Float p_wei_block_double[2 * wei_block_space]; @@ -291,7 +283,7 @@ struct GridwiseConvolutionImplicitGemm_v3_nchw_cyxk_nkhw_lds_double_buffer p_wei_register_clipboard); // LDS double buffer: GEMM on current data - run_blockwise_gemm(p_wei_block_now, p_in_block_now, p_out_thread); + blockwise_gemm.Run(p_wei_block_now, p_in_block_now, p_out_thread); // LDS double buffer: store next data to LDS blockwise_in_copy.RunStoreRegisterClipboard(p_in_register_clipboard, @@ -319,7 +311,7 @@ struct GridwiseConvolutionImplicitGemm_v3_nchw_cyxk_nkhw_lds_double_buffer p_wei_register_clipboard); // LDS double buffer: GEMM on current data - run_blockwise_gemm(p_wei_block_double, p_in_block_double, p_out_thread); + blockwise_gemm.Run(p_wei_block_double, p_in_block_double, p_out_thread); // LDS double buffer: store next data to LDS blockwise_in_copy.RunStoreRegisterClipboard(p_in_register_clipboard, @@ -331,7 +323,7 @@ struct GridwiseConvolutionImplicitGemm_v3_nchw_cyxk_nkhw_lds_double_buffer __syncthreads(); // LDS double buffer: GEMM on current data - run_blockwise_gemm(p_wei_block_double + wei_block_space, + blockwise_gemm.Run(p_wei_block_double + wei_block_space, p_in_block_double + in_block_space, p_out_thread); } diff --git a/composable_kernel/include/utility/config_amd.hpp.in b/composable_kernel/include/utility/config_amd.hpp.in index aa12140fd7..e11ac1b70e 100644 --- a/composable_kernel/include/utility/config_amd.hpp.in +++ b/composable_kernel/include/utility/config_amd.hpp.in @@ -18,7 +18,8 @@ typedef float float4_t __attribute__((ext_vector_type(4))); using index_t = uint32_t; -__device__ void fused_multiply_accumulate(float& d, const float& s0, const float& s1) +template +__device__ void fused_multiply_accumulate(T& d, const T& s0, const T& s1) { d += s0 * s1; } diff --git a/composable_kernel/include/utility/config_nvidia.hpp.in b/composable_kernel/include/utility/config_nvidia.hpp.in index f077db991c..673c2778b1 100644 --- a/composable_kernel/include/utility/config_nvidia.hpp.in +++ b/composable_kernel/include/utility/config_nvidia.hpp.in @@ -22,7 +22,8 @@ using float4_t = float4; using index_t = uint32_t; -__device__ void fused_multiply_accumulate(float& d, const float& s0, const float& s1) +template +__device__ void fused_multiply_accumulate(T& d, const T& s0, const T& s1) { d += s0 * s1; } diff --git a/driver/include/conv_common.hpp b/driver/include/conv_common.hpp index d1ddb42317..636e222901 100644 --- a/driver/include/conv_common.hpp +++ b/driver/include/conv_common.hpp @@ -1,14 +1,14 @@ -#ifndef CK_CONV_COMMON_HPP -#define CK_CONV_COMMON_HPP +#ifndef CONV_COMMON_HPP +#define CONV_COMMON_HPP #include "ConstantTensorDescriptor.hpp" -using namespace ck; - // this is ugly, only for 4d template constexpr auto get_convolution_output_default_4d_tensor_descriptor(InDesc, WeiDesc) { + using namespace ck; + constexpr auto in_desc = InDesc{}; constexpr auto wei_desc = WeiDesc{}; @@ -45,6 +45,8 @@ template constexpr std::size_t calculate_convolution_flops(InDesc, WeiDesc, OutDesc) { + using namespace ck; + constexpr auto wei_desc = WeiDesc{}; constexpr auto out_desc = OutDesc{}; @@ -107,6 +111,8 @@ constexpr std::size_t calculate_convolution_flops(InDesc, WeiDesc, OutDesc) template constexpr std::size_t calculate_convolution_memory_size(Float, InDesc, WeiDesc, OutDesc) { + using namespace ck; + constexpr auto wei_desc = WeiDesc{}; constexpr auto out_desc = OutDesc{}; diff --git a/driver/include/device.hpp b/driver/include/device.hpp index faa4019a09..c43f14b751 100644 --- a/driver/include/device.hpp +++ b/driver/include/device.hpp @@ -1,11 +1,9 @@ -#ifndef CK_DEVICE_HPP -#define CK_DEVICE_HPP +#ifndef DEVICE_HPP +#define DEVICE_HPP #include #include "config.hpp" -using namespace ck; - struct DeviceMem { DeviceMem() = delete; diff --git a/driver/include/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp b/driver/include/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp index e4fa7ef0cd..7e0134069f 100644 --- a/driver/include/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp +++ b/driver/include/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp @@ -6,8 +6,6 @@ #include "gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp" #include "gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp" -using namespace ck; - template void device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw(InDesc, const Tensor& in_nchw, @@ -17,6 +15,8 @@ void device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw(InDesc, Tensor& out_nkhw, index_t nrepeat) { + using namespace ck; + constexpr auto I0 = Number<0>{}; constexpr auto I1 = Number<1>{}; constexpr auto I2 = Number<2>{}; @@ -135,7 +135,6 @@ void device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw(InDesc, WeiBlockCopyClusterLengths_C_K, WeiBlockCopyDataPerAccess_K>{}; -#if 1 float time = launch_kernel(run_gridwise_convolution_kernel, dim3(GridSize), dim3(BlockSize), @@ -149,7 +148,6 @@ void device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw(InDesc, (float)calculate_convolution_flops(InDesc{}, WeiDesc{}, OutDesc{}) / (std::size_t(1000) * 1000 * 1000) / time); usleep(std::min(time * 1000, float(10000))); -#endif } out_nkhw_device_buf.FromDevice(out_nkhw.mData.data()); diff --git a/driver/include/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp b/driver/include/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp index af5711a2dc..be4c41a6db 100644 --- a/driver/include/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp +++ b/driver/include/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp @@ -6,8 +6,6 @@ #include "gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp" #include "gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp" -using namespace ck; - template {}; constexpr auto I1 = Number<1>{}; constexpr auto I2 = Number<2>{}; diff --git a/driver/include/tensor.hpp b/driver/include/tensor.hpp index e7001b1022..804a1ff1ef 100644 --- a/driver/include/tensor.hpp +++ b/driver/include/tensor.hpp @@ -1,5 +1,5 @@ -#ifndef CK_TENSOR_HPP -#define CK_TENSOR_HPP +#ifndef TENSOR_HPP +#define TENSOR_HPP #include #include @@ -8,6 +8,7 @@ #include #include #include +#include "common_header.hpp" template std::ostream& LogRange(std::ostream& os, Range&& range, std::string delim) @@ -269,4 +270,46 @@ struct Tensor std::vector mData; }; +// this is ugly, only for 4d +template +void ostream_ConstantTensorDescriptor(TConstTensorDesc, std::ostream& os = std::cout) +{ + using namespace ck; + + static_assert(TConstTensorDesc::nDim == 4, "nDim is not 4"); + + constexpr auto I0 = Number<0>{}; + constexpr auto I1 = Number<1>{}; + constexpr auto I2 = Number<2>{}; + constexpr auto I3 = Number<3>{}; + constexpr auto desc = TConstTensorDesc{}; + + os << "Lengths: {" << desc.GetLength(I0) << ", " << desc.GetLength(I1) << ", " + << desc.GetLength(I2) << ", " << desc.GetLength(I3) << "}, " + << "Strides: {" << desc.GetStride(I0) << ", " << desc.GetStride(I1) << ", " + << desc.GetStride(I2) << ", " << desc.GetStride(I3) << "}" << std::endl; +} + +// this is ugly, only for 4d +template +auto make_TensorDescriptor(TConstTensorDesc) +{ + using namespace ck; + + static_assert(TConstTensorDesc::nDim == 4, "nDim is not 4"); + + constexpr auto I0 = Number<0>{}; + constexpr auto I1 = Number<1>{}; + constexpr auto I2 = Number<2>{}; + constexpr auto I3 = Number<3>{}; + constexpr auto desc = TConstTensorDesc{}; + + std::initializer_list lengths = { + desc.GetLength(I0), desc.GetLength(I1), desc.GetLength(I2), desc.GetLength(I3)}; + std::initializer_list strides = { + desc.GetStride(I0), desc.GetStride(I1), desc.GetStride(I2), desc.GetStride(I3)}; + + return TensorDescriptor(lengths, strides); +} + #endif diff --git a/driver/src/driver.cpp b/driver/src/driver.cpp index e8977fe6f8..45d6d1a789 100644 --- a/driver/src/driver.cpp +++ b/driver/src/driver.cpp @@ -14,8 +14,6 @@ #include "device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp" #include "device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp" -using namespace ck; - struct GeneratorTensor_1 { template @@ -65,44 +63,6 @@ struct GeneratorTensor_Checkboard } }; -// this is ugly, only for 4d -template -void ostream_ConstantTensorDescriptor(TConstTensorDesc, std::ostream& os = std::cout) -{ - static_assert(TConstTensorDesc::nDim == 4, "nDim is not 4"); - - constexpr auto I0 = Number<0>{}; - constexpr auto I1 = Number<1>{}; - constexpr auto I2 = Number<2>{}; - constexpr auto I3 = Number<3>{}; - constexpr auto desc = TConstTensorDesc{}; - - os << "Lengths: {" << desc.GetLength(I0) << ", " << desc.GetLength(I1) << ", " - << desc.GetLength(I2) << ", " << desc.GetLength(I3) << "}, " - << "Strides: {" << desc.GetStride(I0) << ", " << desc.GetStride(I1) << ", " - << desc.GetStride(I2) << ", " << desc.GetStride(I3) << "}" << std::endl; -} - -// this is ugly, only for 4d -template -auto make_TensorDescriptor(TConstTensorDesc) -{ - static_assert(TConstTensorDesc::nDim == 4, "nDim is not 4"); - - constexpr auto I0 = Number<0>{}; - constexpr auto I1 = Number<1>{}; - constexpr auto I2 = Number<2>{}; - constexpr auto I3 = Number<3>{}; - constexpr auto desc = TConstTensorDesc{}; - - std::initializer_list lengths = { - desc.GetLength(I0), desc.GetLength(I1), desc.GetLength(I2), desc.GetLength(I3)}; - std::initializer_list strides = { - desc.GetStride(I0), desc.GetStride(I1), desc.GetStride(I2), desc.GetStride(I3)}; - - return TensorDescriptor(lengths, strides); -} - template & ref, const Tensor& result) int main(int argc, char* argv[]) { + using namespace ck; + #if 0 constexpr index_t N = 8; constexpr index_t C = 16; @@ -611,7 +573,7 @@ int main(int argc, char* argv[]) constexpr index_t HPad = 0; constexpr index_t WPad = 0; -#elif 0 +#elif 1 // 1x1 filter, 8x8 image // cudnn@V100 77%, ck@V100 76%, ck@P100 79%, ck@VII 51% constexpr index_t N = 128; @@ -787,7 +749,7 @@ int main(int argc, char* argv[]) constexpr index_t HPad = 0; constexpr index_t WPad = 0; -#elif 1 +#elif 0 // 1x1 filter, 7x7 image // cudnn@V100 49%, ck@V100 50%, ck@P100 61%, ck@VII 52% constexpr index_t N = 128; @@ -859,30 +821,31 @@ int main(int argc, char* argv[]) #endif } -#if 1 #if 0 device_convolution_direct_v2_nchw_kcyx_nkhw + in_nchw_desc, in_nchw, wei_kcyx_desc, wei_kcyx, out_nkhw_desc, out_nkhw_device, nrepeat); #elif 0 device_convolution_implicit_gemm_v1_chwn_cyxk_khwn + in_nchw_desc, in_nchw, wei_kcyx_desc, wei_kcyx, out_nkhw_desc, out_nkhw_device, nrepeat); #elif 0 device_convolution_implicit_gemm_v1_nchw_cyxk_nkhw + in_nchw_desc, in_nchw, wei_kcyx_desc, wei_kcyx, out_nkhw_desc, out_nkhw_device, nrepeat); #elif 0 device_convolution_implicit_gemm_v2_chwn_cyxk_khwn + in_nchw_desc, in_nchw, wei_kcyx_desc, wei_kcyx, out_nkhw_desc, out_nkhw_device, nrepeat); #elif 0 - device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw + device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw( + in_nchw_desc, in_nchw, wei_kcyx_desc, wei_kcyx, out_nkhw_desc, out_nkhw_device, nrepeat); #elif 1 - device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw -#endif - (in_nchw_desc, - in_nchw, - wei_kcyx_desc, - wei_kcyx, - out_nkhw_desc, - out_nkhw_device, - ConvStrides{}, - ConvDilations{}, - nrepeat); - + device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw(in_nchw_desc, + in_nchw, + wei_kcyx_desc, + wei_kcyx, + out_nkhw_desc, + out_nkhw_device, + ConvStrides{}, + ConvDilations{}, + nrepeat); #elif 0 device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded(in_nchw_desc, in_nchw,