"...composable_kernel_rocm.git" did not exist on "417f805fb0f161e4afc1edc9410bbdb76fe9b0c8"
Commit ca99f301 authored by Andriy Roshchenko's avatar Andriy Roshchenko
Browse files

Refactoring. Move FP8 definitions into a separate header file.

parent 16a38ff6
This diff is collapsed.
This diff is collapsed.
...@@ -3,6 +3,8 @@ ...@@ -3,6 +3,8 @@
#pragma once #pragma once
#include "ck/ck.hpp"
namespace ck { namespace ck {
// Pseudo random number generator // Pseudo random number generator
......
...@@ -23,8 +23,8 @@ TEST(BF8OCP, NumericLimits) ...@@ -23,8 +23,8 @@ TEST(BF8OCP, NumericLimits)
type_convert<bf8_ocp_t>(0x7D).data); // 0b01111101 type_convert<bf8_ocp_t>(0x7D).data); // 0b01111101
EXPECT_FALSE(ck::NumericLimits<bf8_ocp_t>::QuietNaN() == EXPECT_FALSE(ck::NumericLimits<bf8_ocp_t>::QuietNaN() ==
ck::NumericLimits<bf8_ocp_t>::QuietNaN()); ck::NumericLimits<bf8_ocp_t>::QuietNaN());
EXPECT_TRUE(ck::internal::fp8_is_inf(type_convert<bf8_ocp_t>(0xFC)) && EXPECT_TRUE(ck::fp8_impl::fp8_is_inf(type_convert<bf8_ocp_t>(0xFC)) &&
ck::internal::fp8_is_inf(type_convert<bf8_ocp_t>(0x7C))); ck::fp8_impl::fp8_is_inf(type_convert<bf8_ocp_t>(0x7C)));
} }
TEST(BF8OCP, ConvertFP32Nearest) TEST(BF8OCP, ConvertFP32Nearest)
...@@ -79,7 +79,7 @@ TEST(BF8OCP, ConvertFP32Nearest) ...@@ -79,7 +79,7 @@ TEST(BF8OCP, ConvertFP32Nearest)
// convert quiet NaN to bf8_ocp_t and check if it is quiet NaN // convert quiet NaN to bf8_ocp_t and check if it is quiet NaN
const auto bf8_nan = f8_convert_rne<bf8_ocp_t>(std::numeric_limits<float>::quiet_NaN()); const auto bf8_nan = f8_convert_rne<bf8_ocp_t>(std::numeric_limits<float>::quiet_NaN());
ASSERT_TRUE(ck::internal::ocp_bf8_is_nan(bf8_nan.data)); ASSERT_TRUE(ck::fp8_impl::ocp_bf8_is_nan(bf8_nan.data));
} }
TEST(BF8OCP, ConvertFP32Stochastic) TEST(BF8OCP, ConvertFP32Stochastic)
...@@ -136,7 +136,7 @@ TEST(BF8OCP, ConvertFP32Stochastic) ...@@ -136,7 +136,7 @@ TEST(BF8OCP, ConvertFP32Stochastic)
// convert quiet NaN to bf8_ocp_t and check if it is quiet NaN // convert quiet NaN to bf8_ocp_t and check if it is quiet NaN
const auto bf8_nan = f8_convert_sr<bf8_ocp_t>(std::numeric_limits<float>::quiet_NaN()); const auto bf8_nan = f8_convert_sr<bf8_ocp_t>(std::numeric_limits<float>::quiet_NaN());
ASSERT_TRUE(ck::internal::ocp_bf8_is_nan(bf8_nan.data)); ASSERT_TRUE(ck::fp8_impl::ocp_bf8_is_nan(bf8_nan.data));
} }
TEST(BF8OCP, ConvertFP16Nearest) TEST(BF8OCP, ConvertFP16Nearest)
...@@ -199,7 +199,7 @@ TEST(BF8OCP, ConvertFP16Nearest) ...@@ -199,7 +199,7 @@ TEST(BF8OCP, ConvertFP16Nearest)
// convert quiet NaN to bf8_ocp_t and check if it is quiet NaN // convert quiet NaN to bf8_ocp_t and check if it is quiet NaN
const auto bf8_nan = f8_convert_rne<bf8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN()); const auto bf8_nan = f8_convert_rne<bf8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN());
ASSERT_TRUE(ck::internal::ocp_bf8_is_nan(bf8_nan.data)); ASSERT_TRUE(ck::fp8_impl::ocp_bf8_is_nan(bf8_nan.data));
} }
TEST(BF8OCP, ConvertFP16Stochastic) TEST(BF8OCP, ConvertFP16Stochastic)
...@@ -264,5 +264,5 @@ TEST(BF8OCP, ConvertFP16Stochastic) ...@@ -264,5 +264,5 @@ TEST(BF8OCP, ConvertFP16Stochastic)
// convert quiet NaN to bf8_ocp_t and check if it is quiet NaN // convert quiet NaN to bf8_ocp_t and check if it is quiet NaN
const auto bf8_nan = f8_convert_sr<bf8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN()); const auto bf8_nan = f8_convert_sr<bf8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN());
ASSERT_TRUE(ck::internal::ocp_bf8_is_nan(bf8_nan.data)); ASSERT_TRUE(ck::fp8_impl::ocp_bf8_is_nan(bf8_nan.data));
} }
...@@ -184,7 +184,7 @@ TEST(FP8OCP, ConvertFP16Nearest) ...@@ -184,7 +184,7 @@ TEST(FP8OCP, ConvertFP16Nearest)
// convert quiet NaN to f8_ocp_t and check if it is quiet NaN // convert quiet NaN to f8_ocp_t and check if it is quiet NaN
auto f8_nan = f8_convert_rne<f8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN()); auto f8_nan = f8_convert_rne<f8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN());
ASSERT_TRUE(ck::internal::ocp_f8_is_nan(f8_nan.data)); ASSERT_TRUE(ck::fp8_impl::ocp_f8_is_nan(f8_nan.data));
} }
TEST(FP8OCP, ConvertFP16Stochastic) TEST(FP8OCP, ConvertFP16Stochastic)
...@@ -246,5 +246,5 @@ TEST(FP8OCP, ConvertFP16Stochastic) ...@@ -246,5 +246,5 @@ TEST(FP8OCP, ConvertFP16Stochastic)
// convert quiet NaN to f8_ocp_t and check if it is quiet NaN // convert quiet NaN to f8_ocp_t and check if it is quiet NaN
auto f8_nan = f8_convert_sr<f8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN()); auto f8_nan = f8_convert_sr<f8_ocp_t>(ck::NumericLimits<half_t>::QuietNaN());
ASSERT_TRUE(ck::internal::ocp_f8_is_nan(f8_nan.data)); ASSERT_TRUE(ck::fp8_impl::ocp_f8_is_nan(f8_nan.data));
} }
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