From 289fdf59e175afc266c54a3407f231085ed0dc6f Mon Sep 17 00:00:00 2001 From: "Lin, Qun" Date: Tue, 3 Jun 2025 23:13:44 -0500 Subject: [PATCH] finalize code --- .../device_grouped_conv_bwd_weight_dl_v4.hpp | 442 +++++++++--------- .../grouped_conv_bwd_weight_dl_v4_fp16.cpp | 23 +- script/test_example_conv_bwd.sh | 28 +- 3 files changed, 254 insertions(+), 239 deletions(-) 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 4d61ff3a3e..c6a6eed9d8 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 @@ -16,9 +16,7 @@ using InElementOp = PassThrough; using WeiElementOp = PassThrough; using OutElementOp = PassThrough; -//#define ENABLE_PIPELINE_V2 1 - -template struct Debug; +#define ENABLE_PIPELINE_V2 1 namespace ck { template @@ -36,25 +34,19 @@ __device__ T warp_shuffle_down(const T& v_local, uint32_t lane_delta) #endif } -template +template __global__ void #if CK_USE_LAUNCH_BOUNDS __launch_bounds__(BlockSize, MinimumOccupancy) #endif kernel_grouped_conv_bwd_weight_dl_v4(typename GridwiseConvBwdWeight::Argument arg) { -//#if(!defined(__HIP_DEVICE_COMPILE__)) __shared__ char p_share_in[GridwiseConvBwdWeight::ShareMemInSize * GridwiseConvBwdWeight::NumTilePerBlock]; - __shared__ char - p_share_out[GridwiseConvBwdWeight::ShareMemOutSize * 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 } namespace tensor_operation { @@ -85,7 +77,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 { return (length + pad + pad - ((filter - 1) * dilation + 1)) / stride + 1; } - template + template static constexpr index_t GetAlignedPackW() { constexpr index_t pakced_w = W / ScalarPerVector; @@ -99,11 +91,11 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 } } static constexpr index_t GetBatchPerWave() { return WaveSize / (FilterSize * FilterSize); } - - static constexpr index_t NDimSpatial = 2; + + 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 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); @@ -136,19 +128,20 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 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_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 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 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 = @@ -159,46 +152,83 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 using InDataVector = typename vector_type::type; using OutDataVector = typename vector_type::type; - template + template static void __device__ load_data_from_global(const SrcType* p, - index_t x, - index_t y_offset, - index_t n_stride, - index_t h, - index_t w, - index_t h_stride, - index_t w_stride, - DestVector* p_scratch) + index_t x, + index_t y_offset, + 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; + using SrcVector = typename vector_type::type; 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 % AlignedPackW; - //const index_t y_offset = lane_id / AlignedPackW; + constexpr index_t NumGroup = WaveSize / AlignedPackW; + constexpr index_t AlignedPackH = math::integer_divide_ceil(TileH, NumGroup); + constexpr index_t PackH = TileH / NumGroup; auto get_offset = [&](index_t y_, index_t packed_x_, index_t n_) { - return (y_ * h_stride + packed_x_ * ScalarPerVector * w_stride + n_ * n_stride) / ScalarPerVector; + return (y_ * h_stride + packed_x_ * ScalarPerVector * 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); - ignore = PackW; - //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; + if constexpr(DstScalarPerVector == 1) + { + static_assert(NBatch == 1); + static_for<0, ScalarPerVector, 1>{}( + [&](auto j) { p_scratch_base[j * NumVectorPerPixel] = tmp[0][j.value]; }); + } + else if constexpr(ScalarPerVector == 1) + { + static_assert(DstScalarPerVector > 1); + static_for<0, NBatch, 1>{}([&](auto n) { + p_scratch_base[n / DstScalarPerVector][n % DstScalarPerVector] = tmp[n]; + }); + } + else + { + 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_for<0, PackH, 1>{}([&](auto i) { - const index_t y = y_offset + i * NumGroup; + 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) { @@ -216,7 +246,6 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 } else if constexpr(ScalarPerVector == 1) { - static_assert(DstScalarPerVector > 1); static_for<0, NBatch, 1>{}([&](auto n) { p_scratch_base[n / DstScalarPerVector][n % DstScalarPerVector] = tmp[n]; }); @@ -230,107 +259,62 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 }); }); } - }); - - if constexpr(AlignedPackH != PackH) - { - 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; - if constexpr(DstScalarPerVector == 1) - { - static_assert(NBatch == 1); - static_for<0, ScalarPerVector, 1>{}([&](auto j) { - p_scratch_base[j * NumVectorPerPixel] = tmp[0][j.value]; - }); - } - else if constexpr(ScalarPerVector == 1) - { - static_for<0, NBatch, 1>{}([&](auto n) { - p_scratch_base[n / DstScalarPerVector][n % DstScalarPerVector] = tmp[n]; - }); - } - else - { - 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]; - }); - }); - } - } } } } // todo handle pading in p_sharemem - template + template static void __device__ write_data_to_lds(index_t x, - index_t y_offset, - const DestVector* p_scratch, - DestVector* p_sharemem) + index_t y_offset, + const DestVector* p_scratch, + DestVector* p_sharemem) { - constexpr index_t PackW = TileW / ScalarPerVector; constexpr index_t AlignedPackW = GetAlignedPackW(); - static_assert(AlignedPackW < WaveSize); + static_assert(AlignedPackW <= WaveSize); - 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 % AlignedPackW; - //const index_t y_offset = lane_id / AlignedPackW; + constexpr index_t NumGroup = WaveSize / AlignedPackW; + constexpr index_t AlignedPackH = math::integer_divide_ceil(TileH, NumGroup); + constexpr index_t PackH = TileH / NumGroup; auto get_offset = [&](index_t y_, index_t x_) { return y_ * TileW_Stride * NumVectorPerPixel + x_ * NumVectorPerPixel * ScalarPerVector; }; - ignore = PackW; - //if(x < PackW) + static_for<0, PackH, 1>{}([&](auto i) { + const index_t y = y_offset + i * NumGroup; + 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_for<0, PackH, 1>{}([&](auto i) { + if(y_offset < (TileH - NumGroup * PackH)) + { + constexpr auto i = PackH; const index_t y = y_offset + i * NumGroup; 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 * NumGroup; - 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) + 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; @@ -341,7 +325,8 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 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_]; + return p_share_out[((ho_ + hout_base) * TileOut_Align_W + wo_) * NumVectorPerPixel + + i_]; }; if(x < Filter_X && y < Filter_Y) { @@ -362,6 +347,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 } else { + // for (index_t ho = 0; ho < TileH; ho ++) { static_for<0, TileH, 1>{}([&](auto ho) { static_for<0, TileOut_W, 1>{}([&](auto wo) { // for (index_t wo = 0; wo < TileOut_W; wo ++) { @@ -374,10 +360,9 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 // { // uint32_t * pin = reinterpret_cast(&v_in); // uint32_t * pout = reinterpret_cast(&v_out); - // printf("h w [%d %d] vin %08x %08x %08x %08x vout %08x %08x %08x - // %08x acc = %f\n", ho+ hout_base, wo, pin[0], pin[1], pin[2], - // pin[3], pout[0], pout[1], pout[2],pout[3], acc); - // } + // printf("h w [%d %d] vin %08x vout %08x acc = %f\n", ho+ hout_base, + // wo, pin[0], pout[0], acc); + //} }); }); }); @@ -385,7 +370,8 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 } } template - static void __device__ write_output(const Argument& arg, index_t g, index_t y, index_t x, WeiDataType acc) + 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]; @@ -400,22 +386,21 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 template static void __device__ dump_lds(DstVector* p, index_t totalcount, index_t length) { - for (index_t i = 0; i < totalcount; i++) + for(index_t i = 0; i < totalcount; i++) { - if (i % length ==0) + if(i % length == 0) { - printf("\n [%d]", i/length); + printf("\n [%d]", i / length); } uint32_t* p1 = reinterpret_cast(&p[i]); - static_for<0, sizeof(DstVector)/ sizeof(uint32_t), 1>{}([&](auto j) { - printf("%08x ", p1[j]); - }); + static_for<0, sizeof(DstVector) / sizeof(uint32_t), 1>{}( + [&](auto j) { printf("%08x ", p1[j]); }); } printf("\n"); } static constexpr index_t TotalLdsSize() { - return (ShareMemInSize + ShareMemOutSize) * NumTilePerBlock; + return (ShareMemInSize + ShareMemOutSize) * NumTilePerBlock; } template @@ -427,19 +412,21 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 const index_t lane_id = __lane_id(); constexpr index_t ThreadPerBatch = WaveSize / BatchPerWave; - //Debug> xx3; + // Debug> xx3; static_assert(Tile_H % NumWavePerTile == 0); static_assert(TileOut_H % NumWavePerTile == 0); - InDataVector tmp_in[math::integer_divide_ceil(TileIn_Pack_H, NumWavePerTile) * NumVectorPerPixel * InScalarPerVector] = {}; - OutDataVector tmp_out[math::integer_divide_ceil(TileOut_Pack_H, NumWavePerTile) * NumVectorPerPixel * OutScalarPerVector] = {}; + InDataVector tmp_in[math::integer_divide_ceil(TileIn_Pack_H, NumWavePerTile) * + NumVectorPerPixel * InScalarPerVector] = {}; + OutDataVector tmp_out[math::integer_divide_ceil(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 / NBatch - 1; - index_t n_idx = n / NumTilePerBlock * tile_id; - if constexpr (NumTilePerBlock > 1) + const index_t n = arg.in_g_n_c_wis_lengths_[1]; + index_t num_loop = n / NumTilePerBlock / NBatch - 1; + index_t n_idx = n / NumTilePerBlock * tile_id; + if constexpr(NumTilePerBlock > 1) { if(tile_id == NumTilePerBlock - 1) { @@ -470,19 +457,17 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 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_Tile_H = Tile_H / NumWavePerTile; + constexpr index_t Copy_Tile_H = Tile_H / NumWavePerTile; constexpr index_t Copy_TileOut_H = TileOut_H / NumWavePerTile; - if constexpr (NumWavePerTile > 1) + if constexpr(NumWavePerTile > 1) { static_assert(RequirePadding == false); p_in += Copy_Tile_H * hi_stride * (wave_id % NumWavePerTile); p_out += Copy_TileOut_H * ho_stride * (wave_id % NumWavePerTile); } - InDataVector* share_in = - reinterpret_cast(p_share_in); - OutDataVector* share_out = - reinterpret_cast(p_share_out); + 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); @@ -492,32 +477,45 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 // 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) + static_for<0, math::integer_divide_ceil(count, WaveSize * NumWavePerTile), 1>{}( + [&](auto i) { + if(cluster_id + i * WaveSize * NumWavePerTile < count) { - auto p = share_vec + (cluster_id + i * WaveSize * NumWavePerTile) * stride + j; - *p = {}; + 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 + + (cluster_id + i * WaveSize * NumWavePerTile) * 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; + 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 + Pad_W) * NumVectorPerPixel , Number{}, Number{}, TileIn_Align_W * NumVectorPerPixel); + init_array_pading(share_in + TopPadingSize, + Number{}, + Number{}, + TileIn_Align_W * NumVectorPerPixel); + init_array_pading(share_in + TopPadingSize + (Tile_W + Pad_W) * NumVectorPerPixel, + Number{}, + Number{}, + TileIn_Align_W * NumVectorPerPixel); } if constexpr(Pad_H > 0) { @@ -525,18 +523,19 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 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; + constexpr index_t OutButtomPaddingSize = + (ShareMemOutSize / (sizeof(OutDataType) * NBatch) - TileOutEnd) * NumVectorPerPixel; init_pading(share_out + TileOutEnd * NumVectorPerPixel, Number{}); - if constexpr (NumWavePerTile > 1) + if constexpr(NumWavePerTile > 1) { block_sync_lds(); } - const index_t in_x = lane_id % (Tile_W / InScalarPerVector); - const index_t in_y_offset = lane_id / (Tile_W / InScalarPerVector); - const index_t out_x = lane_id % (TileOut_W / OutScalarPerVector); - const index_t out_y_offset = lane_id / (TileOut_W / OutScalarPerVector); + const index_t in_x = lane_id % TileIn_Pack_W; + const index_t in_y_offset = lane_id / TileIn_Pack_W; + const index_t out_x = lane_id % TileOut_Pack_W; + const index_t out_y_offset = lane_id / TileOut_Pack_W; // prefetch 0 if(in_x < (Tile_W / InScalarPerVector)) @@ -576,7 +575,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 constexpr index_t TileOut_H_batch = math::integer_divide_ceil(Copy_TileOut_H, BatchPerWave); index_t hout_base = lane_id / ThreadPerBatch * TileOut_H_batch; - if constexpr (NumWavePerTile > 1) + if constexpr(NumWavePerTile > 1) { hout_base += Copy_TileOut_H * wave_id; } @@ -584,19 +583,23 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 index_t y = (lane_id % ThreadPerBatch) / Filter_X; float acc = 0; - //if (lane_id == 0) + // if (lane_id == 0) //{ - // dump_lds(tmp_in, sizeof(tmp_in)/sizeof(InDataVector), sizeof(tmp_in)/sizeof(InDataVector)); - // dump_lds(tmp_out, sizeof(tmp_out)/sizeof(OutDataVector), sizeof(tmp_out)/sizeof(OutDataVector)); + // dump_lds(tmp_in, sizeof(tmp_in)/sizeof(InDataVector), + // sizeof(tmp_in)/sizeof(InDataVector)); + // dump_lds(tmp_out, sizeof(tmp_out)/sizeof(OutDataVector), + // sizeof(tmp_out)/sizeof(OutDataVector)); //} - //block_sync_lds(); - //if (threadIdx.x == 0) - //{ - // dump_lds(reinterpret_cast(p_share_in), ShareMemInSize/sizeof(InDataVector), TileIn_Align_W * NumVectorPerPixel); - // dump_lds(reinterpret_cast(p_share_out), ShareMemOutSize/sizeof(OutDataVector), TileOut_Align_W * NumVectorPerPixel); + // block_sync_lds(); + // if (threadIdx.x == 0) + //{ + // dump_lds(reinterpret_cast(p_share_in), + // ShareMemInSize/sizeof(InDataVector), TileIn_Align_W * NumVectorPerPixel); + // dump_lds(reinterpret_cast(p_share_out), + // ShareMemOutSize/sizeof(OutDataVector), TileOut_Align_W * NumVectorPerPixel); //} - //block_sync_lds(); - + // block_sync_lds(); + #if defined(ENABLE_PIPELINE_V2) while(num_loop > 0) { @@ -607,7 +610,10 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 } run_conv_bwd_weight( x, y, ho, wo, hout_base, share_in_base, share_out_base, acc); - + if constexpr(NumWavePerTile > 1) + { + block_sync_lds(); + } if(in_x < (Tile_W / InScalarPerVector)) { load_data_from_global( @@ -666,7 +672,10 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 } run_conv_bwd_weight( x, y, ho, wo, hout_base, share_in_base, share_out_base, acc); - + if constexpr(NumWavePerTile > 1) + { + block_sync_lds(); + } // write 0 if(in_x < (Tile_W / InScalarPerVector)) { @@ -683,11 +692,12 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 #endif // tail { - if constexpr (NumWavePerTile > 1) + if constexpr(NumWavePerTile > 1) { block_sync_lds(); } - run_conv_bwd_weight(x, y, ho, wo, hout_base, share_in_base, share_out_base, acc); + run_conv_bwd_weight( + x, y, ho, wo, hout_base, share_in_base, share_out_base, acc); } if constexpr(ThreadPerBatch == 32) @@ -708,8 +718,10 @@ struct GridwiseGroupedConv2DBwdWeightDlV4 } if constexpr(NumTilePerBlock == 1 && NumWavePerTile == 1) { - - write_output(arg, g_idx, y, x, acc); + if(hout_base == 0) + { + write_output(arg, g_idx, y, x, acc); + } } else { @@ -774,7 +786,7 @@ template + index_t FilterSize, // seqence typename FilterParam, // tuple typename InElementwiseOperation, typename WeiElementwiseOperation, @@ -800,7 +812,7 @@ struct DeviceGroupedConvBwdWeightDlV4 : public DeviceGroupedConvBwdWeight { - using DeviceOp = DeviceGroupedConvBwdWeightDlV4; + using DeviceOp = DeviceGroupedConvBwdWeightDlV4; static constexpr auto I0 = Number<0>{}; static constexpr auto I1 = Number<1>{}; @@ -812,23 +824,22 @@ struct DeviceGroupedConvBwdWeightDlV4 : public DeviceGroupedConvBwdWeight); static_assert(is_same_v); - using GridwiseConvBwdWeight = - GridwiseGroupedConv2DBwdWeightDlV4; + using GridwiseConvBwdWeight = GridwiseGroupedConv2DBwdWeightDlV4; struct Argument : public BaseArgument { @@ -908,16 +919,19 @@ struct DeviceGroupedConvBwdWeightDlV4 : public DeviceGroupedConvBwdWeight (32 * 1024) ? 1 : 2; + constexpr index_t minimum_occupancy = + 1; // GridwiseConvBwdWeight::TotalLdsSize() > (32 * 1024) ? 1 : 2; - const auto kernel = kernel_grouped_conv_bwd_weight_dl_v4; + const auto kernel = kernel_grouped_conv_bwd_weight_dl_v4; ave_time += launch_and_time_kernel( stream_config, kernel, dim3(gdx), dim3(BlockSize), 0, conv_arg); @@ -946,7 +960,6 @@ struct DeviceGroupedConvBwdWeightDlV4 : public DeviceGroupedConvBwdWeight, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 2, 8, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<7, 7>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 16, 1, 1, 1, 8, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<56, 56>, 5, ck::Tuple, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 4, 2, 2, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 5, ck::Tuple, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 1, 8, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 8, 1, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<56, 56>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 4, 4, 2, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 2, 8, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 4, 1, false> - , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 3, ck::Tuple, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 2, 2, false> + ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 2, 8, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<7, 7>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 16, 1, 1, 1, 8, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<56, 56>, 5, ck::Tuple, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 2, 2, 2, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 5, ck::Tuple, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 1, 8, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 8, 1, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<56, 56>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 2, 4, 4, 1, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 2, 4, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<7, 7>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 16, 1, 1, 1, 8, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 4, 1, false> + , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 3, ck::Tuple, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 2, 2, false> >; template diff --git a/script/test_example_conv_bwd.sh b/script/test_example_conv_bwd.sh index 42204de941..1b15d71299 100755 --- a/script/test_example_conv_bwd.sh +++ b/script/test_example_conv_bwd.sh @@ -4,21 +4,21 @@ EXAMPLE="../build/bin/example_grouped_conv_bwd_weight_dl_v4_fp16" set -x # G N K C Y X H W Sy Sx Dy Dx Pad -EXAMPLE 1 2 1 2 480 128 1 1 5 5 28 28 1 1 1 1 2 2 2 2 -EXAMPLE 1 2 1 2 960 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2 -EXAMPLE 1 2 1 2 1344 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2 -EXAMPLE 1 2 1 2 2304 128 1 1 5 5 7 7 1 1 1 1 2 2 2 2 -EXAMPLE 1 2 1 2 288 128 1 1 5 5 56 56 2 2 1 1 2 2 2 2 -EXAMPLE 1 2 1 2 1344 128 1 1 5 5 28 28 2 2 1 1 2 2 2 2 +$EXAMPLE 1 2 1 2 480 128 1 1 5 5 28 28 1 1 1 1 2 2 2 2 +$EXAMPLE 1 2 1 2 960 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2 +$EXAMPLE 1 2 1 2 1344 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2 +$EXAMPLE 1 2 1 2 2304 128 1 1 5 5 7 7 1 1 1 1 2 2 2 2 +$EXAMPLE 1 2 1 2 288 128 1 1 5 5 56 56 2 2 1 1 2 2 2 2 +$EXAMPLE 1 2 1 2 1344 128 1 1 5 5 14 14 2 2 1 1 2 2 2 2 -EXAMPLE 1 2 1 2 288 128 1 1 3 3 56 56 1 1 1 1 1 1 1 1 -EXAMPLE 1 2 1 2 64 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1 -EXAMPLE 1 2 1 2 32 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1 -EXAMPLE 1 2 1 2 960 128 1 1 3 3 14 14 1 1 1 1 1 1 1 1 -EXAMPLE 1 2 1 2 2304 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1 -EXAMPLE 1 2 1 2 3840 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1 -EXAMPLE 1 2 1 2 480 128 1 1 3 3 28 28 2 2 1 1 1 1 1 1 -EXAMPLE 1 2 1 2 192 128 1 1 3 3 112 112 2 2 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 288 128 1 1 3 3 56 56 1 1 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 64 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 32 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 960 128 1 1 3 3 14 14 1 1 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 2304 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 3840 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 480 128 1 1 3 3 28 28 2 2 1 1 1 1 1 1 +$EXAMPLE 1 2 1 2 192 128 1 1 3 3 112 112 2 2 1 1 1 1 1 1 set +x