diff --git a/include/cutlass/platform/platform.h b/include/cutlass/platform/platform.h index 46e1adf483..bf5c4336b4 100644 --- a/include/cutlass/platform/platform.h +++ b/include/cutlass/platform/platform.h @@ -911,6 +911,8 @@ struct numeric_limits { static constexpr float infinity() noexcept { return bit_cast(0x7f800000);} CUTLASS_HOST_DEVICE static constexpr float max() noexcept { return bit_cast(0x7f7fffff);} + CUTLASS_HOST_DEVICE + static constexpr float lowest() noexcept { return -max();} static constexpr bool is_integer = false; static constexpr bool has_infinity = true; }; diff --git a/test/unit/core/functional.cu b/test/unit/core/functional.cu index 607c1e36a8..45013ea269 100644 --- a/test/unit/core/functional.cu +++ b/test/unit/core/functional.cu @@ -36,6 +36,7 @@ #include "cutlass/functional.h" #include "cutlass/core_io.h" +#include "cutlass/epilogue/thread/linear_combination_clamp.h" #include "cutlass/layout/matrix.h" #include "cutlass/util/host_tensor.h" @@ -428,6 +429,31 @@ TEST(Functional, multiply_add_bf16x17) { ///////////////////////////////////////////////////////////////////////////////////////////////// +TEST(Functional, linear_combination_clamp_float) { + using Fragment = cutlass::Array; + using Operator = + cutlass::epilogue::thread::LinearCombinationClamp; + using Tensor = cutlass::HostTensor; + + Fragment accumulator; + accumulator[0] = 2.0f; + EXPECT_EQ(Operator()(accumulator)[0], 2.0f); + + Tensor output({1, 1}); + Tensor input({1, 1}); + input.host_data()[0] = 2.0f; + input.sync_device(); + + test::core::kernel::unary_operator<<<1, 1>>>( + reinterpret_cast(output.device_data()), + reinterpret_cast(input.device_data())); + output.sync_host(); + + EXPECT_EQ(output.host_data()[0], 2.0f); +} + +///////////////////////////////////////////////////////////////////////////////////////////////// + template cutlass::Quaternion random_quaternion(int range) { return cutlass::Quaternion{