From b27dddfb59198a8f6768c091e0d22598a41361ab Mon Sep 17 00:00:00 2001 From: Qun Lin Date: Mon, 2 Jun 2025 10:39:13 +0800 Subject: [PATCH] bwd weight dl v4 - v0 --- .../20_grouped_conv_bwd_weight/CMakeLists.txt | 4 + .../device_grouped_conv_bwd_weight_dl_v4.hpp | 1120 +++++++++++------ .../grouped_conv_bwd_weight_dl_v4_fp16.cpp | 12 +- 3 files changed, 728 insertions(+), 408 deletions(-) diff --git a/example/20_grouped_conv_bwd_weight/CMakeLists.txt b/example/20_grouped_conv_bwd_weight/CMakeLists.txt index de54c4a4b9..6109dd0e07 100644 --- a/example/20_grouped_conv_bwd_weight/CMakeLists.txt +++ b/example/20_grouped_conv_bwd_weight/CMakeLists.txt @@ -20,3 +20,7 @@ add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bw add_example_executable(example_grouped_conv_bwd_weight_v3_xdl_fp16 grouped_conv_bwd_weight_v3_xdl_fp16.cpp) add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bwd_weight_v3_xdl_fp16) + + +add_example_executable(example_grouped_conv_bwd_weight_dl_v4_fp16 grouped_conv_bwd_weight_dl_v4_fp16.cpp) +target_compile_options(example_grouped_conv_bwd_weight_dl_v4_fp16 PRIVATE -save-temps=obj -Wno-gnu-line-marker) \ No newline at end of file diff --git a/example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp b/example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp index 1d34684a86..e1c1cc52bf 100644 --- a/example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp +++ b/example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp @@ -3,11 +3,10 @@ #include "common.hpp" -//#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp" +// #include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp" #include "ck/utility/blkgemmpipe_scheduler.hpp" #include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp" - using InDataType = F16; using WeiDataType = F16; using OutDataType = F16; @@ -17,42 +16,9 @@ using InElementOp = PassThrough; using WeiElementOp = PassThrough; using OutElementOp = PassThrough; +template struct Debug; + namespace ck { -namespace tensor_operation { -namespace device { -static constexpr index_t WaveSize = 64; -static constexpr index_t W_PACK = 2; // WaveSize / arg.input_spatial_lengths_[1]; -static constexpr index_t Tile_H = 32; -static constexpr index_t Tile_W = 32; -static constexpr index_t N_Pack = 2; -static constexpr index_t Pad_H = 2; -static constexpr index_t Pad_W = 2; -static constexpr index_t Filter_X = 5; -static constexpr index_t Filter_Y = 5; - -static constexpr index_t SizeOfType = 2; -static constexpr index_t Tile_Align_W = Tile_W; -static constexpr index_t ShareMemSize = Tile_H * Tile_Align_W * N_Pack * SizeOfType; -static constexpr index_t ScratchSize = ShareMemSize / 64 / 4; -static constexpr index_t Num_Wave = 4; -#define MergeShareMem 0 -template -__device__ T warp_shuffle_up(const T& v_local, uint32_t lane_delta) -{ -#if 0 - return __shfl_up(v_local, lane_delta); -#elif 1 - static_assert(sizeof(T) == sizeof(int32_t), "wrong!"); - - const uint32_t wrap_around_lane_delta = warpSize - lane_delta; - - const int32_t v_remote_tmp = __builtin_amdgcn_ds_bpermute( - (__lane_id() << 2) + (wrap_around_lane_delta << 2), bit_cast(v_local)); - - return bit_cast(v_remote_tmp); -#endif -} - template __device__ T warp_shuffle_down(const T& v_local, uint32_t lane_delta) { @@ -68,316 +34,618 @@ __device__ T warp_shuffle_down(const T& v_local, uint32_t lane_delta) #endif } -#pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wold-style-cast" -template -void __device__ load_data_from_global(DataType* p, - index_t lane_id, - index_t n_stride, - index_t h, - index_t w, - index_t h_stride, - index_t w_stride, - uint32_t* p_scratch) -{ - ignore = h; - ignore = w; - DataType* p_1 = p + n_stride; - static_assert(sizeof(DataType) == 2); - static_assert(Pad_H % W_PACK == 0); - - const index_t x = lane_id % (WaveSize/W_PACK); - const index_t y_base = lane_id / (WaveSize/W_PACK); - - auto get_offset = [&](index_t y_, index_t x_) - { - return y_ * h_stride + x_ * w_stride; - }; - - if(x >= Pad_W && x < w + Pad_W) - { - static_for<0, Tile_H / W_PACK, 1>{}([&](auto i) { - const index_t y = y_base + i * W_PACK; - if constexpr (i * W_PACK >= Pad_H && i * W_PACK < (Tile_H - Pad_H)) - { - const index_t offset = get_offset(y - Pad_H, x - Pad_W); - half2_t tmp = {}; - tmp[0] = p[offset]; - tmp[1] = p_1[offset]; - p_scratch[i] = bit_cast(tmp); - } - }); - } -} - -void __device__ write_data_to_lds(index_t lane_id, const uint32_t* p_scratch, uint32_t* p_sharemem) -{ - const index_t x = lane_id % (WaveSize/W_PACK); - const index_t y_base = lane_id / (WaveSize/W_PACK); - //static_assert(N_Pack * sizeof(InDataType)/ sizeof(uint32_t) == 1); - - auto get_offset = [&](index_t y_, index_t x_) { return (y_ * Tile_Align_W + (x_));}; // + y_ * Filter_X) % Tile_W); }; - static_for<0, Tile_H / W_PACK, 1>{}([&](auto i) { - const index_t y = y_base + i * W_PACK; - const index_t offset = get_offset(y, x); - p_sharemem[offset] = p_scratch[i]; - }); -} - -void __device__ run_conv_bwd_weight(index_t x, index_t y, index_t H, index_t W, index_t H_base, uint32_t* p_share_in, uint32_t* p_share_out,float& acc) -{ - ignore = H; - ignore = W; - auto get_in = [&](int h_, int w_) - { - return p_share_in[(h_ + y + H_base) * Tile_Align_W + (w_ + x)];// + (h_ + y + H_base) * Filter_X) % Tile_W]; - }; - auto get_out = [&](int h_, int w_) - { - return p_share_out[(h_ + Pad_H + H_base) * Tile_Align_W + (w_ + Pad_W)];// + (h_ + Pad_H + H_base) * Filter_X) % Tile_W]; - }; - if (x < Filter_X && y < Filter_Y) - { - //for (int ho = 0; ho < H; ho++) - static_for<0, (Tile_H - Pad_H - Pad_H) / 2, 1>{}([&](auto ho) - { - //for (int wo = 0; wo < W; wo++) - static_for<0, Tile_W - Pad_W - Pad_W, 1>{}([&](auto wo) - { - uint32_t v_in = get_in(ho, wo); - uint32_t v_out = get_out(ho, wo); - acc = __builtin_amdgcn_fdot2(bit_cast(v_in), bit_cast(v_out), acc, false); - }); - }); - } -} -template -void __device__ write_output(const Argument& arg, int g, int y, int x, WeiDataType acc) -{ - const index_t Wei_G_Stride = arg.wei_g_k_c_xs_strides_[0]; - const index_t Y_Stride = arg.wei_g_k_c_xs_strides_[3]; - const index_t X_Stride = arg.wei_g_k_c_xs_strides_[4]; - - if (y < Filter_Y && x < Filter_X) - { - auto p_wei = arg.p_wei_grid_ + Wei_G_Stride * g + y * Y_Stride + x * X_Stride; - *p_wei = acc; - } -} - - -template +template __global__ void #if CK_USE_LAUNCH_BOUNDS - __launch_bounds__(WaveSize * Num_Wave, MinimumOccupancy) +__launch_bounds__(BlockSize, MinimumOccupancy) #endif - kernel_grouped_conv_bwd_weight_dl_v4(Argument arg) + kernel_grouped_conv_bwd_weight_dl_v4(typename GridwiseConvBwdWeight::Argument arg) { -#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx9__)) - const index_t g_idx = __builtin_amdgcn_readfirstlane(blockIdx.x); - const index_t wave_id = __builtin_amdgcn_readfirstlane(threadIdx.x / WaveSize); - const index_t lane_id = __lane_id(); - //index_t n_idx = 0; +//#if(!defined(__HIP_DEVICE_COMPILE__)) + __shared__ char + p_share_in[GridwiseConvBwdWeight::ShareMemInSize * GridwiseConvBwdWeight::NumTilePerBlock]; + __shared__ char + p_share_out[GridwiseConvBwdWeight::ShareMemOutSize * GridwiseConvBwdWeight::NumTilePerBlock]; + + GridwiseConvBwdWeight::template Run(arg, p_share_in, p_share_out); +//#else +// ignore = arg; +//#endif +} -#if MergeShareMem - constexpr index_t ShaderMemSizePerWave =ShareMemSize * 2 - Pad_H * Tile_Align_W * sizeof(uint32_t); - __shared__ uint32_t p_share_mem[ShaderMemSizePerWave/sizeof(uint32_t) * Num_Wave]; - auto* p_share_in = &p_share_mem[0]; - //__shared__ char p_input_1[ShareMemSize]; - auto* p_share_out = &p_share_mem[(ShareMemSize - Pad_H * Tile_Align_W * sizeof(uint32_t))/ sizeof(uint32_t)]; -#else - __shared__ uint32_t p_share_in[ShareMemSize/sizeof(uint32_t) * Num_Wave]; - //__shared__ char p_input_1[ShareMemSize]; - __shared__ uint32_t p_share_out[ShareMemSize/sizeof(uint32_t) * Num_Wave]; - //__shared__ char p_output_1[ShareMemSize]; -#endif - uint32_t p_input_0_scratch[ScratchSize]; - uint32_t p_input_1_scratch[ScratchSize]; - uint32_t p_output_0_scratch[ScratchSize]; - uint32_t p_output_1_scratch[ScratchSize]; +namespace tensor_operation { +namespace device { - static constexpr index_t spatial_offset = 3; - //const index_t G = arg.in_g_n_c_wis_lengths[0]; - const index_t N = arg.in_g_n_c_wis_lengths_[1]; - index_t num_loop = N / Num_Wave / 2 - 1; - index_t n_idx = N / Num_Wave * wave_id; - if (wave_id == Num_Wave - 1) +template + typename InElementwiseOperation, + typename WeiElementwiseOperation, + typename OutElementwiseOperation, + index_t NBatch, + index_t NumWavePerTile, + index_t InScalarPerVector, + index_t OutScalarPerVector, + index_t DstScalarPerVector, + bool RequirePadding, + typename ComputeTypeA = InDataType, + typename ComputeTypeB = ComputeTypeA> +struct GridwiseGroupedConv2DBwdWeightDlV4 +{ + static constexpr index_t + GetConvOut(index_t length, index_t filter, index_t dilation, index_t pad, index_t stride) { - n_idx = N / Num_Wave * (Num_Wave - 1); - num_loop = (N - n_idx)/2 -1; + return (length + pad + pad - ((filter - 1) * dilation + 1)) / stride + 1; } - - // In - const index_t Hi = arg.in_g_n_c_wis_lengths_[spatial_offset + 0]; - const index_t Wi = arg.in_g_n_c_wis_lengths_[spatial_offset + 1]; - - const index_t Hi_Stride = arg.in_g_n_c_wis_strides_[spatial_offset + 0]; - const index_t Wi_Stride = arg.in_g_n_c_wis_strides_[spatial_offset + 1]; - const index_t In_G_Stride = arg.in_g_n_c_wis_strides_[0]; - const index_t In_N_Stride = arg.in_g_n_c_wis_strides_[1]; - - // Out - const index_t Ho = arg.out_g_n_k_wos_lengths_[spatial_offset + 0]; - const index_t Wo = arg.out_g_n_k_wos_lengths_[spatial_offset + 1]; - - const index_t Ho_Stride = arg.out_g_n_k_wos_strides_[spatial_offset + 0]; - const index_t Wo_Stride = arg.out_g_n_k_wos_strides_[spatial_offset + 1]; - const index_t Out_G_Stride = arg.out_g_n_k_wos_strides_[0]; - const index_t Out_N_Stride = arg.out_g_n_k_wos_strides_[1]; - - static_for<0,ScratchSize, 1>{}([&](auto i) + template + static constexpr index_t GetAlignedPackW() { - p_input_0_scratch[i] = 0; - p_output_0_scratch[i] = 0; - p_input_1_scratch[i] = 0; - p_output_1_scratch[i] = 0; - }); - // Wei - - //static_assert(sizeof(InDataType) == 2); - //static_assert(sizeof(OutputDataType) == 2); - // - auto* p_in = arg.p_in_grid_ + g_idx * In_G_Stride + n_idx * In_N_Stride; - auto* p_out = arg.p_out_grid_ + g_idx * Out_G_Stride + n_idx * Out_N_Stride; - - // prefetch 0 - load_data_from_global(p_in, lane_id, In_N_Stride, Hi, Wi, Hi_Stride, Wi_Stride, p_input_0_scratch); - load_data_from_global(p_out, lane_id, Out_N_Stride, Ho, Wo, Ho_Stride, Wo_Stride, p_output_0_scratch); - //load_data_from_global(p_in + In_N_Stride, Hi, Wi, Hi_Stride, Wi_Stride, p_input_1_scratch); - //load_data_from_global(p_out + Out_N_Stride, Ho, Wo, Ho_Stride, Wo_Stride, p_output_1_scratch); - p_in += 2 * In_N_Stride; - p_out += 2 * Out_N_Stride; - - // prefetch 1 - //load_input_from_global(arg, p_in, n_idx + 1, p_input_1_scratch); - //load_output_from_global(arg, p_out, n_idx + 1, p_output_0_scratch); - - // write 0 - auto p_input_0 = p_share_in + ShareMemSize/sizeof(uint32_t) * wave_id; - auto p_output_0 = p_share_out + ShareMemSize/sizeof(uint32_t) * wave_id; - - write_data_to_lds( lane_id,p_input_0_scratch, p_input_0); - write_data_to_lds( lane_id,p_output_0_scratch, p_output_0); - index_t H_base = lane_id >=32 ? (Tile_H - Pad_H - Pad_H) /2 : 0; - index_t x = (lane_id % 32) % Filter_X; - index_t y = (lane_id % 32) / Filter_X; - float acc = 0; - - - while(num_loop > 0) - { - // prefetch 0 - load_data_from_global(p_in, lane_id,In_N_Stride, Hi, Wi, Hi_Stride, Wi_Stride, p_input_0_scratch); - load_data_from_global(p_out, lane_id, Out_N_Stride, Ho, Wo, Ho_Stride, Wo_Stride, p_output_0_scratch); - //load_data_from_global(p_in + In_N_Stride, Hi, Wi, Hi_Stride, Wi_Stride, p_input_1_scratch); - //load_data_from_global(p_out + Out_N_Stride, Ho, Wo, Ho_Stride, Wo_Stride, p_output_1_scratch); - p_in += 2 * In_N_Stride; - p_out += 2 * Out_N_Stride; - - // do conv_bwd on 0 - block_sync_lds(); - run_conv_bwd_weight(x, y, Ho, Wo, H_base, p_input_0, p_output_0, acc); - block_sync_lds(); - // write 1 - //write_input_to_lds(); - //write_output_to_lds(); - - // prefetch 1 - //load_input_from_global(); - //load_output_from_global(); - // do conv_bwd on 1 - //run_conv_bwd_weight(); - - // write 0 - write_data_to_lds( lane_id,p_input_0_scratch, p_input_0); - write_data_to_lds( lane_id,p_output_0_scratch, p_output_0); - - num_loop --; - }; - - // tail - { - run_conv_bwd_weight(x, y, Ho, Wo, H_base, p_input_0, p_output_0, acc); - } - float acc_2 = warp_shuffle_down(acc, 32); - acc += acc_2; - block_sync_lds(); - p_share_in[threadIdx.x] = bit_cast(acc); - block_sync_lds(); - if (H_base == 0 && wave_id == 0) - { - for(int i= 1; i < Num_Wave; i++) + constexpr index_t pakced_w = W / ScalarPerVector; + if constexpr(pakced_w == 7) { - acc += bit_cast(p_share_in[i * WaveSize + lane_id]); + return 8; + } + else + { + return WaveSize / (WaveSize / pakced_w); + } + } + static constexpr index_t GetBatchPerWave() { return WaveSize / (FilterSize * FilterSize); } + + static constexpr index_t NDimSpatial = 2; + static constexpr index_t NumVectorPerPixel = NBatch / DstScalarPerVector; + static constexpr auto I0 = Number<0>{}; + static constexpr auto I1 = Number<1>{}; + + static constexpr index_t WaveSize = 64; + static constexpr index_t Tile_H = BlockTileSize{}.At(I0); + static constexpr index_t Tile_W = BlockTileSize{}.At(I1); + + // Only support pad left = pad right for now + static constexpr index_t Pad_H = tuple_element_t<2, FilterParam>{}.At(I0); + static constexpr index_t Pad_W = tuple_element_t<2, FilterParam>{}.At(I1); + + static constexpr index_t Stride_H = tuple_element_t<1, FilterParam>{}.At(I0); + static constexpr index_t Stride_W = tuple_element_t<1, FilterParam>{}.At(I1); + + // Only support dilation = 1 for now + static constexpr index_t Dilation_Y = tuple_element_t<0, FilterParam>{}.At(I0); + static constexpr index_t Dilation_X = tuple_element_t<0, FilterParam>{}.At(I1); + + static constexpr index_t Filter_Y = FilterSize; + static constexpr index_t Filter_X = FilterSize; + + static constexpr index_t TileIn_H = Tile_H + 2 * Pad_H; + static constexpr index_t TileIn_W = Tile_H + 2 * Pad_W; + + static constexpr index_t TileOut_H = + GetConvOut(Tile_H, FilterSize, Dilation_Y, Pad_H, Stride_H); + static constexpr index_t TileOut_W = + GetConvOut(Tile_W, FilterSize, Dilation_X, Pad_W, Stride_W); + + static constexpr index_t TileIn_Align_W = TileIn_W; + static constexpr index_t TileOut_Align_W = TileOut_W; + + static constexpr index_t TileIn_Pack_W = GetAlignedPackW(); + static constexpr index_t TileIn_Pack_Group = WaveSize / TileIn_Pack_W; + static constexpr index_t TileIn_Pack_H = math::integer_divide_ceil(Tile_H, TileIn_Pack_Group); + static constexpr index_t TileIn_Align_H = + math::max(TileIn_Pack_H * TileIn_Pack_Group + Pad_H, TileIn_H); + + static constexpr index_t TileOut_Pack_W = GetAlignedPackW; + static constexpr index_t TileOut_Pack_Group = WaveSize / TileIn_Pack_W; + static constexpr index_t TileOut_Pack_H = math::integer_divide_ceil(TileOut_H, TileOut_Pack_Group); + + static constexpr index_t BatchPerWave = GetBatchPerWave(); + static constexpr index_t BatchPerTile = BatchPerWave * NumWavePerTile; + static constexpr index_t TileOut_HPerBatch = math::integer_divide_ceil(TileOut_H, BatchPerTile); + static constexpr index_t TileOut_Align_H = + math::max(TileOut_Pack_H * TileOut_Pack_Group, TileOut_HPerBatch* BatchPerWave* NumWavePerTile); + static constexpr index_t ShareMemInSize = + TileIn_Align_H * TileIn_Align_W * sizeof(InDataType) * NBatch; + static constexpr index_t ShareMemOutSize = + TileOut_Align_H * TileOut_Align_W * sizeof(OutDataType) * NBatch; + + static constexpr index_t NumTilePerBlock = BlockSize / WaveSize; + + using InDataVector = typename vector_type::type; + using OutDataVector = typename vector_type::type; + + template + static void __device__ load_data_from_global(const SrcType* p, + index_t lane_id, + index_t n_stride, + index_t h, + index_t w, + index_t h_stride, + index_t w_stride, + DestVector* p_scratch) + { + ignore = h; + ignore = w; + + using SrcVector = typename vector_type::type; + constexpr index_t PackW = TileW / ScalarPerVector; + constexpr index_t AlignedPackW = GetAlignedPackW(); + static_assert(PackW < WaveSize); + + static_assert((AlignedPackW & (AlignedPackW - 1)) == 0, "aligned width is not power 2!"); + + constexpr index_t NumGroup = WaveSize / AlignedPackW; + constexpr index_t AlignedPackH = math::integer_divide_ceil(TileH, NumGroup); + constexpr index_t PackH = TileH / NumGroup; + + const index_t x = lane_id % (WaveSize / NumGroup); + const index_t y_offset = lane_id / (WaveSize / NumGroup); + + auto get_offset = [&](index_t y_, index_t x_, index_t n_) { + return (y_ * h_stride + x_ * w_stride + n_ * n_stride) / ScalarPerVector; + }; + + // todo: check with real width/height + // and use OOB to avoid tynamic control flow. + auto* p_base = reinterpret_cast(p); + if(x < PackW) + { + static_for<0, PackH, 1>{}([&](auto i) { + const index_t y = y_offset + i * NumGroup; + // load data + SrcVector tmp[NBatch]; + static_for<0, NBatch, 1>{}([&](auto n) { + const index_t offset = get_offset(y, x, n); + tmp[n] = p_base[offset]; + }); + + // interleave data + auto* p_scratch_base = p_scratch + i * NumVectorPerPixel * ScalarPerVector; + static_for<0, ScalarPerVector, 1>{}([&](auto j) { + static_for<0, NBatch, 1>{}([&](auto n) { + p_scratch_base[j * NumVectorPerPixel + n / DstScalarPerVector] + [n / DstScalarPerVector] = tmp[n][j.value]; + }); + }); + }); + + if constexpr(AlignedPackH != PackH) + { + static_assert(NumWavePerTile == 1); + if (y_offset < (TileH - NumGroup * PackH)) + { + constexpr auto i = PackH; + const index_t y = y_offset + i * NumGroup; + // load data + SrcVector tmp[NBatch]; + static_for<0, NBatch, 1>{}([&](auto n) { + const index_t offset = get_offset(y, x, n); + tmp[n] = p_base[offset]; + }); + + // interleave data + auto* p_scratch_base = p_scratch + i * NumVectorPerPixel * ScalarPerVector; + static_for<0, ScalarPerVector, 1>{}([&](auto j) { + static_for<0, NBatch, 1>{}([&](auto n) { + p_scratch_base[j * NumVectorPerPixel + n / DstScalarPerVector] + [n / DstScalarPerVector] = tmp[n][j.value]; + }); + }); + } + } } - write_output(arg, g_idx, y, x, acc); } -#else - ignore = karg; -#endif // end of if (defined(__gfx9__)) -} -#pragma clang diagnostic pop + // todo handle pading in p_sharemem + template + static void __device__ write_data_to_lds(index_t lane_id, + const DestVector* p_scratch, + DestVector* p_sharemem) + { + //constexpr index_t PackW = TileW / ScalarPerVector; + constexpr index_t AlignedPackW = GetAlignedPackW(); + static_assert(AlignedPackW < WaveSize); -template {}([&](auto i) { + const index_t y = y_offset + i * AlignedPackW; + const index_t offset = get_offset(y, x); + static_for<0, NumVectorPerPixel * ScalarPerVector, 1>{}([&](auto j) { + p_sharemem[offset + j] = p_scratch[i * NumVectorPerPixel * ScalarPerVector + j]; + }); + }); + if constexpr(AlignedPackH != PackH) + { + static_assert(NumWavePerTile == 1); + if (y_offset < (TileH - NumGroup * PackH)) + { + constexpr auto i = PackH; + const index_t y = y_offset + i * AlignedPackW; + const index_t offset = get_offset(y, x); + static_for<0, NumVectorPerPixel * ScalarPerVector, 1>{}([&](auto j) { + p_sharemem[offset + j] = p_scratch[i * NumVectorPerPixel * ScalarPerVector + j]; + }); + } + } + } + } + + template + static void __device__ run_conv_bwd_weight(index_t x, + index_t y, + index_t h, + index_t w, + index_t hout_base, + InDataVector* p_share_in, + OutDataVector* p_share_out, + AccDataType& acc) + { + ignore = h; + ignore = w; + auto get_in = [&](index_t ho_, index_t wo_, index_t i_) { + // p_share_in has been shift with left_pad + index_t hi = (ho_ + hout_base) * Stride_H + y * Dilation_Y; + index_t wi = wo_ * Stride_W + x * Dilation_X; + return p_share_in[(hi * TileIn_Align_W + wi) * NumVectorPerPixel + i_]; + }; + auto get_out = [&](index_t ho_, index_t wo_, index_t i_) { + return p_share_out[((ho_ + hout_base) * TileOut_Align_W + wo_) * NumVectorPerPixel + i_]; + }; + if(x < Filter_X && y < Filter_Y) + { + static_for<0, TileH, 1>{}([&](auto ho) { + static_for<0, TileOut_W, 1>{}([&](auto wo) { + static_for<0, NumVectorPerPixel, 1>{}([&](auto i) { + auto v_in = get_in(ho, wo, i); + auto v_out = get_out(ho, wo, i); + inner_product(v_in, v_out, acc); + }); + }); + }); + } + } + template + static void __device__ write_output(const Argument& arg, index_t g, index_t y, index_t x, WeiDataType acc) + { + const index_t Wei_G_Stride = arg.wei_g_k_c_xs_strides_[0]; + const index_t Y_Stride = arg.wei_g_k_c_xs_strides_[3]; + const index_t X_Stride = arg.wei_g_k_c_xs_strides_[4]; + + if(y < Filter_Y && x < Filter_X) + { + auto p_wei = arg.p_wei_grid_ + Wei_G_Stride * g + y * Y_Stride + x * X_Stride; + *p_wei = acc; + } + } + + + template + static void __device__ Run(Argument arg, char* p_share_in, char* p_share_out) + { + const index_t g_idx = __builtin_amdgcn_readfirstlane(blockIdx.x); + const index_t wave_id = __builtin_amdgcn_readfirstlane(threadIdx.x / WaveSize); + const index_t lane_id = __lane_id(); + + constexpr index_t ThreadPerBatch = WaveSize / BatchPerWave; + + static_assert(TileIn_Pack_H % NumWavePerTile == 0); + static_assert(TileOut_Pack_H % NumWavePerTile == 0); + InDataVector tmp_in[TileIn_Pack_H / NumWavePerTile * NumVectorPerPixel * InScalarPerVector] = {}; + OutDataVector tmp_out[TileOut_Pack_H / NumWavePerTile * NumVectorPerPixel * OutScalarPerVector] = {}; + + static_assert(NumTilePerBlock == 1 || NumWavePerTile == 1); + + static constexpr index_t spatial_offset = 3; + const index_t n = arg.in_g_n_c_wis_lengths_[1]; + index_t num_loop = n / NumTilePerBlock / BatchPerWave - 1; + index_t n_idx = n / NumTilePerBlock * wave_id; + + if(wave_id == NumTilePerBlock - 1) + { + n_idx = n / NumTilePerBlock * (NumTilePerBlock - 1); + num_loop = (n - n_idx) / BatchPerWave - 1; + } + + // In + const index_t hi = arg.in_g_n_c_wis_lengths_[spatial_offset + 0]; + const index_t wi = arg.in_g_n_c_wis_lengths_[spatial_offset + 1]; + + const index_t hi_stride = arg.in_g_n_c_wis_strides_[spatial_offset + 0]; + const index_t wi_stride = arg.in_g_n_c_wis_strides_[spatial_offset + 1]; + const index_t in_g_stride = arg.in_g_n_c_wis_strides_[0]; + const index_t in_n_stride = arg.in_g_n_c_wis_strides_[1]; + + // Out + const index_t ho = arg.out_g_n_k_wos_lengths_[spatial_offset + 0]; + const index_t wo = arg.out_g_n_k_wos_lengths_[spatial_offset + 1]; + + const index_t ho_stride = arg.out_g_n_k_wos_strides_[spatial_offset + 0]; + const index_t wo_stride = arg.out_g_n_k_wos_strides_[spatial_offset + 1]; + const index_t out_g_stride = arg.out_g_n_k_wos_strides_[0]; + const index_t out_n_stride = arg.out_g_n_k_wos_strides_[1]; + + // Wei + auto* p_in = arg.p_in_grid_ + g_idx * in_g_stride + n_idx * in_n_stride; + auto* p_out = arg.p_out_grid_ + g_idx * out_g_stride + n_idx * out_n_stride; + + constexpr index_t Copy_TileIn_H = TileIn_H / NumWavePerTile; + constexpr index_t Copy_TileOut_H = TileOut_H / NumWavePerTile; + if constexpr (NumWavePerTile > 1) + { + static_assert(RequirePadding == false); + p_in += Copy_TileIn_H * hi_stride * wave_id; + p_out += Copy_TileOut_H * ho_stride * wave_id; + } + + InDataVector* share_in = + reinterpret_cast(p_share_in); + OutDataVector* share_out = + reinterpret_cast(p_share_out); + if constexpr(NumTilePerBlock > 1) + { + share_in = reinterpret_cast(p_share_in + ShareMemInSize * wave_id); + share_out = reinterpret_cast(p_share_out + ShareMemOutSize * wave_id); + } + + // init lds 0 + index_t cluster_id = threadIdx.x % (WaveSize * NumWavePerTile); + auto init_pading = [&](auto* share_vec, auto count) { + static_for<0, math::integer_divide_ceil(count, WaveSize * NumWavePerTile), 1>{}([&](auto i) { + if(cluster_id + i * WaveSize * NumWavePerTile < count) + { + share_vec[cluster_id + i * WaveSize * NumWavePerTile] = {}; + } + }); + }; + auto init_array_pading = [&](auto* share_vec, auto element_count, auto array_count, index_t stride) { + static_for<0, math::integer_divide_ceil(array_count, WaveSize * NumWavePerTile), 1>{}([&](auto i) { + static_for<0, element_count, 1>{}([&](auto j) { + if(cluster_id + i * WaveSize * NumWavePerTile < array_count) + { + auto p = share_vec + i * stride + j; + *p = {}; + } + }); + }); + }; + constexpr index_t TopPadingSize = Pad_H * TileIn_Align_W * NumVectorPerPixel; + constexpr index_t TileInEnd = (Tile_H + Pad_H) * TileIn_Align_W; + constexpr index_t ButtomPaddingSize = (ShareMemInSize / (sizeof(InDataType) * NBatch) - TileInEnd) * NumVectorPerPixel; + if constexpr(Pad_W > 0) + { + static_assert(ButtomPaddingSize >= 0); + init_array_pading(share_in + TopPadingSize, Number{}, Number{}, TileIn_Align_W * NumVectorPerPixel); + init_array_pading(share_in + TopPadingSize + Tile_W * NumVectorPerPixel , Number{}, Number{}, TileIn_Align_W * NumVectorPerPixel); + } + if constexpr(Pad_H > 0) + { + init_pading(share_in, Number{}); + init_pading(share_in + TileInEnd * NumVectorPerPixel, Number{}); + } + constexpr index_t TileOutEnd = TileOut_H * TileOut_Align_W; + constexpr index_t OutButtomPaddingSize = (ShareMemOutSize / (sizeof(OutDataType) * NBatch) - TileOutEnd) * NumVectorPerPixel; + init_pading(share_out + TileOutEnd * NumVectorPerPixel, Number{}); + + if constexpr (NumWavePerTile > 1) + { + block_sync_lds(); + } + + // prefetch 0 + load_data_from_global( + p_in, lane_id, in_n_stride, hi, wi, hi_stride, wi_stride, tmp_in); + load_data_from_global( + p_out, lane_id, out_n_stride, ho, wo, ho_stride, wo_stride, tmp_out); + p_in += NBatch * in_n_stride; + p_out += NBatch * out_n_stride; + + // adjust share memory offset + share_in += (TileIn_Align_W * Pad_H + Pad_W) * NumVectorPerPixel; + if constexpr (NumWavePerTile > 1) + { + static_assert(RequirePadding == false); + share_in += Copy_TileIn_H * TileIn_Align_W * NumVectorPerPixel * wave_id; + share_out += Copy_TileOut_H * TileOut_Align_W * NumVectorPerPixel * wave_id; + } + + write_data_to_lds( + lane_id, tmp_in, share_in); + write_data_to_lds( + lane_id, tmp_out, share_out); + + constexpr index_t TileOut_H_batch = math::integer_divide_ceil(Copy_TileOut_H, ThreadPerBatch); + index_t hout_base = lane_id / ThreadPerBatch * TileOut_H_batch; + if constexpr (NumWavePerTile > 1) + { + hout_base += Copy_TileOut_H * wave_id; + } + index_t x = (lane_id % ThreadPerBatch) % Filter_X; + index_t y = (lane_id % ThreadPerBatch) / Filter_X; + float acc = 0; + + while(num_loop > 0) + { + + load_data_from_global( + p_in, lane_id, in_n_stride, hi, wi, hi_stride, wi_stride, tmp_in); + load_data_from_global( + p_out, lane_id, out_n_stride, ho, wo, ho_stride, wo_stride, tmp_out); + p_in += NBatch * in_n_stride; + p_out += NBatch * out_n_stride; + + // do conv_bwd on 0 + run_conv_bwd_weight( + x, y, ho, wo, hout_base, share_in, share_out, acc); + + // write 0 + write_data_to_lds( + lane_id, tmp_in, share_in); + write_data_to_lds( + lane_id, tmp_out, share_out); + num_loop--; + }; + + // tail + { + run_conv_bwd_weight(x, y, ho, wo, hout_base, share_in, share_out, acc); + } + + if constexpr(ThreadPerBatch == 32) + { + //static_assert(0); + float acc_2 = warp_shuffle_down(acc, ThreadPerBatch); + acc += acc_2; + } + else if constexpr(ThreadPerBatch == 9) + { + // todo optimization reduce operation. + float acc_2 = warp_shuffle_down(acc, ThreadPerBatch); + float acc_3 = warp_shuffle_down(acc, 2 * ThreadPerBatch); + float acc_4 = warp_shuffle_down(acc, 3 * ThreadPerBatch); + float acc_5 = warp_shuffle_down(acc, 4 * ThreadPerBatch); + float acc_6 = warp_shuffle_down(acc, 5 * ThreadPerBatch); + float acc_7 = warp_shuffle_down(acc, 6 * ThreadPerBatch); + acc += acc_2 + acc_3 + acc_4 + acc_5 + acc_6 + acc_7; + } + if constexpr(NumTilePerBlock == 1 && NumWavePerTile == 1) + { + + write_output(arg, g_idx, y, x, acc); + } + else + { + uint32_t* p_share_acc = reinterpret_cast(p_share_in); + block_sync_lds(); + p_share_acc[threadIdx.x] = bit_cast(acc); + block_sync_lds(); + if(hout_base == 0 && wave_id == 0) + { + for(int i = 1; i < NumTilePerBlock * NumWavePerTile; i++) + { + acc += bit_cast(p_share_acc[i * WaveSize + lane_id]); + } + write_output(arg, g_idx, y, x, type_convert(acc)); + } + } + } + struct Argument + { + Argument(const InDataType* p_in_grid, + WeiDataType* p_wei_grid, + const OutDataType* p_out_grid, + const std::array& in_g_n_c_wis_lengths, // input + const std::array& in_g_n_c_wis_strides, + const std::array& wei_g_k_c_xs_lengths, // weight + const std::array& wei_g_k_c_xs_strides, + const std::array& out_g_n_k_wos_lengths, // output + const std::array& out_g_n_k_wos_strides) + : p_in_grid_{p_in_grid}, + p_wei_grid_{p_wei_grid}, + p_out_grid_{p_out_grid}, + in_g_n_c_wis_lengths_(in_g_n_c_wis_lengths), + in_g_n_c_wis_strides_(in_g_n_c_wis_strides), + wei_g_k_c_xs_lengths_(wei_g_k_c_xs_lengths), + wei_g_k_c_xs_strides_(wei_g_k_c_xs_strides), + out_g_n_k_wos_lengths_(out_g_n_k_wos_lengths), + out_g_n_k_wos_strides_(out_g_n_k_wos_strides) + { + } + + std::size_t GetWorkspaceSizeBytes() const { return 0; } + + const InDataType* p_in_grid_; + WeiDataType* p_wei_grid_; + const OutDataType* p_out_grid_; + + std::array in_g_n_c_wis_lengths_; + std::array in_g_n_c_wis_strides_; + std::array wei_g_k_c_xs_lengths_; + std::array wei_g_k_c_xs_strides_; + std::array out_g_n_k_wos_lengths_; + std::array out_g_n_k_wos_strides_; + }; +}; + +template + typename BlockTileSize, // input, without include pading + index_t FilterSize, // seqence typename FilterParam, // tuple typename InElementwiseOperation, typename WeiElementwiseOperation, typename OutElementwiseOperation, + index_t NBatch, + index_t NumWavePerTile, + index_t InScalarPerVector, + index_t OutScalarPerVector, + index_t DstScalarPerVector, + bool RequirePadding, typename ComputeTypeA = InDataType, - typename ComputeTypeB = ComputeTypeA - > -struct GridwiseGroupedConvBwdWeightDlV4 -{ - -} - - -template - typename FilterParam, // tuple - typename InElementwiseOperation, - typename WeiElementwiseOperation, - typename OutElementwiseOperation, - typename ComputeTypeA = InDataType, - typename ComputeTypeB = ComputeTypeA - > -struct DeviceGroupedConvBwdWeightDlV4 - : public DeviceGroupedConvBwdWeight + typename ComputeTypeB = ComputeTypeA> +struct DeviceGroupedConvBwdWeightDlV4 : public DeviceGroupedConvBwdWeight { using DeviceOp = DeviceGroupedConvBwdWeightDlV4; + static constexpr auto I0 = Number<0>{}; + static constexpr auto I1 = Number<1>{}; + + static_assert(NDimSpatial == 2); + static_assert(FilterSize * FilterSize < 64); + static_assert(RequirePadding == false); + static_assert(NBatch % DstScalarPerVector == 0); static_assert(is_same_v); static_assert(is_same_v); static_assert(is_same_v); + using GridwiseConvBwdWeight = + GridwiseGroupedConv2DBwdWeightDlV4; + struct Argument : public BaseArgument { Argument(const InDataType* p_in_grid, @@ -417,10 +685,7 @@ struct DeviceGroupedConvBwdWeightDlV4 { } - std::size_t GetWorkspaceSizeBytes() const - { - return 0; - } + std::size_t GetWorkspaceSizeBytes() const { return 0; } const InDataType* p_in_grid_; WeiDataType* p_wei_grid_; @@ -432,7 +697,7 @@ struct DeviceGroupedConvBwdWeightDlV4 std::array in_g_n_c_wis_lengths_; std::array in_g_n_c_wis_strides_; - std::array wei_g_k_c_xs_lengths_; + std::array wei_g_k_c_xs_lengths_; std::array wei_g_k_c_xs_strides_; std::array out_g_n_k_wos_lengths_; std::array out_g_n_k_wos_strides_; @@ -448,48 +713,34 @@ struct DeviceGroupedConvBwdWeightDlV4 { using Argument = DeviceOp::Argument; - void ShowInfo(const Argument&) - { + void ShowInfo(const Argument&) {} + index_t CalculateGridSize(const Argument& arg) { return arg.in_g_n_c_wis_lengths_[0]; } - } - index_t CalculateGridSize(const Argument& arg) + float Run(const Argument& arg, const StreamConfig& stream_config = StreamConfig{}) { - return arg.in_g_n_c_wis_lengths_[0];; - } - - float RunGemmV3(const Argument& arg, const StreamConfig& stream_config = StreamConfig{}) - { - index_t gdx = CalculateGridSize(arg); float ave_time = 0; - typename GridwiseConvBwdWeight::Argument conv_arg{ - p_a_grid, p_b_grid, arg.p_c_grid_, GemmM, GemmN, GemmK, I0, I0, I0, arg.k_batch_}; + typename GridwiseConvBwdWeight::Argument conv_arg{arg.p_in_grid_, + arg.p_wei_grid_, + arg.p_out_grid_, + arg.in_g_n_c_wis_lengths_, + arg.in_g_n_c_wis_strides_, + arg.wei_g_k_c_xs_lengths_, + arg.wei_g_k_c_xs_strides_, + arg.out_g_n_k_wos_lengths_, + arg.out_g_n_k_wos_strides_}; - constexpr index_t minimum_occupancy = 1; - // BlkGemmPipeSched == BlockGemmPipelineScheduler::Intrawave ? 1 : 2; - constexpr index_t BlockSize = Num_Wave * WaveSize; - const auto kernel = kernel_grouped_conv_bwd_weight_dl_v4< - Argument, minimum_occupancy>; + constexpr index_t minimum_occupancy = 1; + + const auto kernel = kernel_grouped_conv_bwd_weight_dl_v4; ave_time += launch_and_time_kernel( - stream_config, - kernel, - dim3(gdx), - dim3(BlockSize), - 0, - arg); + stream_config, kernel, dim3(gdx), dim3(BlockSize), 0, conv_arg); return ave_time; } - float Run(const Argument& arg, const StreamConfig& stream_config = StreamConfig{}) - { - float avg_time = 0.f; - avg_time += RunGemmV3(arg, stream_config); - return avg_time; - } - float Run(const BaseArgument* p_arg, const StreamConfig& stream_config = StreamConfig{}) override { @@ -503,8 +754,82 @@ struct DeviceGroupedConvBwdWeightDlV4 return true; } - static bool IsSupportedArgument(const Argument&) + static bool IsSupportedArgument(const Argument& arg) { + constexpr index_t spatial_offset = 3; + // In + const index_t hi = arg.in_g_n_c_wis_lengths_[spatial_offset + 0]; + const index_t wi = arg.in_g_n_c_wis_lengths_[spatial_offset + 1]; + const index_t wi_stride = arg.in_g_n_c_wis_strides_[spatial_offset + 1]; + // Out + //const index_t ho = arg.out_g_n_k_wos_lengths_[spatial_offset + 0]; + const index_t wo = arg.out_g_n_k_wos_lengths_[spatial_offset + 1]; + const index_t wo_stride = arg.out_g_n_k_wos_strides_[spatial_offset + 1]; + // Wei + const index_t filter_y = arg.wei_g_k_c_xs_lengths_[spatial_offset + 0]; + const index_t filter_x = arg.wei_g_k_c_xs_lengths_[spatial_offset + 1]; + const index_t filter_k = arg.wei_g_k_c_xs_lengths_[1]; + const index_t filter_c = arg.wei_g_k_c_xs_lengths_[2]; + + static constexpr index_t Tile_H = BlockTileSize{}.At(I0); + static constexpr index_t Tile_W = BlockTileSize{}.At(I1); + static constexpr index_t Pad_H = tuple_element_t<2, FilterParam>{}.At(I0); + static constexpr index_t Pad_W = tuple_element_t<2, FilterParam>{}.At(I1); + static constexpr index_t Stride_H = tuple_element_t<1, FilterParam>{}.At(I0); + static constexpr index_t Stride_W = tuple_element_t<1, FilterParam>{}.At(I1); + static constexpr index_t Dilation_Y = tuple_element_t<0, FilterParam>{}.At(I0); + static constexpr index_t Dilation_X = tuple_element_t<0, FilterParam>{}.At(I1); + + if(filter_k != 1 || filter_c != 1) + { + return false; + } + if constexpr(RequirePadding == false) + { + if(hi != Tile_H || wi != Tile_W) + { + return false; + } + } + if(filter_y != FilterSize || filter_x != FilterSize) + { + return false; + } + if(Pad_H != arg.input_left_pads_[0] || Pad_W != arg.input_left_pads_[1] || + Pad_H != arg.input_right_pads_[0] || Pad_W != arg.input_right_pads_[1]) + { + return false; + } + if(Stride_H != arg.conv_filter_strides_[0] || Stride_W != arg.conv_filter_strides_[1]) + { + return false; + } + if(Dilation_Y != arg.conv_filter_dilations_[0] || Dilation_X != arg.conv_filter_dilations_[1]) + { + return false; + } + if(InScalarPerVector > 1) + { + if(wi % InScalarPerVector != 0) + { + return false; + } + if(wi_stride != 1) + { + return false; + } + } + if(OutScalarPerVector > 1) + { + if(wo % OutScalarPerVector != 0) + { + return false; + } + if(wo_stride != 1) + { + return false; + } + } return true; } @@ -598,7 +923,39 @@ struct DeviceGroupedConvBwdWeightDlV4 std::string GetTypeString() const override { - return ""; + auto str = std::stringstream(); + + index_t Pad_H = tuple_element_t<2, FilterParam>{}.At(I0); + index_t Pad_W = tuple_element_t<2, FilterParam>{}.At(I1); + + index_t Stride_H = tuple_element_t<1, FilterParam>{}.At(I0); + index_t Stride_W = tuple_element_t<1, FilterParam>{}.At(I1); + + index_t Dilation_H = tuple_element_t<1, FilterParam>{}.At(I0); + index_t Dilation_W = tuple_element_t<1, FilterParam>{}.At(I1); + + // clang-format off + str << "DeviceGroupedConvBwdWeightDlV4<" + << NDimSpatial << ", " + << BlockSize << ", " + << InLayout::name << ", " + << WeiLayout::name << ", " + << OutLayout::name << ", " + << "BlockTileSize<" << BlockTileSize{}.At(I0) << BlockTileSize{}.At(I1) << ">, " + << "FilterSize<" << FilterSize << ","<< FilterSize << ">, " + << "Dilation<" << Dilation_H << ", " << Dilation_W<< ">, " + << "Stride<" << Stride_H << ", " << Stride_W<< ">, " + << "Pad<" << Pad_H << ", " << Pad_W<< ">, " + << "NBatch: " << NBatch<< ", " + << "NumWavePerTile: " << NumWavePerTile<< ", " + << "InScalarPerVector: " << InScalarPerVector<< ", " + << "OutScalarPerVector: " << OutScalarPerVector<< ", " + << "DstScalarPerVector: " << DstScalarPerVector<< ", " + << "RequirePadding: " << RequirePadding << ">" + << std::endl; + // clang-format on + + return str.str(); } size_t GetWorkSpaceSize(const BaseArgument* p_arg) const override @@ -629,57 +986,6 @@ struct DeviceGroupedConvBwdWeightDlV4 "DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument structure!"); } }; - - -} -} -} - -using ALayout = ck::tensor_layout::convolution::GNHWC; -using BLayout = ck::tensor_layout::convolution::GKYXC; -using ELayout = ck::tensor_layout::convolution::GNHWK; - -template -using DeviceConvBwdWeightInstance = - ck::tensor_operation::device::DeviceGroupedConvBwdWeightNaive; - -template -using HostConvBwdWeightInstance = ck::tensor_operation::host::ReferenceConvBwdWeight; - -#include "run_grouped_conv_bwd_weight_example.inc" - -int main(int argc, char* argv[]) -{ - ExecutionConfig config; - ck::utils::conv::ConvParam conv_param = DefaultConvParam; - - if(!parse_cmd_args(argc, argv, config, conv_param)) - { - return 1; - } - - switch(conv_param.num_dim_spatial_) - { - case 1: break;//return !run_grouped_conv_bwd_weight<1>(config, conv_param); - case 2: return !run_grouped_conv_bwd_weight<2>(config, conv_param); - case 3: break;//return !run_grouped_conv_bwd_weight<3>(config, conv_param); - default: break; - } - - return 1; -} +} // namesapce device +} // namespace tensor_operation +} // namespace ck diff --git a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_dl_v4_fp16.cpp b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_dl_v4_fp16.cpp index 18ca451ac2..469bd4b97d 100644 --- a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_dl_v4_fp16.cpp +++ b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_dl_v4_fp16.cpp @@ -23,15 +23,25 @@ using ELayout = ck::tensor_layout::convolution::GNHWK; template using DeviceConvBwdWeightInstance = ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4, + 5, + ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, - OutElementOp>; + OutElementOp, + 2, // N batch + 1, // NumWavePerTile + 4, // InScalarPerVector + 4, // OutScalarPerVector + 2, // DstScalarPerVector + false>; template using HostConvBwdWeightInstance = ck::tensor_operation::host::ReferenceConvBwdWeight