diff --git a/example/30_grouped_conv_fwd_multiple_d/device_grouped_conv_fwd_dl_v4.hpp b/example/30_grouped_conv_fwd_multiple_d/device_grouped_conv_fwd_dl_v4.hpp index 3354db70af..855e520508 100644 --- a/example/30_grouped_conv_fwd_multiple_d/device_grouped_conv_fwd_dl_v4.hpp +++ b/example/30_grouped_conv_fwd_multiple_d/device_grouped_conv_fwd_dl_v4.hpp @@ -127,7 +127,7 @@ struct GridwiseGroupedConv2DFwdDlV4 using WeiDataVector = typename vector_type::type; // constants for data load/store - static constexpr index_t TileIn_Pack_W = GetAlignedPackW(); + 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 = @@ -140,11 +140,12 @@ struct GridwiseGroupedConv2DFwdDlV4 static constexpr index_t TileOut_Align_H = TileOut_Pack_H * TileOut_Pack_Group; // constants for internal subtile - static constexpr index_t TilePerWave = 4; - static constexpr index_t ThreadPerTile = WaveSize / TilePerWave; static constexpr index_t HRepeate = math::integer_divide_ceil(TileOut_H, SubTileH); static constexpr index_t WRepeate = math::integer_divide_ceil(TileOut_W, SubTileW); - static_assert(HRepeate * WRepeate <= ThreadPerTile); + static constexpr index_t TilePerWave = WaveSize / (HRepeate * WRepeate); + static constexpr index_t ThreadPerTile = WaveSize / TilePerWave; + static_assert(NBatch % TilePerWave == 0); + static constexpr index_t TileIn_Max_W = SubTileW * Stride_W * (WRepeate - 1) + math::integer_least_multiple(SubTileW * Stride_W + (Filter_X - 1) * Dilation_X, @@ -167,8 +168,6 @@ struct GridwiseGroupedConv2DFwdDlV4 #endif static_assert(BlockSize == WaveSize); - // static constexpr index_t NumTilePerBlock = BlockSize / WaveSize; - template {}([&](auto wo_) { OutShareVector output = {}; - static_for<0, OutScalarPerVector_Internal, 1>{}([&](auto i) { - output[i.value] = - type_convert(acc[wo_ * OutScalarPerVector_Internal + i]); - }); + if constexpr(OutScalarPerVector_Internal == 1) + { + output = type_convert(acc[wo_ * OutScalarPerVector_Internal]); + } + else + { + static_for<0, OutScalarPerVector_Internal, 1>{}([&](auto i) { + output[i.value] = + type_convert(acc[wo_ * OutScalarPerVector_Internal + i]); + }); + } //if (threadIdx.x == 4) { // printf("threadIdx %u outut %04x %04x\n", threadIdx.x, bit_cast(output[0]), bit_cast(output[1])); @@ -331,7 +336,7 @@ struct GridwiseGroupedConv2DFwdDlV4 }; // constexpr auto SubTileInH = SubTileH * Stride_H + (Filter_Y - 1) * Dilation_Y; - constexpr auto SubTileInW = SubTileW * Stride_W + (Filter_X - 1) * Dilation_X; + constexpr auto SubTileInW = math::integer_least_multiple(SubTileW * Stride_W + (Filter_X - 1) * Dilation_X, InScalarPerVector_Internal); static_assert(SubTileInW % InScalarPerVector_Internal == 0); static_assert(SubTileW % OutScalarPerVector_Internal == 0); diff --git a/example/30_grouped_conv_fwd_multiple_d/grouped_conv_fwd_dl_fp16.cpp b/example/30_grouped_conv_fwd_multiple_d/grouped_conv_fwd_dl_fp16.cpp index 6b70aab62e..d808150329 100644 --- a/example/30_grouped_conv_fwd_multiple_d/grouped_conv_fwd_dl_fp16.cpp +++ b/example/30_grouped_conv_fwd_multiple_d/grouped_conv_fwd_dl_fp16.cpp @@ -50,9 +50,22 @@ using AccType = AccDataType; using OutType = OutKernelDataType; using DeviceConvFwdFactory = std::tuple< + // NDimSpatial BlockSize In Wei Acc Out BlockTileSize FilterSize FilterParam (dilation, stride, padding) NBatch SubTileH W ScalarPerVector(in out) RequirePadding> - // ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<28, 28>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 4, 4, 4, 4, 4, false> - ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<28, 28>, 5, ck::Tuple, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 4, 2, false> + ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<7, 7>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 1, 1, false> + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<14, 14>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 2, 2, false> + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<28, 28>, 5, ck::Tuple, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 4, 4, false> + + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<28, 28>, 5, ck::Tuple, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 4, 2, false> + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<56, 56>, 5, ck::Tuple, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 8, 4, false> + + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<7, 7>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 1, 1, false> + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<14, 14>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 2, 2, false> + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<56, 56>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 8, 7, 8, 8, 8, false> + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<112, 112>, 3, ck::Tuple, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 14, 16, 8, 8, false> + + , ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<2, 64, InType, WeiType, AccType, OutType, S<28, 28>, 3, ck::Tuple, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 32, 4, 4, 4, 2, false> +//#if 0#endif >; #include "run_grouped_conv_fwd_example.inc" diff --git a/script/test_example_conv_fwd.sh b/script/test_example_conv_fwd.sh new file mode 100755 index 0000000000..df561fe651 --- /dev/null +++ b/script/test_example_conv_fwd.sh @@ -0,0 +1,23 @@ +#!/bin/bash +EXAMPLE="../build/bin/example_grouped_conv_fwd_dl_fp16" + +set -x + +# G N K C Y X H W Sy Sx Dy Dx Pad +$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 960 128 1 1 5 5 14 14 1 1 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 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 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 960 128 1 1 3 3 14 14 1 1 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 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 +