ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 02:02:46 +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>

[ROCm/composable_kernel commit: 6834e5ee74]
This commit is contained in:
Mateusz Ozga
2024-09-16 10:15:06 +02:00
committed by GitHub
parent db557ada8b
commit d7326fb525
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>)

1
library/src/tensor_operation_instance/gpu/max_pool_bwd/CMakeLists.txt
View File

@@ -2,5 +2,6 @@ set(DEVICE_MAXPOOL_BWD_INSTANCES)
list(APPEND DEVICE_MAXPOOL_BWD_INSTANCES device_max_pool_bwd_f16_instance.cpp
device_max_pool_bwd_bf16_instance.cpp
device_max_pool_bwd_f32_instance.cpp
device_max_pool_bwd_f8_instance.cpp
device_max_pool_bwd_int8_instance.cpp)
add_instance_library(device_max_pool_bwd_instance ${DEVICE_MAXPOOL_BWD_INSTANCES})

20
library/src/tensor_operation_instance/gpu/max_pool_bwd/device_max_pool_bwd_f8_instance.cpp Normal file
View File

@@ -0,0 +1,20 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "max_pool_bwd_instance_common.hpp"
namespace ck {
namespace tensor_operation {
namespace device {
namespace instance {
void add_device_maxpool_bwd_f8_instances(
std::vector<std::unique_ptr<DeviceMaxPoolBwd<F8, I32, F8>>>& instances)
{
add_device_operation_instances(instances, device_maxpool_bwd_instances<F8, I32, F8>{});
}
} // namespace instance
} // namespace device
} // namespace tensor_operation
} // namespace ck

17
test/pool/test_max_pool2d_bwd.cpp
View File

@@ -55,6 +55,7 @@ using Max_Pool_2D_f32_types = ::testing::Types<std::tuple<F32, F32, I32>>;
using Max_Pool_2D_int8_types = ::testing::Types<std::tuple<I8, I8, I32>>;
using Max_Pool_2D_f16_types = ::testing::Types<std::tuple<F16, F16, I32>>;
using Max_Pool_2D_bf16_types = ::testing::Types<std::tuple<BF16, BF16, I32>>;
using Max_Pool_2D_f8_types = ::testing::Types<std::tuple<F8, F8, I32>>;
template <typename TType>
class MaxPool2D_f32 : public MaxPool2dBWDTest<TType>
@@ -108,10 +109,24 @@ class MaxPool2D_bf16 : public MaxPool2dBWDTest<TType>
}
};
template <typename TType>
class MaxPool2D_f8 : public MaxPool2dBWDTest<TType>
{
protected:
void SetUp() override
{
if(!CK_ENABLE_FP8)
{
GTEST_SKIP() << "Skipping MaxPool2D_f8 tests because CK_ENABLE_FP8 is not enabled";
}
}
};
TYPED_TEST_SUITE(MaxPool2D_f32, Max_Pool_2D_f32_types);
TYPED_TEST_SUITE(MaxPool2D_int8, Max_Pool_2D_int8_types);
TYPED_TEST_SUITE(MaxPool2D_f16, Max_Pool_2D_f16_types);
TYPED_TEST_SUITE(MaxPool2D_bf16, Max_Pool_2D_bf16_types);
TYPED_TEST_SUITE(MaxPool2D_f8, Max_Pool_2D_f8_types);
TYPED_TEST(MaxPool2D_f32, MaxPool2DTest_f32) { this->Run(); }
@@ -120,3 +135,5 @@ TYPED_TEST(MaxPool2D_int8, MaxPool2DTest_int8) { this->Run(); }
TYPED_TEST(MaxPool2D_f16, MaxPool2DTest_f16) { this->Run(); }
TYPED_TEST(MaxPool2D_bf16, MaxPool2DTest_bf16) { this->Run(); }
TYPED_TEST(MaxPool2D_f8, MaxPool2DTest_f8) { this->Run(); }