From 2c93b3057dbfd46600674820888570c7c45857cc Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Sun, 15 Sep 2019 16:31:54 -0500 Subject: [PATCH] initial implementation for nchw v4r4 padding --- ...tion_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp | 6 +- ..._v4r1_nchw_kcyx_nkhw_lds_double_buffer.hpp | 6 +- ...plicit_gemm_v4r1_nchw_kcyx_nkhw_padded.hpp | 6 +- ...tion_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp | 4 +- ..._v4r4_nchw_kcyx_nkhw_lds_double_buffer.hpp | 4 +- ...plicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp | 457 ++++++++++++++++ .../tensor_coordinate_v2.hpp | 6 +- .../tensor_description/tensor_descriptor.hpp | 22 +- ...plicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp | 225 ++++++++ driver/src/driver.cpp | 1 + driver/src/driver.cu | 487 +----------------- 11 files changed, 721 insertions(+), 503 deletions(-) create mode 100644 composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp create mode 100644 driver/include/device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp mode change 100644 => 120000 driver/src/driver.cu diff --git a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp index 04785a4a06..d00c92ff9e 100644 --- a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp +++ b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp @@ -100,10 +100,8 @@ struct GridwiseConvolutionImplicitGemm_v4r1_nchw_kcyx_nkhw constexpr index_t E = C * Y * X; // sanity-check for vectorized memory load - static_assert(ConvStrideW == 1 || InBlockCopySrcDataPerRead_B == 1, - "wrong! global vector load of input tensor is wrong"); - - static_assert((X == 1 || ConvDilationW % InBlockCopySrcDataPerRead_B == 0), + static_assert((Ho == 1 || ConvStrideW % InBlockCopySrcDataPerRead_B == 0) && + (X == 1 || ConvDilationW % InBlockCopySrcDataPerRead_B == 0), "wrong! aligment requirement for vectorized global load of input tensor will " "be violated"); 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 629e822dcb..4191de7880 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 @@ -100,10 +100,8 @@ struct GridwiseConvolutionImplicitGemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer constexpr index_t E = C * Y * X; // sanity-check for vectorized memory load - static_assert(ConvStrideW == 1 || InBlockCopySrcDataPerRead_B == 1, - "wrong! global vector load of input tensor is wrong"); - - static_assert((X == 1 || ConvDilationW % InBlockCopySrcDataPerRead_B == 0), + static_assert((Ho == 1 || ConvStrideW % InBlockCopySrcDataPerRead_B == 0) && + (X == 1 || ConvDilationW % InBlockCopySrcDataPerRead_B == 0), "wrong! aligment requirement for vectorized global load of input tensor will " "be violated"); diff --git a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_padded.hpp b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_padded.hpp index cde9ce509a..959a9112d3 100644 --- a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_padded.hpp +++ b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_padded.hpp @@ -107,10 +107,8 @@ struct GridwiseConvolutionImplicitGemm_v4r1_nchw_kcyx_nkhw_padded constexpr index_t E = C * Y * X; // sanity-check for vectorized memory load - static_assert(ConvStrideW == 1 || InBlockCopySrcDataPerRead_B == 1, - "wrong! global vector load of input tensor is wrong"); - - static_assert((X == 1 || ConvDilationW % InBlockCopySrcDataPerRead_B == 0), + static_assert((Ho == 1 || ConvStrideW % InBlockCopySrcDataPerRead_B == 0) && + (X == 1 || ConvDilationW % InBlockCopySrcDataPerRead_B == 0), "wrong! aligment requirement for vectorized global load of input tensor will " "be violated"); diff --git a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp index 7989f7fb19..1ec519aad6 100644 --- a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp +++ b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp @@ -83,7 +83,9 @@ struct GridwiseConvolutionImplicitGemm_v4r4_nchw_kcyx_nkhw constexpr index_t E = C * Y * X; constexpr index_t B = N * Ho * Wo; - static_assert((X == 1 || ConvDilationW % InBlockCopyDataPerAccess_B == 0), + // sanity-check for vectorized memory load + static_assert((Ho == 1 || ConvStrideW % InBlockCopyDataPerAccess_B == 0) && + (X == 1 || ConvDilationW % InBlockCopyDataPerAccess_B == 0), "wrong! aligment requirement for vectorized global load of input tensor will " "be violated"); diff --git a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer.hpp b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer.hpp index eeb63bf8c8..ad411be6db 100644 --- a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer.hpp +++ b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer.hpp @@ -83,7 +83,9 @@ struct GridwiseConvolutionImplicitGemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer constexpr index_t E = C * Y * X; constexpr index_t B = N * Ho * Wo; - static_assert((X == 1 || ConvDilationW % InBlockCopyDataPerAccess_B == 0), + // sanity-check for vectorized memory load + static_assert((Ho == 1 || ConvStrideW % InBlockCopyDataPerAccess_B == 0) && + (X == 1 || ConvDilationW % InBlockCopyDataPerAccess_B == 0), "wrong! aligment requirement for vectorized global load of input tensor will " "be violated"); diff --git a/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp new file mode 100644 index 0000000000..0ebfa08f35 --- /dev/null +++ b/composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp @@ -0,0 +1,457 @@ +#ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V4R4_NCHW_KCYX_NKHW_PADDED_HPP +#define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V4R4_NCHW_KCYX_NKHW_PADDED_HPP + +#include "common_header.hpp" +#include "ConstantTensorDescriptor.hpp" +#include "ConstantMergedTensorDescriptor.hpp" +#include "ConstantMatrixDescriptor.hpp" +#include "blockwise_generic_tensor_slice_copy.hpp" +#include "blockwise_gemm.hpp" +#include "threadwise_generic_tensor_slice_copy.hpp" + +namespace ck { + +// B = merge(N, Ho, Wo) +template +struct GridwiseConvolutionImplicitGemm_v4r4_nchw_kcyx_nkhw_padded +{ +#if 1 + __device__ void Run(const Float* const __restrict__ p_in_global, + const Float* const __restrict__ p_wei_global, + Float* const __restrict__ p_out_global) const + { + constexpr auto I0 = Number<0>{}; + constexpr auto I1 = Number<1>{}; + constexpr auto I2 = Number<2>{}; + constexpr auto I3 = Number<3>{}; + constexpr auto I5 = Number<5>{}; + + constexpr auto True = integral_constant{}; + + constexpr auto in_n_c_hi_wi_global_desc = + make_native_tensor_descriptor(InGlobalDesc::GetLengths(), InGlobalDesc::GetStrides()); + constexpr auto wei_k_c_y_x_global_desc = + make_native_tensor_descriptor(WeiGlobalDesc::GetLengths(), WeiGlobalDesc::GetStrides()); + constexpr auto out_n_k_ho_wo_global_desc = + make_native_tensor_descriptor(OutGlobalDesc::GetLengths(), OutGlobalDesc::GetStrides()); + + constexpr index_t N = in_n_c_hi_wi_global_desc.GetLength(I0); + constexpr index_t C = in_n_c_hi_wi_global_desc.GetLength(I1); + constexpr index_t Hi = in_n_c_hi_wi_global_desc.GetLength(I2); + constexpr index_t Wi = in_n_c_hi_wi_global_desc.GetLength(I3); + + constexpr index_t K = out_n_k_ho_wo_global_desc.GetLength(I1); + constexpr index_t Ho = out_n_k_ho_wo_global_desc.GetLength(I2); + constexpr index_t Wo = out_n_k_ho_wo_global_desc.GetLength(I3); + + constexpr index_t Y = wei_k_c_y_x_global_desc.GetLength(I2); + constexpr index_t X = wei_k_c_y_x_global_desc.GetLength(I3); + + constexpr index_t ConvStrideH = ConvStrides{}[0]; + constexpr index_t ConvStrideW = ConvStrides{}[1]; + + constexpr index_t ConvDilationH = ConvDilations{}[0]; + constexpr index_t ConvDilationW = ConvDilations{}[1]; + + constexpr index_t E = C * Y * X; + constexpr index_t B = N * Ho * Wo; + + // sanity-check for vectorized memory load + static_assert((Ho == 1 || ConvStrideW % InBlockCopyDataPerAccess_B == 0) && + (X == 1 || ConvDilationW % InBlockCopyDataPerAccess_B == 0), + "wrong! aligment requirement for vectorized global load of input tensor will " + "be violated"); + + // divide block work by [K, B] + static_assert(K % KPerBlock == 0 && B % BPerBlock == 0 && E % EPerBlock == 0, + "wrong! cannot divide work evenly among block"); + + constexpr index_t KBlockWork = K / KPerBlock; + constexpr index_t BBlockWork = B / BPerBlock; + + constexpr auto block_work_desc = + make_ConstantTensorDescriptor_packed(Sequence{}); + + const auto block_work_multi_id = + block_work_desc.GetMultiIndexFrom1dIndex(get_block_1d_id()); + + const index_t k_block_data_on_global = block_work_multi_id[0] * KPerBlock; + const index_t b_block_data_on_global = block_work_multi_id[1] * BPerBlock; + + // input tensor + // global mem + constexpr auto in_n_c_hip_wip_global_desc = transform_tensor_descriptor( + in_n_c_hi_wi_global_desc, + make_tuple( + PassThrough{}, PassThrough{}, Pad, LeftPads, RightPads>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{})); + + constexpr auto in_n_c_y_ho_x_wo_global_desc = transform_tensor_descriptor( + in_n_c_hip_wip_global_desc, + make_tuple(PassThrough{}, + PassThrough{}, + Embed, Sequence>{}, + Embed, Sequence>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{})); + + constexpr auto in_e_b_global_desc = transform_tensor_descriptor( + in_n_c_y_ho_x_wo_global_desc, + make_tuple(Merge>{}, Merge>{}), + make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}), + make_tuple(Sequence<0>{}, Sequence<1>{})); + + // LDS mem + // be careful of LDS alignment + constexpr auto in_e_b_block_desc = + make_native_tensor_descriptor_packed(Sequence{}); + + // input blockwise copy + auto blockwise_in_copy = + BlockwiseGenericTensorSliceCopy_v4( + {0, b_block_data_on_global}, {0, 0}); + + // weight tensor + // global mem + constexpr auto wei_e_k_global_desc = + transform_tensor_descriptor(wei_k_c_y_x_global_desc, + make_tuple(Merge>{}, PassThrough{}), + make_tuple(Sequence<1, 2, 3>{}, Sequence<0>{}), + make_tuple(Sequence<0>{}, Sequence<1>{})); + + // LDS + // be careful of LDS alignment + constexpr auto wei_e_k_block_desc = make_native_tensor_descriptor_aligned( + Sequence{}, + Number{}); + + // weight blockwise copy + auto blockwise_wei_copy = + BlockwiseGenericTensorSliceCopy_v4( + {0, k_block_data_on_global}, {0, 0}); + + // GEMM definition + // c_mtx += transpose(a_mtx) * b_mtx + // a_mtx[EPerBlock, KPerBlock] is in LDS + // b_mtx[EPerBlocl, BPerBlock] is in LDS + // c_mtx[KPerBlock, BPerBlock] is distributed among threads, and saved in + // register + constexpr auto a_e_k_block_mtx_desc = make_ConstantMatrixDescriptor(wei_e_k_block_desc); + + constexpr auto b_e_b_block_mtx_desc = make_ConstantMatrixDescriptor(in_e_b_block_desc); + + // sanity check + static_assert( + KPerBlock % (GemmMPerThreadSubC * GemmMLevel0Cluster * GemmMLevel1Cluster) == 0 && + BPerBlock % (GemmNPerThreadSubC * GemmNLevel0Cluster * GemmNLevel1Cluster) == 0, + "wrong!"); + + constexpr index_t GemmMRepeat = + KPerBlock / (GemmMPerThreadSubC * GemmMLevel0Cluster * GemmMLevel1Cluster); + + constexpr index_t GemmNRepeat = + BPerBlock / (GemmNPerThreadSubC * GemmNLevel0Cluster * GemmNLevel1Cluster); + + // c_thread_mtx definition: this is a mess + // TODO:: more elegent way of defining c_thread_mtx + constexpr auto c_k0k1_b0b1_thread_mtx_desc = make_ConstantMatrixDescriptor_packed( + Number{}, Number{}); + + const auto blockwise_gemm = BlockwiseGemmBlockABlockBThreadCTransANormalBNormalC_v2< + BlockSize, + decltype(a_e_k_block_mtx_desc), + decltype(b_e_b_block_mtx_desc), + decltype(c_k0k1_b0b1_thread_mtx_desc), + GemmMPerThreadSubC, + GemmNPerThreadSubC, + GemmMLevel0Cluster, + GemmNLevel0Cluster, + GemmMLevel1Cluster, + GemmNLevel1Cluster, + GemmKPerThreadLoop, + GemmDataPerReadA, + GemmDataPerReadB>{}; + + // LDS allocation for input and weight: be careful of alignment + constexpr index_t max_align = math::lcm(InBlockCopyDataPerAccess_B, + WeiBlockCopyDstDataPerWrite_K, + GemmDataPerReadA, + GemmDataPerReadB); + + constexpr index_t in_block_space = + math::integer_least_multiple(in_e_b_block_desc.GetElementSpace(), max_align); + + constexpr index_t wei_block_space = + math::integer_least_multiple(wei_e_k_block_desc.GetElementSpace(), max_align); + + __shared__ Float p_in_block[in_block_space]; + __shared__ Float p_wei_block[wei_block_space]; + + // register allocation for output + Float p_out_thread[c_k0k1_b0b1_thread_mtx_desc.GetElementSpace()]; + + // zero out threadwise output + threadwise_matrix_set_zero(c_k0k1_b0b1_thread_mtx_desc, p_out_thread); + + for(index_t e_block_data_begin = 0; e_block_data_begin < E; e_block_data_begin += EPerBlock) + { + blockwise_in_copy.Run(p_in_global, p_in_block); + blockwise_wei_copy.Run(p_wei_global, p_wei_block); + + __syncthreads(); + + blockwise_gemm.Run(p_wei_block, p_in_block, p_out_thread); + + __syncthreads(); + + blockwise_in_copy.MoveSrcSliceWindow(make_multi_index(EPerBlock, 0), True); + blockwise_wei_copy.MoveSrcSliceWindow(make_multi_index(EPerBlock, 0), True); + } + + // copy output: register to global memory + { + // calculate origin of thread output tensor on global memory + // blockwise GEMM c matrix starting index + const auto c_thread_mtx_on_block = + blockwise_gemm.GetBeginOfThreadMatrixC(get_thread_local_1d_id()); + + const index_t k_thread_data_on_global = + k_block_data_on_global + c_thread_mtx_on_block.row; + + const index_t b_thread_data_on_global = + b_block_data_on_global + c_thread_mtx_on_block.col; + + // src descriptor + constexpr auto out_k0_k1_b0_b1_thread_desc = make_native_tensor_descriptor_packed( + Sequence{}); + + // dst descriptor + constexpr index_t K1 = GemmMPerThreadSubC * GemmMLevel0Cluster * GemmMLevel1Cluster; + constexpr index_t B1 = GemmNPerThreadSubC * GemmNLevel0Cluster * GemmNLevel1Cluster; + + constexpr index_t K0 = K / K1; + constexpr index_t B0 = B / B1; + + constexpr auto out_k_b_global_desc = transform_tensor_descriptor( + out_n_k_ho_wo_global_desc, + make_tuple(PassThrough{}, Merge>{}), + make_tuple(Sequence<1>{}, Sequence<0, 2, 3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{})); + + constexpr auto out_k0_k1_b0_b1_global_desc = transform_tensor_descriptor( + out_k_b_global_desc, + make_tuple(Unmerge>{}, Unmerge>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}), + make_tuple(Sequence<0, 1>{}, Sequence<2, 3>{})); + + // output threadwise copy + auto threadwise_out_copy = ThreadwiseGenericTensorSliceCopy_v4r2< + decltype(out_k0_k1_b0_b1_thread_desc), + decltype(out_k0_k1_b0_b1_global_desc), + decltype(out_k0_k1_b0_b1_thread_desc.GetLengths()), + arithmetic_sequence_gen<0, 4, 1>::type, + 3, + OutThreadCopyDataPerAccess_B, + OutThreadCopyDataPerAccess_B>({0, 0, 0, 0}, + {k_thread_data_on_global / K1, + k_thread_data_on_global % K1, + b_thread_data_on_global / B1, + b_thread_data_on_global % B1}); + + threadwise_out_copy.Run(p_out_thread, p_out_global); + } + } +#else + __device__ void Run(const Float* const __restrict__ p_in_global, + const Float* const __restrict__ p_wei_global, + Float* const __restrict__ p_out_global) const + { + constexpr auto I0 = Number<0>{}; + constexpr auto I1 = Number<1>{}; + constexpr auto I2 = Number<2>{}; + constexpr auto I3 = Number<3>{}; + constexpr auto I5 = Number<5>{}; + + constexpr auto True = integral_constant{}; + + constexpr auto in_n_c_hi_wi_global_desc = + make_native_tensor_descriptor(InGlobalDesc::GetLengths(), InGlobalDesc::GetStrides()); + constexpr auto wei_k_c_y_x_global_desc = + make_native_tensor_descriptor(WeiGlobalDesc::GetLengths(), WeiGlobalDesc::GetStrides()); + constexpr auto out_n_k_ho_wo_global_desc = + make_native_tensor_descriptor(OutGlobalDesc::GetLengths(), OutGlobalDesc::GetStrides()); + + constexpr index_t N = in_n_c_hi_wi_global_desc.GetLength(I0); + constexpr index_t C = in_n_c_hi_wi_global_desc.GetLength(I1); + constexpr index_t Hi = in_n_c_hi_wi_global_desc.GetLength(I2); + constexpr index_t Wi = in_n_c_hi_wi_global_desc.GetLength(I3); + + constexpr index_t K = out_n_k_ho_wo_global_desc.GetLength(I1); + constexpr index_t Ho = out_n_k_ho_wo_global_desc.GetLength(I2); + constexpr index_t Wo = out_n_k_ho_wo_global_desc.GetLength(I3); + + constexpr index_t Y = wei_k_c_y_x_global_desc.GetLength(I2); + constexpr index_t X = wei_k_c_y_x_global_desc.GetLength(I3); + + constexpr index_t ConvStrideH = ConvStrides{}[0]; + constexpr index_t ConvStrideW = ConvStrides{}[1]; + + constexpr index_t ConvDilationH = ConvDilations{}[0]; + constexpr index_t ConvDilationW = ConvDilations{}[1]; + + constexpr index_t E = C * Y * X; + constexpr index_t B = N * Ho * Wo; + + // sanity-check for vectorized memory load + static_assert((Ho == 1 || ConvStrideW % InBlockCopyDataPerAccess_B == 0) && + (X == 1 || ConvDilationW % InBlockCopyDataPerAccess_B == 0), + "wrong! aligment requirement for vectorized global load of input tensor will " + "be violated"); + + // input tensor + constexpr auto in_n_c_hip_wip_global_desc = transform_tensor_descriptor( + in_n_c_hi_wi_global_desc, + make_tuple( + PassThrough{}, PassThrough{}, Pad, LeftPads, RightPads>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{})); + + constexpr auto in_n_c_y_ho_x_wo_global_desc = transform_tensor_descriptor( + in_n_c_hip_wip_global_desc, + make_tuple(PassThrough{}, + PassThrough{}, + Embed, Sequence>{}, + Embed, Sequence>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{})); + + constexpr auto in_e_b_global_desc = transform_tensor_descriptor( + in_n_c_y_ho_x_wo_global_desc, + make_tuple(Merge>{}, Merge>{}), + make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}), + make_tuple(Sequence<0>{}, Sequence<1>{})); + + // output tensor + constexpr auto out_k_b_global_desc = + transform_tensor_descriptor(out_n_k_ho_wo_global_desc, + make_tuple(PassThrough{}, Merge>{}), + make_tuple(Sequence<1>{}, Sequence<0, 2, 3>{}), + make_tuple(Sequence<0>{}, Sequence<1>{})); + + constexpr index_t K1 = GemmMPerThreadSubC * GemmMLevel0Cluster * GemmMLevel1Cluster; + constexpr index_t B1 = GemmNPerThreadSubC * GemmNLevel0Cluster * GemmNLevel1Cluster; + + constexpr index_t K0 = K / K1; + constexpr index_t B0 = B / B1; + + constexpr auto out_k0_k1_b0_b1_global_desc = transform_tensor_descriptor( + out_k_b_global_desc, + make_tuple(Unmerge>{}, Unmerge>{}), + make_tuple(Sequence<0>{}, Sequence<1>{}), + make_tuple(Sequence<0, 1>{}, Sequence<2, 3>{})); + +#if 1 + if(get_thread_local_1d_id() == 0 && get_block_1d_id() == 0) + { + print_tensor_descriptor("in_e_b_global_desc: ", in_e_b_global_desc); + print_tensor_descriptor("in_n_c_y_ho_x_wo_global_desc: ", in_n_c_y_ho_x_wo_global_desc); + print_tensor_descriptor("in_n_c_hip_wip_global_desc: ", in_n_c_hip_wip_global_desc); + print_tensor_descriptor("in_n_c_hi_wi_global_desc: ", in_n_c_hi_wi_global_desc); + + auto coord3 = make_tensor_coordinate_v2(in_e_b_global_desc, {1, 1}); + + auto idx3 = coord3.GetIndex(); + auto idx2 = coord3.GetLowerCoordinate().GetIndex(); + auto idx1 = coord3.GetLowerCoordinate().GetLowerCoordinate().GetIndex(); + auto idx0 = + coord3.GetLowerCoordinate().GetLowerCoordinate().GetLowerCoordinate().GetIndex(); + + print_array("idx3: ", idx3); + print_array("idx2: ", idx2); + print_array("idx1: ", idx1); + print_array("idx0: ", idx0); + } + + if(get_thread_local_1d_id() == 0 && get_block_1d_id() == 0) + { + print_tensor_descriptor("out_k0_k1_b0_b1_global_desc: ", out_k0_k1_b0_b1_global_desc); + print_tensor_descriptor("out_k_b_global_desc: ", out_k_b_global_desc); + print_tensor_descriptor("out_n_k_ho_wo_global_desc: ", out_n_k_ho_wo_global_desc); + + auto coord2 = make_tensor_coordinate_v2(out_k0_k1_b0_b1_global_desc, {1, 1, 1, 1}); + + auto idx2 = coord2.GetIndex(); + auto idx1 = coord2.GetLowerCoordinate().GetIndex(); + auto idx0 = coord2.GetLowerCoordinate().GetLowerCoordinate().GetIndex(); + + print_array("idx2: ", idx2); + print_array("idx1: ", idx1); + print_array("idx0: ", idx0); + } +#endif + } +#endif +}; + +} // namespace ck +#endif diff --git a/composable_kernel/include/tensor_description/tensor_coordinate_v2.hpp b/composable_kernel/include/tensor_description/tensor_coordinate_v2.hpp index b7191f6c8b..2a35457bda 100644 --- a/composable_kernel/include/tensor_description/tensor_coordinate_v2.hpp +++ b/composable_kernel/include/tensor_description/tensor_coordinate_v2.hpp @@ -190,14 +190,16 @@ struct TensorCoordinate_v2 __host__ __device__ static constexpr auto MakeDummyTensorCoordinate(NativeTensorDescriptor) { - return NativeTensorCoordinate>(); + return NativeTensorCoordinate>( + make_zero_array()); } template __host__ __device__ static constexpr auto MakeDummyTensorCoordinate(TransformedTensorDescriptor) { - return TransformedTensorCoordinate>(); + return TransformedTensorCoordinate>( + make_zero_array()); } public: diff --git a/composable_kernel/include/tensor_description/tensor_descriptor.hpp b/composable_kernel/include/tensor_description/tensor_descriptor.hpp index 0113997a80..a41dc1da67 100644 --- a/composable_kernel/include/tensor_description/tensor_descriptor.hpp +++ b/composable_kernel/include/tensor_description/tensor_descriptor.hpp @@ -187,8 +187,28 @@ struct TransformedTensorDescriptor nTransform == UpDimensionIds::Size(), "wrong! # of transformations not the same"); - // TODO: sanity check: LowDimensionIds should include all low-dimensions, + // sanity check: + // LowDimensionIds should include all low-dimensions, // UpDimensionIds should include all up-dimensions + using mingled_up_dimension_ids = + decltype(unpack(lambda_merge_sequences{}, UpDimensionIds{})); + + using sorted_up_dimension_ids = + typename sequence_sort>::type; + + static_assert(sorted_up_dimension_ids::Size() == nDimUp && + is_valid_sequence_map{}, + "wrong! UpDimensionIds is not configured correctly"); + + using mingled_low_dimension_ids = + decltype(unpack(lambda_merge_sequences{}, LowDimensionIds{})); + + using sorted_low_dimension_ids = + typename sequence_sort>::type; + + static_assert(sorted_low_dimension_ids::Size() == nDimLow && + is_valid_sequence_map{}, + "wrong! LowDimensionIds is not configured correctly"); // TODO: sanity check: while a up-dimension could be associated with multille // transformation, a low-dimension should be associated with only one transformation diff --git a/driver/include/device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp b/driver/include/device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp new file mode 100644 index 0000000000..d43008c0df --- /dev/null +++ b/driver/include/device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp @@ -0,0 +1,225 @@ +#pragma once +#include +#include "device.hpp" +#include "tensor.hpp" +#include "gridwise_convolution_kernel_wrapper.hpp" +#include "gridwise_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp" + +template +void device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded(InDesc, + const Tensor& in_nchw, + WeiDesc, + const Tensor& wei_kcyx, + OutDesc, + Tensor& out_nkhw, + ConvStrides, + ConvDilations, + LeftPads, + RightPads, + index_t nrepeat) +{ + using namespace ck; + + constexpr auto I0 = Number<0>{}; + constexpr auto I1 = Number<1>{}; + constexpr auto I2 = Number<2>{}; + constexpr auto I3 = Number<3>{}; + + constexpr auto in_nchw_desc = InDesc{}; + constexpr auto wei_kcyx_desc = WeiDesc{}; + constexpr auto out_nkhw_desc = OutDesc{}; + + constexpr index_t N = out_nkhw_desc.GetLength(I0); + constexpr index_t K = out_nkhw_desc.GetLength(I1); + constexpr index_t Ho = out_nkhw_desc.GetLength(I2); + constexpr index_t Wo = out_nkhw_desc.GetLength(I3); + + std::size_t data_sz = sizeof(T); + DeviceMem in_nchw_device_buf(data_sz * in_nchw.mDesc.GetElementSpace()); + DeviceMem wei_kcyx_device_buf(data_sz * wei_kcyx.mDesc.GetElementSpace()); + DeviceMem out_nkhw_device_buf(data_sz * out_nkhw.mDesc.GetElementSpace()); + + in_nchw_device_buf.ToDevice(in_nchw.mData.data()); + wei_kcyx_device_buf.ToDevice(wei_kcyx.mData.data()); + out_nkhw_device_buf.ToDevice(out_nkhw.mData.data()); + +#if 1 + constexpr index_t BlockSize = 256; + + constexpr index_t BPerBlock = 128; + constexpr index_t KPerBlock = 128; + constexpr index_t EPerBlock = 8; + + constexpr index_t GemmMPerThreadSubC = 4; + constexpr index_t GemmNPerThreadSubC = 4; + constexpr index_t GemmMLevel0Cluster = 4; + constexpr index_t GemmNLevel0Cluster = 4; + constexpr index_t GemmMLevel1Cluster = 4; + constexpr index_t GemmNLevel1Cluster = 4; + constexpr index_t GemmKPerThreadLoop = 1; + constexpr index_t GemmDataPerReadA = 4; + constexpr index_t GemmDataPerReadB = 4; + + using InBlockCopySubLengths_E_B = Sequence<4, 1>; + using InBlockCopyClusterLengths_E_B = Sequence<2, 128>; + using InBlockCopyThreadClusterArrangeOrder = Sequence<0, 1>; // [E, B] + using InBlockCopySrcAccessOrder = Sequence<0, 1>; // [E, B] + using InBlockCopyDstAccessOrder = Sequence<0, 1>; // [E, B] + + constexpr index_t InBlockCopyDataPerAccess_B = 1; + + using WeiBlockCopySubLengths_E_K = Sequence<4, 1>; + using WeiBlockCopyClusterLengths_E_K = Sequence<2, 128>; + using WeiBlockCopyThreadClusterArrangeOrder = Sequence<1, 0>; // [K, E] + using WeiBlockCopySrcAccessOrder = Sequence<1, 0>; // [K, E] + using WeiBlockCopyDstAccessOrder = Sequence<0, 1>; // [E, K] + + constexpr index_t WeiBlockCopySrcDataPerRead_E = 4; + constexpr index_t WeiBlockCopyDstDataPerWrite_K = 1; + + constexpr index_t OutThreadCopyDataPerAccess_B = 1; +#elif 1 + // 1x1 filter, 8x8 image + constexpr index_t BlockSize = 256; + + constexpr index_t BPerBlock = 128; + constexpr index_t KPerBlock = 128; + constexpr index_t EPerBlock = 8; + + constexpr index_t GemmMPerThreadSubC = 4; + constexpr index_t GemmNPerThreadSubC = 4; + constexpr index_t GemmMLevel0Cluster = 4; + constexpr index_t GemmNLevel0Cluster = 4; + constexpr index_t GemmMLevel1Cluster = 4; + constexpr index_t GemmNLevel1Cluster = 4; + constexpr index_t GemmKPerThreadLoop = 1; + constexpr index_t GemmDataPerReadA = 4; + constexpr index_t GemmDataPerReadB = 4; + + using InBlockCopySubLengths_E_B = Sequence<1, 4>; + using InBlockCopyClusterLengths_E_B = Sequence<8, 32>; + using InBlockCopyThreadClusterArrangeOrder = Sequence<0, 1>; // [E, B] + using InBlockCopySrcAccessOrder = Sequence<0, 1>; // [E, B] + using InBlockCopyDstAccessOrder = Sequence<0, 1>; // [E, B] + + constexpr index_t InBlockCopyDataPerAccess_B = 4; + + using WeiBlockCopySubLengths_E_K = Sequence<4, 1>; + using WeiBlockCopyClusterLengths_E_K = Sequence<2, 128>; + using WeiBlockCopyThreadClusterArrangeOrder = Sequence<1, 0>; // [K, E] + using WeiBlockCopySrcAccessOrder = Sequence<1, 0>; // [K, E] + using WeiBlockCopyDstAccessOrder = Sequence<0, 1>; // [E, K] + + constexpr index_t WeiBlockCopySrcDataPerRead_E = 4; + constexpr index_t WeiBlockCopyDstDataPerWrite_K = 1; + + constexpr index_t OutThreadCopyDataPerAccess_B = 4; +#elif 0 + // 1x1 filter, 14x14 image + constexpr index_t BlockSize = 256; + + constexpr index_t BPerBlock = 128; + constexpr index_t KPerBlock = 128; + constexpr index_t EPerBlock = 8; + + constexpr index_t GemmMPerThreadSubC = 4; + constexpr index_t GemmNPerThreadSubC = 4; + constexpr index_t GemmMLevel0Cluster = 4; + constexpr index_t GemmNLevel0Cluster = 4; + constexpr index_t GemmMLevel1Cluster = 4; + constexpr index_t GemmNLevel1Cluster = 4; + constexpr index_t GemmKPerThreadLoop = 1; + constexpr index_t GemmDataPerReadA = 4; + constexpr index_t GemmDataPerReadB = 4; + + using InBlockCopySubLengths_E_B = Sequence<2, 2>; + using InBlockCopyClusterLengths_E_B = Sequence<4, 64>; + using InBlockCopyThreadClusterArrangeOrder = Sequence<0, 1>; // [E, B] + using InBlockCopySrcAccessOrder = Sequence<0, 1>; // [E, B] + using InBlockCopyDstAccessOrder = Sequence<0, 1>; // [E, B] + + constexpr index_t InBlockCopyDataPerAccess_B = 2; + + using WeiBlockCopySubLengths_E_K = Sequence<4, 1>; + using WeiBlockCopyClusterLengths_E_K = Sequence<2, 128>; + using WeiBlockCopyThreadClusterArrangeOrder = Sequence<1, 0>; // [K, E] + using WeiBlockCopySrcAccessOrder = Sequence<1, 0>; // [K, E] + using WeiBlockCopyDstAccessOrder = Sequence<0, 1>; // [E, K] + + constexpr index_t WeiBlockCopySrcDataPerRead_E = 4; + constexpr index_t WeiBlockCopyDstDataPerWrite_K = 1; + + constexpr index_t OutThreadCopyDataPerAccess_B = 2; +#endif + + constexpr index_t B = N * Ho * Wo; + + constexpr index_t GridSize = + ((B + BPerBlock - 1) / BPerBlock) * ((K + KPerBlock - 1) / KPerBlock); + + printf("%s: BlockSize %u, GridSize %u \n", __func__, BlockSize, GridSize); + + constexpr auto gridwise_conv = GridwiseConvolutionImplicitGemm_v4r4_nchw_kcyx_nkhw_padded< + GridSize, + BlockSize, + T, + decltype(in_nchw_desc), + decltype(wei_kcyx_desc), + decltype(out_nkhw_desc), + ConvStrides, + ConvDilations, + LeftPads, + RightPads, + BPerBlock, + KPerBlock, + EPerBlock, + GemmMPerThreadSubC, + GemmNPerThreadSubC, + GemmMLevel0Cluster, + GemmNLevel0Cluster, + GemmMLevel1Cluster, + GemmNLevel1Cluster, + GemmKPerThreadLoop, + GemmDataPerReadA, + GemmDataPerReadB, + InBlockCopySubLengths_E_B, + InBlockCopyClusterLengths_E_B, + InBlockCopyThreadClusterArrangeOrder, + InBlockCopySrcAccessOrder, + InBlockCopyDstAccessOrder, + InBlockCopyDataPerAccess_B, + WeiBlockCopySubLengths_E_K, + WeiBlockCopyClusterLengths_E_K, + WeiBlockCopyThreadClusterArrangeOrder, + WeiBlockCopySrcAccessOrder, + WeiBlockCopyDstAccessOrder, + WeiBlockCopySrcDataPerRead_E, + WeiBlockCopyDstDataPerWrite_K, + OutThreadCopyDataPerAccess_B>{}; + + for(index_t i = 0; i < nrepeat; ++i) + { + float time = launch_kernel(run_gridwise_convolution_kernel, + dim3(GridSize), + dim3(BlockSize), + 0, + static_cast(in_nchw_device_buf.GetDeviceBuffer()), + static_cast(wei_kcyx_device_buf.GetDeviceBuffer()), + static_cast(out_nkhw_device_buf.GetDeviceBuffer())); + + printf("Elapsed time : %f ms, %f TFlop/s\n", + time, + (float)calculate_convolution_flops(InDesc{}, WeiDesc{}, OutDesc{}) / + (std::size_t(1000) * 1000 * 1000) / time); + usleep(std::min(time * 1000, float(10000))); + } + + out_nkhw_device_buf.FromDevice(out_nkhw.mData.data()); +} diff --git a/driver/src/driver.cpp b/driver/src/driver.cpp index c792ead481..cd8325f0f8 100644 --- a/driver/src/driver.cpp +++ b/driver/src/driver.cpp @@ -19,6 +19,7 @@ //#include "device_convolution_implicit_gemm_v4r2_nchw_kcyx_nkhw.hpp" //#include "device_convolution_implicit_gemm_v4r3_nchw_kcyx_nkhw.hpp" #include "device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp" +#include "device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded.hpp" struct GeneratorTensor_1 { diff --git a/driver/src/driver.cu b/driver/src/driver.cu deleted file mode 100644 index 65f9afdf1f..0000000000 --- a/driver/src/driver.cu +++ /dev/null @@ -1,486 +0,0 @@ -#include -#include -#include -#include -#include -#include "config.hpp" -#include "ConstantTensorDescriptor.hpp" -#include "device.hpp" -#include "conv_common.hpp" -#include "host_conv.hpp" -#include "device_convolution_direct_v2_nchw_kcyx_nkhw.hpp" -#include "device_convolution_implicit_gemm_v1_chwn_cyxk_khwn.hpp" -#include "device_convolution_implicit_gemm_v1_chwn_cyxk_khwn_padded.hpp" -//#include "device_convolution_implicit_gemm_v1_nchw_cyxk_nkhw.hpp" -//#include "device_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp" -//#include "device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp" -#include "device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp" -#include "device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_padded.hpp" -//#include "device_convolution_implicit_gemm_v4r2_nchw_kcyx_nkhw.hpp" -//#include "device_convolution_implicit_gemm_v4r3_nchw_kcyx_nkhw.hpp" -#include "device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp" - -struct GeneratorTensor_1 -{ - template - double operator()(Is... is) - { - return 1; - } -}; - -struct GeneratorTensor_2 -{ - int min_value = 0; - int max_value = 1; - - template - double operator()(Is...) - { - return (std::rand() % (max_value - min_value)) + min_value; - } -}; - -struct GeneratorTensor_3 -{ - template - double operator()(Is... is) - { - std::array dims = {{static_cast(is)...}}; - - auto f_acc = [](auto a, auto b) { return 100 * a + b; }; - - return std::accumulate(dims.begin(), dims.end(), index_t(0), f_acc); - } -}; - -struct GeneratorTensor_Checkboard -{ - template - double operator()(Ts... Xs) const - { - std::array dims = {{Xs...}}; - return std::accumulate(dims.begin(), - dims.end(), - true, - [](bool init, index_t x) -> int { return init != (x % 2); }) - ? 1 - : -1; - } -}; - -int main(int argc, char* argv[]) -{ - using namespace ck; - -#if 0 - constexpr index_t N = 32; - constexpr index_t C = 8; - constexpr index_t HI = 1; - constexpr index_t WI = 1; - constexpr index_t K = 128; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - using LeftPads = Sequence<1, 1>; - using RightPads = Sequence<0, 0>; -#elif 1 - // 3x3, 34x34 - constexpr index_t N = 64; - constexpr index_t C = 256; - constexpr index_t HI = 34; - constexpr index_t WI = 34; - constexpr index_t K = 128; - constexpr index_t Y = 3; - constexpr index_t X = 3; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - using LeftPads = Sequence<0, 0>; - using RightPads = Sequence<0, 0>; -#elif 0 - // 1x1 filter, 8x8 image - // cudnn@V100 68%, ck@V100 72%, ck@P100 52%, ck@VII 42% - constexpr index_t N = 64; - constexpr index_t C = 1536; - constexpr index_t HI = 8; - constexpr index_t WI = 8; - constexpr index_t K = 256; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 8x8 image - // cudnn@V100 77%, ck@V100 76%, ck@P100 79%, ck@VII 51% - constexpr index_t N = 128; - constexpr index_t C = 2048; - constexpr index_t HI = 8; - constexpr index_t WI = 8; - constexpr index_t K = 384; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 7x7 image - // cudnn@V100 82%, ck@V100 76%, ck@P100 67%, ck@VII 64% - constexpr index_t N = 128; - constexpr index_t C = 832; - constexpr index_t HI = 7; - constexpr index_t WI = 7; - constexpr index_t K = 384; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 8x8 image - // cudnn@V100 83%, ck@V100 75%, ck@P100 78%, ck@VII 65% - constexpr index_t N = 128; - constexpr index_t C = 1280; - constexpr index_t HI = 8; - constexpr index_t WI = 8; - constexpr index_t K = 384; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 14x14 image - // cudnn@V100 62%, ck@V100 68%, ck@P100 70%, ck@VII 50% - constexpr index_t N = 128; - constexpr index_t C = 512; - constexpr index_t HI = 14; - constexpr index_t WI = 14; - constexpr index_t K = 128; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 8x8 image - // cudnn@V100 74%, ck@V100 57%, ck@P100 78%, ck@VII 61% - constexpr index_t N = 64; - constexpr index_t C = 1536; - constexpr index_t HI = 8; - constexpr index_t WI = 8; - constexpr index_t K = 384; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 28x28 image - // cudnn@V100 86%, ck@V100 84%, ck@P100 80%, ck@VII 69% - constexpr index_t N = 128; - constexpr index_t C = 256; - constexpr index_t HI = 28; - constexpr index_t WI = 28; - constexpr index_t K = 128; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 7x7 image - // cudnn@V100 71%, ck@V100 55%, ck@P100 70%, ck@VII 62% - constexpr index_t N = 128; - constexpr index_t C = 832; - constexpr index_t HI = 7; - constexpr index_t WI = 7; - constexpr index_t K = 256; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 3x3 filter, 2x2 stride, 35x35 input, 17x17 output - // cudnn@V100 90%, ck@V100 93%, ck@P100 83%, ck@VII 81% - constexpr index_t N = 128; - constexpr index_t C = 288; - constexpr index_t HI = 35; - constexpr index_t WI = 35; - constexpr index_t K = 384; - constexpr index_t Y = 3; - constexpr index_t X = 3; - - using ConvStrides = Sequence<2, 2>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 1 - // 1x1 filter, 17x17 input - // cudnn@V100 81%, ck@V100 76%, ck@P100 70%, ck@VII 76% - constexpr index_t N = 128; - constexpr index_t C = 768; - constexpr index_t HI = 17; - constexpr index_t WI = 17; - constexpr index_t K = 128; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 14x14 image - // cudnn@V100 73%, ck@V100 71%, ck@P100 70%, ck@VII 64% - constexpr index_t N = 128; - constexpr index_t C = 528; - constexpr index_t HI = 14; - constexpr index_t WI = 14; - constexpr index_t K = 128; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 14x14 image - // cudnn@V100 73%, ck@V100 72%, ck@P100 79%, ck@VII 75% - constexpr index_t N = 128; - constexpr index_t C = 528; - constexpr index_t HI = 14; - constexpr index_t WI = 14; - constexpr index_t K = 256; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#elif 0 - // 1x1 filter, 7x7 image - // cudnn@V100 49%, ck@V100 50%, ck@P100 61%, ck@VII 52% - constexpr index_t N = 128; - constexpr index_t C = 832; - constexpr index_t HI = 7; - constexpr index_t WI = 7; - constexpr index_t K = 128; - constexpr index_t Y = 1; - constexpr index_t X = 1; - - using ConvStrides = Sequence<1, 1>; - using ConvDilations = Sequence<1, 1>; - - constexpr index_t HPad = 0; - constexpr index_t WPad = 0; -#endif - - auto in_nchw_desc = make_ConstantTensorDescriptor_packed(Sequence{}); - auto wei_kcyx_desc = make_ConstantTensorDescriptor_packed(Sequence{}); - auto out_nkhw_desc = get_convolution_with_padding_output_default_4d_tensor_descriptor( - in_nchw_desc, wei_kcyx_desc, ConvStrides{}, ConvDilations{}, LeftPads{}, RightPads{}); - - ostream_ConstantTensorDescriptor(in_nchw_desc, std::cout << "in_nchw_desc: "); - ostream_ConstantTensorDescriptor(wei_kcyx_desc, std::cout << "wei_kcyx_desc: "); - ostream_ConstantTensorDescriptor(out_nkhw_desc, std::cout << "out_nkhw_desc: "); - - using in_data_t = float; - using out_data_t = float; - Tensor in_nchw(make_TensorDescriptor(in_nchw_desc)); - Tensor wei_kcyx(make_TensorDescriptor(wei_kcyx_desc)); - Tensor out_nkhw_host(make_TensorDescriptor(out_nkhw_desc)); - Tensor out_nkhw_device(make_TensorDescriptor(out_nkhw_desc)); - - std::size_t num_thread = std::thread::hardware_concurrency(); - - if(argc != 3) - { - printf("arg1: do_verification, arg2: nrepeat\n"); - exit(1); - } - - bool do_verification = atoi(argv[1]); - index_t nrepeat = atoi(argv[2]); - - if(do_verification) - { -#if 0 - in_nchw.GenerateTensorValue(GeneratorTensor_1{}, num_thread); - wei_kcyx.GenerateTensorValue(GeneratorTensor_1{}, num_thread); -#elif 0 - in_nchw.GenerateTensorValue(GeneratorTensor_1{}, num_thread); - wei_kcyx.GenerateTensorValue(GeneratorTensor_2{-5, 5}, num_thread); -#elif 0 - in_nchw.GenerateTensorValue(GeneratorTensor_3{}, num_thread); - wei_kcyx.GenerateTensorValue(GeneratorTensor_1{}, num_thread); -#elif 1 - in_nchw.GenerateTensorValue(GeneratorTensor_2{-5, 5}, num_thread); - wei_kcyx.GenerateTensorValue(GeneratorTensor_2{-5, 5}, num_thread); -#elif 0 - in_nchw.GenerateTensorValue(GeneratorTensor_2{1, 5}, num_thread); - - auto gen_wei = [](auto... is) { - return GeneratorTensor_2{1, 5}(is...) * GeneratorTensor_Checkboard{}(is...); - }; - wei_kcyx.GenerateTensorValue(gen_wei, num_thread); -#endif - } - -#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_chwn_cyxk_khwn_padded(in_nchw_desc, - in_nchw, - wei_kcyx_desc, - wei_kcyx, - out_nkhw_desc, - out_nkhw_device, - LeftPads{}, - RightPads{}, - 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( - (in_nchw_desc, in_nchw, wei_kcyx_desc, wei_kcyx, out_nkhw_desc, out_nkhw_device, nrepeat); -#elif 0 - device_convolution_implicit_gemm_v4r1_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_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_padded(in_nchw_desc, - in_nchw, - wei_kcyx_desc, - wei_kcyx, - out_nkhw_desc, - out_nkhw_device, - ConvStrides{}, - ConvDilations{}, - LeftPads{}, - RightPads{}, - nrepeat); -#elif 0 - device_convolution_implicit_gemm_v4r2_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_convolution_implicit_gemm_v4r3_nchw_kcyx_nkhw(in_nchw_desc, - in_nchw, - wei_kcyx_desc, - wei_kcyx, - out_nkhw_desc, - out_nkhw_device, - ConvStrides{}, - ConvDilations{}, - nrepeat); -#elif 1 - device_convolution_implicit_gemm_v4r4_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_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw_padded(in_nchw_desc, - in_nchw, - wei_kcyx_desc, - wei_kcyx, - out_nkhw_desc, - out_nkhw_device, - ConvStrides{}, - ConvDilations{}, - LeftPads{}, - RightPads{}, - nrepeat); -#endif - - if(do_verification) - { -#if 1 - if(Y == 3 && X == 3 && ConvStrides{}[0] == 1 && ConvStrides{}[1] == 1 && - ConvDilations{}[0] == 1 && ConvDilations{}[1] == 1) - { - host_winograd_3x3_convolution( - in_nchw, wei_kcyx, out_nkhw_host, LeftPads{}, RightPads{}); - } - else -#endif - { - host_direct_convolution(in_nchw, - wei_kcyx, - out_nkhw_host, - ConvStrides{}, - ConvDilations{}, - LeftPads{}, - RightPads{}); - } - check_error(out_nkhw_host, out_nkhw_device); - -#if 0 - LogRange(std::cout << "in_nchw : ", in_nchw.mData, ",") << std::endl; - LogRange(std::cout << "wei_kcyx: ", wei_kcyx.mData, ",") << std::endl; - LogRange(std::cout << "out_nkhw_host : ", out_nkhw_host.mData, ",") << std::endl; - LogRange(std::cout << "out_nkhw_device: ", out_nkhw_device.mData, ",") << std::endl; -#endif - } -} diff --git a/driver/src/driver.cu b/driver/src/driver.cu new file mode 120000 index 0000000000..1ca4fea9d7 --- /dev/null +++ b/driver/src/driver.cu @@ -0,0 +1 @@ +driver.cpp \ No newline at end of file