ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-11 17:00:18 +00:00
Code Issues Packages Projects Releases Wiki Activity

This commit contains implementation of max pool2d for f8 type (#1506)

Browse Source 
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
This commit is contained in:
Mateusz Ozga
2024-09-16 10:15:06 +02:00
committed by GitHub
parent 844f5a1712
commit 6834e5ee74
4 changed files with 47 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
library/include/ck/library/tensor_operation_instance/gpu/max_pool_bwd.hpp
View File

@@ -23,6 +23,10 @@ void add_device_maxpool_bwd_bf16_instances(
void add_device_maxpool_bwd_f32_instances(
std::vector<std::unique_ptr<DeviceMaxPoolBwd<F32, I32, F32>>>&);
#endif
#ifdef CK_ENABLE_FP8
void add_device_maxpool_bwd_f8_instances(
std::vector<std::unique_ptr<DeviceMaxPoolBwd<F8, I32, F8>>>&);
#endif
#ifdef CK_ENABLE_INT8
void add_device_maxpool_bwd_int8_instances(
std::vector<std::unique_ptr<DeviceMaxPoolBwd<I8, I32, I8>>>&);
@@ -53,6 +57,11 @@ struct DeviceOperationInstanceFactory<
is_same_v<IndexDataType, I32>)
add_device_maxpool_bwd_f32_instances(op_ptrs);
#endif
#ifdef CK_ENABLE_FP8
else if constexpr(is_same_v<DOutDataType, F8> && is_same_v<DInDataType, F8> &&
is_same_v<IndexDataType, I32>)
add_device_maxpool_bwd_f8_instances(op_ptrs);
#endif
#ifdef CK_ENABLE_INT8
else if constexpr(is_same_v<DOutDataType, I8> && is_same_v<DInDataType, I8> &&
is_same_v<IndexDataType, I32>)