ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 03:30:11 +00:00
Code Issues Packages Projects Releases Wiki Activity

add test script and basic instance

Browse Source
This commit is contained in:
Qun Lin
2025-06-12 08:54:51 +08:00
parent 746925d4c1
commit fdcbf3dd19
3 changed files with 55 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

29
example/30_grouped_conv_fwd_multiple_d/device_grouped_conv_fwd_dl_v4.hpp
View File

@@ -127,7 +127,7 @@ struct GridwiseGroupedConv2DFwdDlV4
using WeiDataVector = typename vector_type<WeiDataType, WeiScalarPerVector>::type;
// constants for data load/store
static constexpr index_t TileIn_Pack_W = GetAlignedPackW<TileIn_W, InScalarPerVector>();
static constexpr index_t TileIn_Pack_W = GetAlignedPackW<Tile_W, InScalarPerVector>();
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 <index_t TileH,
index_t AlignedPackW,
index_t ScalarPerVector,
@@ -290,7 +289,6 @@ struct GridwiseGroupedConv2DFwdDlV4
ignore = w_max;
#endif
static_assert(SubTileW % OutScalarPerVector == 0);
static_assert((SubTileW * Stride_W + (Filter_X - 1) * Dilation_X) % InScalarPerVector == 0);
static_assert(WeiScalarPerVector == 2);
static_assert(Filter_X % 2 == 1);
static_assert(Dilation_X == 1);
@@ -303,10 +301,17 @@ struct GridwiseGroupedConv2DFwdDlV4
auto set_out = [&](index_t ho_, auto count, float* acc) {
static_for<0, count / OutScalarPerVector_Internal, 1>{}([&](auto wo_) {
OutShareVector output = {};
static_for<0, OutScalarPerVector_Internal, 1>{}([&](auto i) {
output[i.value] =
type_convert<OutDataType>(acc[wo_ * OutScalarPerVector_Internal + i]);
});
if constexpr(OutScalarPerVector_Internal == 1)
{
output = type_convert<OutDataType>(acc[wo_ * OutScalarPerVector_Internal]);
}
else
{
static_for<0, OutScalarPerVector_Internal, 1>{}([&](auto i) {
output[i.value] =
type_convert<OutDataType>(acc[wo_ * OutScalarPerVector_Internal + i]);
});
}
//if (threadIdx.x == 4)
{
// printf("threadIdx %u outut %04x %04x\n", threadIdx.x, bit_cast<uint16_t>(output[0]), bit_cast<uint16_t>(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);

17
example/30_grouped_conv_fwd_multiple_d/grouped_conv_fwd_dl_fp16.cpp
View File

@@ -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<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<1,1>, 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<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<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<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<1,1>, 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<1,1>, 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>, 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>, 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>, 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>, 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<1,1>, 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"