ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-14 19:18:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

add device factory and fix build error.

Browse Source
This commit is contained in:
Qun Lin
2025-06-02 15:05:10 +08:00
parent b27dddfb59
commit eb158d9b86
3 changed files with 167 additions and 62 deletions
Download Patch File Download Diff File Expand all files Collapse all files

88
example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp
View File

@@ -152,7 +152,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
static constexpr index_t ShareMemOutSize =
TileOut_Align_H * TileOut_Align_W * sizeof(OutDataType) * NBatch;
static constexpr index_t NumTilePerBlock = BlockSize / WaveSize;
static constexpr index_t NumTilePerBlock = BlockSize / WaveSize / NumWavePerTile;
using InDataVector = typename vector_type<InDataType, DstScalarPerVector>::type;
using OutDataVector = typename vector_type<OutDataType, DstScalarPerVector>::type;
@@ -204,21 +204,37 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
// interleave data
auto* p_scratch_base = p_scratch + i * NumVectorPerPixel * ScalarPerVector;
static_for<0, ScalarPerVector, 1>{}([&](auto j) {
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[j * NumVectorPerPixel + n / DstScalarPerVector]
[n / DstScalarPerVector] = tmp[n][j.value];
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_assert(NumWavePerTile == 1);
if (y_offset < (TileH - NumGroup * 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 y = y_offset + i * NumGroup;
// load data
SrcVector tmp[NBatch];
static_for<0, NBatch, 1>{}([&](auto n) {
@@ -228,12 +244,28 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
// 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(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];
});
});
}
}
}
}
@@ -245,7 +277,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
const DestVector* p_scratch,
DestVector* p_sharemem)
{
//constexpr index_t PackW = TileW / ScalarPerVector;
constexpr index_t PackW = TileW / ScalarPerVector;
constexpr index_t AlignedPackW = GetAlignedPackW<TileW, ScalarPerVector>();
static_assert(AlignedPackW < WaveSize);
@@ -260,7 +292,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
return (y_ * TileW_Stride + x_) * NumVectorPerPixel;
};
//if(x < PackW)
if(x < PackW)
{
static_for<0, PackH, 1>{}([&](auto i) {
const index_t y = y_offset + i * AlignedPackW;
@@ -271,7 +303,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
});
if constexpr(AlignedPackH != PackH)
{
static_assert(NumWavePerTile == 1);
//static_assert(NumWavePerTile == 1);
if (y_offset < (TileH - NumGroup * PackH))
{
constexpr auto i = PackH;
@@ -342,11 +374,11 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
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] = {};
//Debug<Sequence<TileIn_Pack_H, TileOut_Pack_H>> 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] = {};
static_assert(NumTilePerBlock == 1 || NumWavePerTile == 1);
@@ -383,12 +415,12 @@ 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_TileIn_H = TileIn_H / NumWavePerTile;
constexpr index_t Copy_Tile_H = Tile_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_in += Copy_Tile_H * hi_stride * wave_id;
p_out += Copy_TileOut_H * ho_stride * wave_id;
}
@@ -447,7 +479,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
}
// prefetch 0
load_data_from_global<Copy_TileIn_H, TileIn_W, InScalarPerVector>(
load_data_from_global<Copy_Tile_H, Tile_W, InScalarPerVector>(
p_in, lane_id, in_n_stride, hi, wi, hi_stride, wi_stride, tmp_in);
load_data_from_global<Copy_TileOut_H, TileOut_W, OutScalarPerVector>(
p_out, lane_id, out_n_stride, ho, wo, ho_stride, wo_stride, tmp_out);
@@ -459,11 +491,11 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
if constexpr (NumWavePerTile > 1)
{
static_assert(RequirePadding == false);
share_in += Copy_TileIn_H * TileIn_Align_W * NumVectorPerPixel * wave_id;
share_in += Copy_Tile_H * TileIn_Align_W * NumVectorPerPixel * wave_id;
share_out += Copy_TileOut_H * TileOut_Align_W * NumVectorPerPixel * wave_id;
}
write_data_to_lds<Copy_TileIn_H, TileIn_W, TileIn_Align_W, InScalarPerVector>(
write_data_to_lds<Copy_Tile_H, Tile_W, TileIn_Align_W, InScalarPerVector>(
lane_id, tmp_in, share_in);
write_data_to_lds<Copy_TileOut_H, TileOut_W, TileOut_Align_W, OutScalarPerVector>(
lane_id, tmp_out, share_out);
@@ -481,7 +513,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
while(num_loop > 0)
{
load_data_from_global<Copy_TileIn_H, TileIn_W, InScalarPerVector>(
load_data_from_global<Copy_Tile_H, Tile_W, InScalarPerVector>(
p_in, lane_id, in_n_stride, hi, wi, hi_stride, wi_stride, tmp_in);
load_data_from_global<Copy_TileOut_H, TileOut_W, OutScalarPerVector>(
p_out, lane_id, out_n_stride, ho, wo, ho_stride, wo_stride, tmp_out);
@@ -493,7 +525,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
x, y, ho, wo, hout_base, share_in, share_out, acc);
// write 0
write_data_to_lds<Copy_TileIn_H, TileIn_W, TileIn_Align_W, InScalarPerVector>(
write_data_to_lds<Copy_Tile_H, Tile_W, TileIn_Align_W, InScalarPerVector>(
lane_id, tmp_in, share_in);
write_data_to_lds<Copy_TileOut_H, TileOut_W, TileOut_Align_W, OutScalarPerVector>(
lane_id, tmp_out, share_out);

16
example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_dl_v4_fp16.cpp
View File

@@ -6,6 +6,7 @@
#include "ck/utility/blkgemmpipe_scheduler.hpp"
#include "device_grouped_conv_bwd_weight_dl_v4.hpp"
#define ENABLE_CONV_FACTORY 1
using InDataType = F16;
using WeiDataType = F16;
@@ -43,6 +44,21 @@ using DeviceConvBwdWeightInstance =
2, // DstScalarPerVector
false>;
using DeviceConvBwdWeightFactory = std::tuple<
// NDimSpatial BlockSize InLayout WeiLayout OutLayout InDataType WeiDataType OutDataType BlockTileSize FilterSize FilterParam(dilation, stride, pad) NBatch NumWavePerTile InScalarPerVector OutScalarPerVector DstScalarPerVector RequirePadding
ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 5, ck::Tuple<S<1,1>, 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<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<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<1,1>, 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<1,1>, 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>, 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>, 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>, 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>, 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<1,1>, 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<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 2, 2, false>
>;
template <ck::index_t NDimSpatial>
using HostConvBwdWeightInstance = ck::tensor_operation::host::ReferenceConvBwdWeight<NDimSpatial,
InDataType,

125
example/20_grouped_conv_bwd_weight/run_grouped_conv_bwd_weight_example.inc
View File

@@ -75,7 +75,82 @@ bool run_grouped_conv_bwd_weight(const ExecutionConfig& config,
range_copy(conv_param.input_left_pads_, begin(input_left_pads));
range_copy(conv_param.input_right_pads_, begin(input_right_pads));
// do GEMM
// do CONV
if(config.do_verification)
{
auto ref_conv = HostConvBwdWeightInstance<NDimSpatial>{};
auto ref_invoker = ref_conv.MakeInvoker();
auto ref_argument = ref_conv.MakeArgument(in,
wei_host_result,
out,
conv_param.conv_filter_strides_,
conv_param.conv_filter_dilations_,
conv_param.input_left_pads_,
conv_param.input_right_pads_,
InElementOp{},
WeiElementOp{},
OutElementOp{},
{},
{},
{});
ref_invoker.Run(ref_argument);
}
#if ENABLE_CONV_FACTORY
ck::static_for<0, std::tuple_size_v<DeviceConvBwdWeightFactory>, 1>{}([&](auto i) -> void {
const auto device_conv_bwd_weight_instance = std::get<i>(DeviceConvBwdWeightFactory{});
using DeviceConvBwdWeightInstance = ck::remove_cvref_t<decltype(device_conv_bwd_weight_instance)>;
auto conv = DeviceConvBwdWeightInstance{};
auto invoker = conv.MakeInvoker();
auto argument = conv.MakeArgument(static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
input_lengths,
input_strides,
filter_lengths,
weights_strides,
output_lengths,
output_strides,
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads,
InElementOp{},
WeiElementOp{},
OutElementOp{},
split_k);
DeviceMem gemm_workspace_dev(conv.GetWorkSpaceSize(&argument));
conv.SetWorkSpacePointer(&argument, gemm_workspace_dev.GetDeviceBuffer());
if(conv.IsSupportedArgument(argument))
{
std::cout << "Run conv :" << conv.GetTypeString() << std::endl;
invoker.ShowInfo(argument);
float avg_time = invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
if(config.time_kernel)
{
std::size_t flop = conv_param.GetFlops();
std::size_t num_btype = conv_param.GetByte<InDataType, WeiDataType, OutDataType>();
float tflops = static_cast<float>(flop) / 1.E9 / avg_time;
float gb_per_sec = num_btype / 1.E6 / avg_time;
std::cerr << "Perf: " << avg_time << " ms, " << tflops << " TFlops, " << gb_per_sec
<< " GB/s" << std::endl
<< "DeviceOp: " << conv.GetTypeString() << std::endl;
}
if(config.do_verification)
{
wei_device_buf.FromDevice(wei_device_result.mData.data());
ck::utils::check_err(wei_device_result.mData, wei_host_result.mData);
}
}
});
#else
auto conv = DeviceConvBwdWeightInstance<NDimSpatial>{};
auto invoker = conv.MakeInvoker();
auto argument = conv.MakeArgument(static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
@@ -107,45 +182,27 @@ bool run_grouped_conv_bwd_weight(const ExecutionConfig& config,
return true;
}
invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
float avg_time = invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
if(config.time_kernel)
{
std::size_t flop = conv_param.GetFlops();
std::size_t num_btype = conv_param.GetByte<InDataType, WeiDataType, OutDataType>();
float tflops = static_cast<float>(flop) / 1.E9 / avg_time;
float gb_per_sec = num_btype / 1.E6 / avg_time;
std::cerr << "Perf: " << avg_time << " ms, " << tflops << " TFlops, " << gb_per_sec
<< " GB/s" << std::endl
<< "DeviceOp: " << conv.GetTypeString() << std::endl;
}
if(config.do_verification)
{
auto ref_conv = HostConvBwdWeightInstance<NDimSpatial>{};
auto ref_invoker = ref_conv.MakeInvoker();
auto ref_argument = ref_conv.MakeArgument(in,
wei_host_result,
out,
conv_param.conv_filter_strides_,
conv_param.conv_filter_dilations_,
conv_param.input_left_pads_,
conv_param.input_right_pads_,
InElementOp{},
WeiElementOp{},
OutElementOp{},
{},
{},
{});
ref_invoker.Run(ref_argument);
wei_device_buf.FromDevice(wei_device_result.mData.data());
return ck::utils::check_err(wei_device_result.mData, wei_host_result.mData);
}
float avg_time = invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
std::size_t flop = conv_param.GetFlops();
std::size_t num_btype = conv_param.GetByte<InDataType, WeiDataType, OutDataType>();
float tflops = static_cast<float>(flop) / 1.E9 / avg_time;
float gb_per_sec = num_btype / 1.E6 / avg_time;
std::cerr << "Perf: " << avg_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s"
<< std::endl
<< "DeviceOp: " << conv.GetTypeString() << std::endl;
#endif
return true;
}