Skip to content

GitLab

"include/composable_kernel/utility/amd_inline_asm.hpp" did not exist on "fbc7817bbbe4580eb460c78cdffb497a5a226a82"
Commit 159e1354 authored by Jun Liu's avatar Jun Liu
Browse files

Fix more places for workaround

parent 0f58a333
Hide whitespace changes
Inline Side-by-side
Showing with 23 additions and 20 deletions
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scaleadd_ab.hpp
View file @ 159e1354
...@@ -23,19 +23,20 @@ using ScaleAdd = ck::tensor_operation::element_wise::ScaleAdd; ...@@ -23,19 +23,20 @@ using ScaleAdd = ck::tensor_operation::element_wise::ScaleAdd;
#ifdef CK_ENABLE_BF16 #ifdef CK_ENABLE_BF16
// grouped conv3d forward multi AB scaleadd, NDHWGC/GKZYXC/NDHWGK // grouped conv3d forward multi AB scaleadd, NDHWGC/GKZYXC/NDHWGK
void add_device_grouped_conv3d_fwd_xdl_scaleadd_ab_ndhwgc_gkzyxc_ndhwgk_bf16_instances( // TODO: Workaround for https://ontrack-internal.amd.com/browse/SWDEV-435347
std::vector<std::unique_ptr<DeviceGroupedConvFwdMultipleABD<3, // void add_device_grouped_conv3d_fwd_xdl_scaleadd_ab_ndhwgc_gkzyxc_ndhwgk_bf16_instances(
NDHWGC, // std::vector<std::unique_ptr<DeviceGroupedConvFwdMultipleABD<3,
GKZYXC, // NDHWGC,
ck::Tuple<>, // GKZYXC,
NDHWGK, // ck::Tuple<>,
ck::Tuple<BF16, BF16>, // NDHWGK,
ck::Tuple<BF16, BF16>, // ck::Tuple<BF16, BF16>,
ck::Tuple<>, // ck::Tuple<BF16, BF16>,
BF16, // ck::Tuple<>,
ScaleAdd, // BF16,
ScaleAdd, // ScaleAdd,
PassThrough>>>& instances); // ScaleAdd,
// PassThrough>>>& instances);
#endif #endif
#ifdef CK_ENABLE_FP16 #ifdef CK_ENABLE_FP16
...@@ -151,13 +152,15 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupe ...@@ -151,13 +152,15 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupe
} }
#endif #endif
#ifdef CK_ENABLE_BF16 #ifdef CK_ENABLE_BF16
if constexpr(is_same_v<InDataType, ck::Tuple<ck::bhalf_t, ck::bhalf_t>> && // TODO: Workaround for https://ontrack-internal.amd.com/browse/SWDEV-435347
is_same_v<WeiDataType, ck::Tuple<ck::bhalf_t, ck::bhalf_t>> && // if constexpr(is_same_v<InDataType, ck::Tuple<ck::bhalf_t, ck::bhalf_t>> &&
is_same_v<OutDataType, ck::bhalf_t> && is_same_v<ComputeType, ck::bhalf_t>) // is_same_v<WeiDataType, ck::Tuple<ck::bhalf_t, ck::bhalf_t>> &&
{ // is_same_v<OutDataType, ck::bhalf_t> && is_same_v<ComputeType,
add_device_grouped_conv3d_fwd_xdl_scaleadd_ab_ndhwgc_gkzyxc_ndhwgk_bf16_instances( // ck::bhalf_t>)
op_ptrs); // {
} // add_device_grouped_conv3d_fwd_xdl_scaleadd_ab_ndhwgc_gkzyxc_ndhwgk_bf16_instances(
// op_ptrs);
// }
#endif #endif
#ifdef CK_ENABLE_INT8 #ifdef CK_ENABLE_INT8
if constexpr(is_same_v<InDataType, ck::Tuple<int8_t, int8_t>> && if constexpr(is_same_v<InDataType, ck::Tuple<int8_t, int8_t>> &&
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment