mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-07-13 18:51:13 +00:00
fix clang format
This commit is contained in:
@@ -22,13 +22,13 @@ struct filter_tuple_by_modulo
|
||||
|
||||
// Generate filtered indices for this stride and offset.
|
||||
static constexpr int new_size = (std::tuple_size_v<Tuple> + Stride - Offset - 1) / Stride;
|
||||
|
||||
|
||||
template <std::size_t... Is>
|
||||
static constexpr auto to_index(std::index_sequence<Is...>)
|
||||
{
|
||||
return std::index_sequence<(Offset + Is * Stride)...>{};
|
||||
}
|
||||
|
||||
|
||||
using filtered_indices = decltype(to_index(std::make_index_sequence<new_size>{}));
|
||||
|
||||
// Helper struct to construct the new tuple type from the filtered indices.
|
||||
@@ -36,12 +36,12 @@ struct filter_tuple_by_modulo
|
||||
struct make_filtered_tuple_type_impl;
|
||||
|
||||
template <typename T, std::size_t... Is>
|
||||
struct make_filtered_tuple_type_impl<T, std::index_sequence<Is...>> {
|
||||
struct make_filtered_tuple_type_impl<T, std::index_sequence<Is...>>
|
||||
{
|
||||
using type = std::tuple<std::tuple_element_t<Is, T>...>;
|
||||
};
|
||||
|
||||
using type = typename make_filtered_tuple_type_impl<Tuple, filtered_indices>::type;
|
||||
|
||||
};
|
||||
|
||||
// Filter a tuple with a stride and offset.
|
||||
|
||||
@@ -31,35 +31,34 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_comp_instances_s
|
||||
instances,
|
||||
util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
util::filter_tuple_by_modulo_t<device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
}
|
||||
|
||||
@@ -31,35 +31,34 @@ void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_comp_instances_s
|
||||
instances,
|
||||
util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
util::filter_tuple_by_modulo_t<device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_comp_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
}
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 0>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 1>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 2>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 3>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 4>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 5>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 5>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 6>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 7>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances_sharded<8, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 0>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 1>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 2>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 3>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 4>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 5>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 5>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 6>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 7>(
|
||||
template void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances_sharded<8, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 9>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 9>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 1>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 2>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 3>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 4>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 6>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 7>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 8>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances_sharded<10, 8>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 9>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 9>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 1>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 2>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 3>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 4>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 6>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 7>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 8>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances_sharded<10, 8>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_bf16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -27,41 +27,39 @@ template <int Shards, int ShardIndex>
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances)
|
||||
{
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Interwave>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
}
|
||||
|
||||
} // namespace instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 9>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 9>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 1>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 2>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 3>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 4>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 6>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 7>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 8>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances_sharded<10, 8>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -27,41 +27,39 @@ template <int Shards, int ShardIndex>
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances)
|
||||
{
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Intrawave>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f16_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
}
|
||||
|
||||
} // namespace instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 9>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 9>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 1>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 2>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 3>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 4>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 6>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 7>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 8>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances_sharded<10, 8>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f16_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -27,41 +27,39 @@ template <int Shards, int ShardIndex>
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances)
|
||||
{
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Interwave>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Interwave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
}
|
||||
|
||||
} // namespace instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 9>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 9>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 1>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 2>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 3>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 4>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 6>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 7>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 8>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances_sharded<10, 8>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_inter_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -27,41 +27,39 @@ template <int Shards, int ShardIndex>
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances)
|
||||
{
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Intrawave>,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwdDefault,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1P0,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
add_device_operation_instances(
|
||||
instances,
|
||||
ck::util::filter_tuple_by_modulo_t<
|
||||
device_grouped_conv_fwd_xdl_f32_mem_instances<3,
|
||||
NGCDHW,
|
||||
GKCZYX,
|
||||
Empty_Tuple,
|
||||
NGKDHW,
|
||||
ConvFwd1x1S1P0,
|
||||
Intrawave>,
|
||||
Shards,
|
||||
ShardIndex>{});
|
||||
}
|
||||
|
||||
} // namespace instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 9>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 9>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 1>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 1>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 2>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 2>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 3>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 3>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 4>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 4>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 0>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 0>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 6>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 6>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 7>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 7>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
@@ -4,7 +4,7 @@
|
||||
#include "device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instance.inc"
|
||||
|
||||
namespace ck::tensor_operation::device::instance {
|
||||
template
|
||||
void add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 8>(
|
||||
template void
|
||||
add_device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances_sharded<10, 8>(
|
||||
device_grouped_conv3d_fwd_xdl_ngcdhw_gkczyx_ngkdhw_f32_mem_intra_instances& instances);
|
||||
} // namespace ck::tensor_operation::device::instance
|
||||
|
||||
Reference in New Issue
Block a user