Skip to content

NHWC conv 2d: fwd bfp16/int8, Device level tuning and host API #73

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
asroy merged 18 commits into from
Feb 12, 2022
Merged

NHWC conv 2d: fwd bfp16/int8, Device level tuning and host API #73

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
18 commits
Select commit Hold shift + click to select a range
4297ee1
add fwd bf16 conv
ltqin Jan 19, 2022
7a11971
change tunning parametor
ltqin Jan 19, 2022
3b6d956
add int8 for conv fwd
ltqin Jan 22, 2022
7c0da10
remove comments
ltqin Jan 22, 2022
8438d93
change tunning parametor for int8
ltqin Jan 22, 2022
4fb8be7
change init int8 example
ltqin Jan 24, 2022
01f13b7
add test for conv2d fwd
ltqin Jan 26, 2022
d49e71c
Merge branch 'develop' into ck_conv_fwd_bf16
ltqin Feb 7, 2022
525cb46
change device operation file pos because merge develop
ltqin Feb 7, 2022
feda990
fwd int8 use reference
ltqin Feb 8, 2022
2647dd6
test_conv_fwd use reference
ltqin Feb 8, 2022
430d5c2
add braket for if statement
ltqin Feb 10, 2022
a6eca54
Merge branch 'develop' into ck_conv_fwd_bf16
ltqin Feb 11, 2022
d9a1f3d
rename fwd example name
ltqin Feb 11, 2022
47fb3d8
Merge remote-tracking branch 'origin/develop' into ck_conv_fwd_bf16
Feb 11, 2022
206d5e5
remove StaticBufferOfVectorTypeV2
Feb 11, 2022
dfb8024
Merge remote-tracking branch 'origin/develop' into ck_conv_fwd_bf16
Feb 11, 2022
6af131f
tweak example
Feb 12, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions composable_kernel/include/tensor_operation/element_wise_operation.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,12 @@ struct PassThrough
__host__ __device__ void operator()(float& y, const float& x) const { y = x; }

__host__ __device__ void operator()(half_t& y, const half_t& x) const { y = x; }

__host__ __device__ void operator()(ushort& y, const ushort& x) const { y = x; }

__host__ __device__ void operator()(int32_t& y, const int32_t& x) const { y = x; }

__host__ __device__ void operator()(int8_t& y, const int8_t& x) const { y = x; }
};

struct Add
Expand Down
2 changes: 0 additions & 2 deletions composable_kernel/include/utility/common_header.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,8 +30,6 @@
#include "amd_address_space.hpp"
#include "amd_buffer_addressing.hpp"
#include "static_buffer.hpp"
// TODO remove this
#include "static_buffer_of_vector_type_v2.hpp"
#include "dynamic_buffer.hpp"
#include "is_known_at_compile_time.hpp"
#include "transpose_vectors.hpp"
Expand Down
10 changes: 10 additions & 0 deletions composable_kernel/include/utility/dynamic_buffer.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -171,6 +171,8 @@ struct DynamicBuffer
is_same<remove_cvref_t<X>, int8x4_t>::value) ||
(is_same<remove_cvref_t<T>, int8_t>::value &&
is_same<remove_cvref_t<X>, int8x8_t>::value) ||
(is_same<remove_cvref_t<T>, int8_t>::value &&
is_same<remove_cvref_t<X>, int8x16_t>::value) ||
(is_same<remove_cvref_t<T>, int8x4_t>::value &&
is_same<remove_cvref_t<X>, int8x4_t>::value) ||
(is_same<remove_cvref_t<T>, int8x8_t>::value &&
Expand Down Expand Up @@ -212,6 +214,14 @@ struct DynamicBuffer
*c_style_pointer_cast<int32x2_t*>(&p_data_[i]) =
*c_style_pointer_cast<const int32x2_t*>(&x);
}
else if constexpr(is_same<remove_cvref_t<T>, int8_t>::value &&
is_same<remove_cvref_t<X>, int8x16_t>::value)
{
// HACK: cast pointer of x is bad
// TODO: remove this after compiler fix
*c_style_pointer_cast<int32x4_t*>(&p_data_[i]) =
*c_style_pointer_cast<const int32x4_t*>(&x);
}
else if constexpr(is_same<remove_cvref_t<T>, int8x4_t>::value &&
is_same<remove_cvref_t<X>, int8x4_t>::value)
{
Expand Down
118 changes: 0 additions & 118 deletions composable_kernel/include/utility/static_buffer_of_vector_type_v2.hpp
View file
Open in desktop

This file was deleted.

2 changes: 2 additions & 0 deletions device_operation/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,8 @@ set(DEVICE_BATCHED_GEMM_INSTANCE_SOURCE
set(DEVICE_CONV2D_FWD_INSTANCE_SOURCE
${PROJECT_SOURCE_DIR}/device_operation/src/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_f32_instance.cpp;
${PROJECT_SOURCE_DIR}/device_operation/src/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_f16_instance.cpp;
${PROJECT_SOURCE_DIR}/device_operation/src/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_bf16_instance.cpp;
${PROJECT_SOURCE_DIR}/device_operation/src/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk_int8_instance.cpp;
${PROJECT_SOURCE_DIR}/device_operation/src/device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk_f16_instance.cpp;
)

Expand Down
Loading