wgsl: Print abstract-floats with full precision.
Change-Id: Ie95547f065b896983b90ffd5455538fdd843b81a Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/104824 Reviewed-by: David Neto <dneto@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
parent
9f513ca541
commit
d6daefc379
|
@ -25,9 +25,34 @@
|
|||
|
||||
namespace tint::writer {
|
||||
|
||||
std::string FloatToString(float f) {
|
||||
// Try printing the float in fixed point, with a smallish limit on the
|
||||
// precision
|
||||
namespace {
|
||||
|
||||
template <typename T>
|
||||
struct Traits;
|
||||
|
||||
template <>
|
||||
struct Traits<float> {
|
||||
using uint_t = uint32_t;
|
||||
static constexpr int kExponentBias = 127;
|
||||
static constexpr uint_t kExponentMask = 0x7f800000;
|
||||
static constexpr uint_t kMantissaMask = 0x007fffff;
|
||||
static constexpr uint_t kSignMask = 0x80000000;
|
||||
static constexpr int kMantissaBits = 23;
|
||||
};
|
||||
|
||||
template <>
|
||||
struct Traits<double> {
|
||||
using uint_t = uint64_t;
|
||||
static constexpr int kExponentBias = 1023;
|
||||
static constexpr uint_t kExponentMask = 0x7ff0000000000000;
|
||||
static constexpr uint_t kMantissaMask = 0x000fffffffffffff;
|
||||
static constexpr uint_t kSignMask = 0x8000000000000000;
|
||||
static constexpr int kMantissaBits = 52;
|
||||
};
|
||||
|
||||
template <typename F>
|
||||
std::string ToString(F f) {
|
||||
// Try printing the float in fixed point, with a smallish limit on the precision
|
||||
std::stringstream fixed;
|
||||
fixed.flags(fixed.flags() | std::ios_base::showpoint | std::ios_base::fixed);
|
||||
fixed.imbue(std::locale::classic());
|
||||
|
@ -36,13 +61,13 @@ std::string FloatToString(float f) {
|
|||
std::string str = fixed.str();
|
||||
|
||||
// If this string can be parsed without loss of information, use it.
|
||||
// (Use double here to dodge a bug in older libc++ versions which
|
||||
// would incorrectly read back FLT_MAX as INF.)
|
||||
// (Use double here to dodge a bug in older libc++ versions which would incorrectly read back
|
||||
// FLT_MAX as INF.)
|
||||
double roundtripped;
|
||||
fixed >> roundtripped;
|
||||
|
||||
auto float_equal_no_warning = std::equal_to<float>();
|
||||
if (float_equal_no_warning(f, static_cast<float>(roundtripped))) {
|
||||
auto float_equal_no_warning = std::equal_to<F>();
|
||||
if (float_equal_no_warning(f, static_cast<F>(roundtripped))) {
|
||||
while (str.length() >= 2 && str[str.size() - 1] == '0' && str[str.size() - 2] != '.') {
|
||||
str.pop_back();
|
||||
}
|
||||
|
@ -50,38 +75,41 @@ std::string FloatToString(float f) {
|
|||
return str;
|
||||
}
|
||||
|
||||
// Resort to scientific, with the minimum precision needed to preserve the
|
||||
// whole float
|
||||
// Resort to scientific, with the minimum precision needed to preserve the whole float
|
||||
std::stringstream sci;
|
||||
sci.imbue(std::locale::classic());
|
||||
sci.precision(std::numeric_limits<float>::max_digits10);
|
||||
sci.precision(std::numeric_limits<F>::max_digits10);
|
||||
sci << f;
|
||||
return sci.str();
|
||||
}
|
||||
|
||||
std::string FloatToBitPreservingString(float f) {
|
||||
template <typename F>
|
||||
std::string ToBitPreservingString(F f) {
|
||||
using T = Traits<F>;
|
||||
using uint_t = typename T::uint_t;
|
||||
|
||||
// For the NaN case, avoid handling the number as a floating point value.
|
||||
// Some machines will modify the top bit in the mantissa of a NaN.
|
||||
|
||||
std::stringstream ss;
|
||||
|
||||
uint32_t float_bits = 0u;
|
||||
typename T::uint_t float_bits = 0u;
|
||||
static_assert(sizeof(float_bits) == sizeof(f));
|
||||
std::memcpy(&float_bits, &f, sizeof(float_bits));
|
||||
|
||||
// Handle the sign.
|
||||
const uint32_t kSignMask = 1u << 31;
|
||||
if (float_bits & kSignMask) {
|
||||
if (float_bits & T::kSignMask) {
|
||||
// If `f` is -0.0 print -0.0.
|
||||
ss << '-';
|
||||
// Strip sign bit.
|
||||
float_bits = float_bits & (~kSignMask);
|
||||
float_bits = float_bits & (~T::kSignMask);
|
||||
}
|
||||
|
||||
switch (std::fpclassify(f)) {
|
||||
case FP_ZERO:
|
||||
case FP_NORMAL:
|
||||
std::memcpy(&f, &float_bits, sizeof(float_bits));
|
||||
ss << FloatToString(f);
|
||||
ss << ToString(f);
|
||||
break;
|
||||
|
||||
default: {
|
||||
|
@ -89,46 +117,39 @@ std::string FloatToBitPreservingString(float f) {
|
|||
// TODO(dneto): It's unclear how Infinity and NaN should be handled.
|
||||
// See https://github.com/gpuweb/gpuweb/issues/1769
|
||||
|
||||
// std::hexfloat prints 'nan' and 'inf' instead of an
|
||||
// explicit representation like we want. Split it out
|
||||
// manually.
|
||||
const int kExponentBias = 127;
|
||||
const int kExponentMask = 0x7f800000;
|
||||
const int kMantissaMask = 0x007fffff;
|
||||
const int kMantissaBits = 23;
|
||||
|
||||
int mantissaNibbles = (kMantissaBits + 3) / 4;
|
||||
// std::hexfloat prints 'nan' and 'inf' instead of an explicit representation like we
|
||||
// want. Split it out manually.
|
||||
int mantissa_nibbles = (T::kMantissaBits + 3) / 4;
|
||||
|
||||
const int biased_exponent =
|
||||
static_cast<int>((float_bits & kExponentMask) >> kMantissaBits);
|
||||
int exponent = biased_exponent - kExponentBias;
|
||||
uint32_t mantissa = float_bits & kMantissaMask;
|
||||
static_cast<int>((float_bits & T::kExponentMask) >> T::kMantissaBits);
|
||||
int exponent = biased_exponent - T::kExponentBias;
|
||||
uint_t mantissa = float_bits & T::kMantissaMask;
|
||||
|
||||
ss << "0x";
|
||||
|
||||
if (exponent == 128) {
|
||||
if (exponent == T::kExponentBias + 1) {
|
||||
if (mantissa == 0) {
|
||||
// Infinity case.
|
||||
ss << "1p+128";
|
||||
ss << "1p+" << exponent;
|
||||
} else {
|
||||
// NaN case.
|
||||
// Emit the mantissa bits as if they are left-justified after the
|
||||
// binary point. This is what SPIRV-Tools hex float emitter does,
|
||||
// and it's a justifiable choice independent of the bit width
|
||||
// of the mantissa.
|
||||
mantissa <<= (4 - (kMantissaBits % 4));
|
||||
// Remove trailing zeroes, for tidyness.
|
||||
// NaN case.
|
||||
// Emit the mantissa bits as if they are left-justified after the binary point.
|
||||
// This is what SPIRV-Tools hex float emitter does, and it's a justifiable
|
||||
// choice independent of the bit width of the mantissa.
|
||||
mantissa <<= (4 - (T::kMantissaBits % 4));
|
||||
// Remove trailing zeroes, for tidiness.
|
||||
while (0 == (0xf & mantissa)) {
|
||||
mantissa >>= 4;
|
||||
mantissaNibbles--;
|
||||
mantissa_nibbles--;
|
||||
}
|
||||
ss << "1." << std::hex << std::setfill('0') << std::setw(mantissaNibbles)
|
||||
<< mantissa << "p+128";
|
||||
ss << "1." << std::hex << std::setfill('0') << std::setw(mantissa_nibbles)
|
||||
<< mantissa << "p+" << std::dec << exponent;
|
||||
}
|
||||
} else {
|
||||
// Subnormal, and not zero.
|
||||
TINT_ASSERT(Writer, mantissa != 0);
|
||||
const int kTopBit = (1 << kMantissaBits);
|
||||
const auto kTopBit = static_cast<uint_t>(1u) << T::kMantissaBits;
|
||||
|
||||
// Shift left until we get 1.x
|
||||
while (0 == (kTopBit & mantissa)) {
|
||||
|
@ -138,17 +159,19 @@ std::string FloatToBitPreservingString(float f) {
|
|||
// Emit the leading 1, and remove it from the mantissa.
|
||||
ss << "1";
|
||||
mantissa = mantissa ^ kTopBit;
|
||||
mantissa <<= 1;
|
||||
exponent++;
|
||||
|
||||
// Left-justify mantissa to whole nibble.
|
||||
mantissa <<= (4 - (T::kMantissaBits % 4));
|
||||
|
||||
// Emit the fractional part.
|
||||
if (mantissa) {
|
||||
// Remove trailing zeroes, for tidyness
|
||||
// Remove trailing zeroes, for tidiness
|
||||
while (0 == (0xf & mantissa)) {
|
||||
mantissa >>= 4;
|
||||
mantissaNibbles--;
|
||||
mantissa_nibbles--;
|
||||
}
|
||||
ss << "." << std::hex << std::setfill('0') << std::setw(mantissaNibbles)
|
||||
ss << "." << std::hex << std::setfill('0') << std::setw(mantissa_nibbles)
|
||||
<< mantissa;
|
||||
}
|
||||
// Emit the exponent
|
||||
|
@ -159,4 +182,22 @@ std::string FloatToBitPreservingString(float f) {
|
|||
return ss.str();
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
std::string FloatToString(float f) {
|
||||
return ToString(f);
|
||||
}
|
||||
|
||||
std::string FloatToBitPreservingString(float f) {
|
||||
return ToBitPreservingString(f);
|
||||
}
|
||||
|
||||
std::string DoubleToString(double f) {
|
||||
return ToString(f);
|
||||
}
|
||||
|
||||
std::string DoubleToBitPreservingString(double f) {
|
||||
return ToBitPreservingString(f);
|
||||
}
|
||||
|
||||
} // namespace tint::writer
|
||||
|
|
|
@ -27,11 +27,24 @@ namespace tint::writer {
|
|||
/// @return the float f formatted to a string
|
||||
std::string FloatToString(float f);
|
||||
|
||||
/// Converts the double `f` to a string using fixed-point notation (not
|
||||
/// scientific). The double will be printed with the full precision required to
|
||||
/// describe the double. All trailing `0`s will be omitted after the last
|
||||
/// non-zero fractional number, unless the fractional is zero, in which case the
|
||||
/// number will end with `.0`.
|
||||
/// @return the double f formatted to a string
|
||||
std::string DoubleToString(double f);
|
||||
|
||||
/// Converts the float `f` to a string, using hex float notation for infinities,
|
||||
/// NaNs, or subnormal numbers. Otherwise behaves as FloatToString.
|
||||
/// @return the float f formatted to a string
|
||||
std::string FloatToBitPreservingString(float f);
|
||||
|
||||
/// Converts the double `f` to a string, using hex double notation for infinities,
|
||||
/// NaNs, or subnormal numbers. Otherwise behaves as FloatToString.
|
||||
/// @return the double f formatted to a string
|
||||
std::string DoubleToBitPreservingString(double f);
|
||||
|
||||
} // namespace tint::writer
|
||||
|
||||
#endif // SRC_TINT_WRITER_FLOAT_TO_STRING_H_
|
||||
|
|
|
@ -14,33 +14,19 @@
|
|||
|
||||
#include "src/tint/writer/float_to_string.h"
|
||||
|
||||
#include <cmath>
|
||||
#include <math.h>
|
||||
#include <cstring>
|
||||
#include <limits>
|
||||
|
||||
#include "gtest/gtest.h"
|
||||
#include "src/tint/utils/bitcast.h"
|
||||
|
||||
namespace tint::writer {
|
||||
namespace {
|
||||
|
||||
// Makes an IEEE 754 binary32 floating point number with
|
||||
// - 0 sign if sign is 0, 1 otherwise
|
||||
// - 'exponent_bits' is placed in the exponent space.
|
||||
// So, the exponent bias must already be included.
|
||||
float MakeFloat(uint32_t sign, uint32_t biased_exponent, uint32_t mantissa) {
|
||||
const uint32_t sign_bit = sign ? 0x80000000u : 0u;
|
||||
// The binary32 exponent is 8 bits, just below the sign.
|
||||
const uint32_t exponent_bits = (biased_exponent & 0xffu) << 23;
|
||||
// The mantissa is the bottom 23 bits.
|
||||
const uint32_t mantissa_bits = (mantissa & 0x7fffffu);
|
||||
|
||||
uint32_t bits = sign_bit | exponent_bits | mantissa_bits;
|
||||
float result = 0.0f;
|
||||
static_assert(sizeof(result) == sizeof(bits),
|
||||
"expected float and uint32_t to be the same size");
|
||||
std::memcpy(&result, &bits, sizeof(bits));
|
||||
return result;
|
||||
}
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// FloatToString //
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
TEST(FloatToStringTest, Zero) {
|
||||
EXPECT_EQ(FloatToString(0.0f), "0.0");
|
||||
|
@ -93,14 +79,18 @@ TEST(FloatToStringTest, Precision) {
|
|||
EXPECT_EQ(FloatToString(1e-20f), "9.99999968e-21");
|
||||
}
|
||||
|
||||
// FloatToBitPreservingString
|
||||
//
|
||||
// First replicate the tests for FloatToString
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// FloatToBitPreservingString //
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, Zero) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(0.0f), "0.0");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, NegativeZero) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(-0.0f), "-0.0");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, One) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(1.0f), "1.0");
|
||||
}
|
||||
|
@ -141,49 +131,204 @@ TEST(FloatToBitPreservingStringTest, Lowest) {
|
|||
"-340282346638528859811704183484516925440.0");
|
||||
}
|
||||
|
||||
// Special cases for bit-preserving output.
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, NegativeZero) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(std::copysign(0.0f, -5.0f)), "-0.0");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, ZeroAsBits) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 0, 0)), "0.0");
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(1, 0, 0)), "-0.0");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, OneBits) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 127, 0)), "1.0");
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(1, 127, 0)), "-1.0");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, SmallestDenormal) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 0, 1)), "0x1p-149");
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(1, 0, 1)), "-0x1p-149");
|
||||
EXPECT_EQ(FloatToBitPreservingString(0x1p-149f), "0x1p-149");
|
||||
EXPECT_EQ(FloatToBitPreservingString(-0x1p-149f), "-0x1p-149");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, BiggerDenormal) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 0, 2)), "0x1p-148");
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(1, 0, 2)), "-0x1p-148");
|
||||
EXPECT_EQ(FloatToBitPreservingString(0x1p-148f), "0x1p-148");
|
||||
EXPECT_EQ(FloatToBitPreservingString(-0x1p-148f), "-0x1p-148");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, LargestDenormal) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 0, 0x7fffff)), "0x1.fffffcp-127");
|
||||
static_assert(0x0.fffffep-126f == 0x1.fffffcp-127f);
|
||||
EXPECT_EQ(FloatToBitPreservingString(0x0.fffffep-126f), "0x1.fffffcp-127");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, Subnormal_cafebe) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 0, 0xcafebe)), "0x1.2bfaf8p-127");
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(1, 0, 0xcafebe)), "-0x1.2bfaf8p-127");
|
||||
EXPECT_EQ(FloatToBitPreservingString(0x1.2bfaf8p-127f), "0x1.2bfaf8p-127");
|
||||
EXPECT_EQ(FloatToBitPreservingString(-0x1.2bfaf8p-127f), "-0x1.2bfaf8p-127");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, Subnormal_aaaaa) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 0, 0xaaaaa)), "0x1.55554p-130");
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(1, 0, 0xaaaaa)), "-0x1.55554p-130");
|
||||
EXPECT_EQ(FloatToBitPreservingString(0x1.55554p-130f), "0x1.55554p-130");
|
||||
EXPECT_EQ(FloatToBitPreservingString(-0x1.55554p-130f), "-0x1.55554p-130");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, Infinity) {
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(0, 255, 0)), "0x1p+128");
|
||||
EXPECT_EQ(FloatToBitPreservingString(MakeFloat(1, 255, 0)), "-0x1p+128");
|
||||
EXPECT_EQ(FloatToBitPreservingString(INFINITY), "0x1p+128");
|
||||
EXPECT_EQ(FloatToBitPreservingString(-INFINITY), "-0x1p+128");
|
||||
}
|
||||
|
||||
TEST(FloatToBitPreservingStringTest, NaN) {
|
||||
// TODO(crbug.com/tint/1714): On x86, this bitcast will set bit 22 (the highest mantissa bit) to
|
||||
// 1, regardless of the bit value in the integer. This is likely due to IEEE 754's
|
||||
// recommendation that that the highest mantissa bit differentiates quiet NaNs from signalling
|
||||
// NaNs. On x86, float return values usually go via the FPU which can transform the signalling
|
||||
// NaN bit (0) to quiet NaN (1). As NaN floating point numbers can be silently modified by the
|
||||
// architecture, and the signalling bit is architecture defined, this test may fail on other
|
||||
// architectures.
|
||||
auto nan = utils::Bitcast<float>(0x7fc0beef);
|
||||
EXPECT_EQ(FloatToBitPreservingString(nan), "0x1.817ddep+128");
|
||||
EXPECT_EQ(FloatToBitPreservingString(-nan), "-0x1.817ddep+128");
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// DoubleToString //
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
TEST(DoubleToStringTest, Zero) {
|
||||
EXPECT_EQ(DoubleToString(0.0), "0.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToStringTest, One) {
|
||||
EXPECT_EQ(DoubleToString(1.0), "1.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToStringTest, MinusOne) {
|
||||
EXPECT_EQ(DoubleToString(-1.0), "-1.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToStringTest, Billion) {
|
||||
EXPECT_EQ(DoubleToString(1e9), "1000000000.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToStringTest, Small) {
|
||||
EXPECT_NE(DoubleToString(std::numeric_limits<double>::epsilon()), "0.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToStringTest, Highest) {
|
||||
const auto highest = std::numeric_limits<double>::max();
|
||||
const auto expected_highest = 1.797693134862315708e+308;
|
||||
if (highest < expected_highest || highest > expected_highest) {
|
||||
GTEST_SKIP() << "std::numeric_limits<double>::max() is not as expected for "
|
||||
"this target";
|
||||
}
|
||||
EXPECT_EQ(DoubleToString(std::numeric_limits<double>::max()),
|
||||
"179769313486231570814527423731704356798070567525844996598917476803157260780028538760"
|
||||
"589558632766878171540458953514382464234321326889464182768467546703537516986049910576"
|
||||
"551282076245490090389328944075868508455133942304583236903222948165808559332123348274"
|
||||
"797826204144723168738177180919299881250404026184124858368.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToStringTest, Lowest) {
|
||||
// Some compilers complain if you test floating point numbers for equality.
|
||||
// So say it via two inequalities.
|
||||
const auto lowest = std::numeric_limits<double>::lowest();
|
||||
const auto expected_lowest = -1.797693134862315708e+308;
|
||||
if (lowest < expected_lowest || lowest > expected_lowest) {
|
||||
GTEST_SKIP() << "std::numeric_limits<double>::lowest() is not as expected for "
|
||||
"this target";
|
||||
}
|
||||
EXPECT_EQ(DoubleToString(std::numeric_limits<double>::lowest()),
|
||||
"-17976931348623157081452742373170435679807056752584499659891747680315726078002853876"
|
||||
"058955863276687817154045895351438246423432132688946418276846754670353751698604991057"
|
||||
"655128207624549009038932894407586850845513394230458323690322294816580855933212334827"
|
||||
"4797826204144723168738177180919299881250404026184124858368.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToStringTest, Precision) {
|
||||
EXPECT_EQ(DoubleToString(1e-8), "0.00000001");
|
||||
EXPECT_EQ(DoubleToString(1e-9), "0.000000001");
|
||||
EXPECT_EQ(DoubleToString(1e-10), "1e-10");
|
||||
EXPECT_EQ(DoubleToString(1e-15), "1.0000000000000001e-15");
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// DoubleToBitPreservingString //
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Zero) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(0.0), "0.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, NegativeZero) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(-0.0), "-0.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, One) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(1.0), "1.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, MinusOne) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(-1.0), "-1.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Billion) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(1e9), "1000000000.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Small) {
|
||||
EXPECT_NE(DoubleToBitPreservingString(std::numeric_limits<double>::epsilon()), "0.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Highest) {
|
||||
const auto highest = std::numeric_limits<double>::max();
|
||||
const auto expected_highest = 1.797693134862315708e+308;
|
||||
if (highest < expected_highest || highest > expected_highest) {
|
||||
GTEST_SKIP() << "std::numeric_limits<float>::max() is not as expected for "
|
||||
"this target";
|
||||
}
|
||||
EXPECT_EQ(DoubleToBitPreservingString(std::numeric_limits<double>::max()),
|
||||
"179769313486231570814527423731704356798070567525844996598917476803157260780028538760"
|
||||
"589558632766878171540458953514382464234321326889464182768467546703537516986049910576"
|
||||
"551282076245490090389328944075868508455133942304583236903222948165808559332123348274"
|
||||
"797826204144723168738177180919299881250404026184124858368.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Lowest) {
|
||||
// Some compilers complain if you test floating point numbers for equality.
|
||||
// So say it via two inequalities.
|
||||
const auto lowest = std::numeric_limits<double>::lowest();
|
||||
const auto expected_lowest = -1.797693134862315708e+308;
|
||||
if (lowest < expected_lowest || lowest > expected_lowest) {
|
||||
GTEST_SKIP() << "std::numeric_limits<float>::lowest() is not as expected for "
|
||||
"this target";
|
||||
}
|
||||
EXPECT_EQ(DoubleToBitPreservingString(std::numeric_limits<double>::lowest()),
|
||||
"-17976931348623157081452742373170435679807056752584499659891747680315726078002853876"
|
||||
"058955863276687817154045895351438246423432132688946418276846754670353751698604991057"
|
||||
"655128207624549009038932894407586850845513394230458323690322294816580855933212334827"
|
||||
"4797826204144723168738177180919299881250404026184124858368.0");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, SmallestDenormal) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(0x1p-1074), "0x1p-1074");
|
||||
EXPECT_EQ(DoubleToBitPreservingString(-0x1p-1074), "-0x1p-1074");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, BiggerDenormal) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(0x1p-1073), "0x1p-1073");
|
||||
EXPECT_EQ(DoubleToBitPreservingString(-0x1p-1073), "-0x1p-1073");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, LargestDenormal) {
|
||||
static_assert(0x0.fffffffffffffp-1022 == 0x1.ffffffffffffep-1023);
|
||||
EXPECT_EQ(DoubleToBitPreservingString(0x0.fffffffffffffp-1022), "0x1.ffffffffffffep-1023");
|
||||
EXPECT_EQ(DoubleToBitPreservingString(-0x0.fffffffffffffp-1022), "-0x1.ffffffffffffep-1023");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Subnormal_cafef00dbeef) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(0x1.cafef00dbeefp-1023), "0x1.cafef00dbeefp-1023");
|
||||
EXPECT_EQ(DoubleToBitPreservingString(-0x1.cafef00dbeefp-1023), "-0x1.cafef00dbeefp-1023");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Subnormal_aaaaaaaaaaaaap) {
|
||||
static_assert(0x0.aaaaaaaaaaaaap-1023 == 0x1.5555555555554p-1024);
|
||||
EXPECT_EQ(DoubleToBitPreservingString(0x0.aaaaaaaaaaaaap-1023), "0x1.5555555555554p-1024");
|
||||
EXPECT_EQ(DoubleToBitPreservingString(-0x0.aaaaaaaaaaaaap-1023), "-0x1.5555555555554p-1024");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, Infinity) {
|
||||
EXPECT_EQ(DoubleToBitPreservingString(static_cast<double>(INFINITY)), "0x1p+1024");
|
||||
EXPECT_EQ(DoubleToBitPreservingString(static_cast<double>(-INFINITY)), "-0x1p+1024");
|
||||
}
|
||||
|
||||
TEST(DoubleToBitPreservingStringTest, NaN) {
|
||||
auto nan = utils::Bitcast<double>(0x7ff8cafef00dbeefull);
|
||||
EXPECT_EQ(DoubleToBitPreservingString(static_cast<double>(nan)), "0x1.8cafef00dbeefp+1024");
|
||||
EXPECT_EQ(DoubleToBitPreservingString(static_cast<double>(-nan)), "-0x1.8cafef00dbeefp+1024");
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
|
|
@ -263,7 +263,11 @@ bool GeneratorImpl::EmitLiteral(std::ostream& out, const ast::LiteralExpression*
|
|||
// Note that all normal and subnormal f16 values are normal f32 values, and since NaN
|
||||
// and Inf are not allowed to be spelled in literal, it should be fine to emit f16
|
||||
// literals in this way.
|
||||
out << FloatToBitPreservingString(static_cast<float>(l->value)) << l->suffix;
|
||||
if (l->suffix == ast::FloatLiteralExpression::Suffix::kNone) {
|
||||
out << DoubleToBitPreservingString(l->value);
|
||||
} else {
|
||||
out << FloatToBitPreservingString(static_cast<float>(l->value)) << l->suffix;
|
||||
}
|
||||
return true;
|
||||
},
|
||||
[&](const ast::IntLiteralExpression* l) { //
|
||||
|
|
|
@ -1,18 +1,18 @@
|
|||
fn original_clusterfuzz_code() {
|
||||
atan2(1, 0.100000001);
|
||||
atan2(1, 0.1);
|
||||
}
|
||||
|
||||
fn more_tests_that_would_fail() {
|
||||
{
|
||||
let a = atan2(1, 0.100000001);
|
||||
let b = atan2(0.100000001, 1);
|
||||
let a = atan2(1, 0.1);
|
||||
let b = atan2(0.1, 1);
|
||||
}
|
||||
{
|
||||
let a = (1 + 1.5);
|
||||
let b = (1.5 + 1);
|
||||
}
|
||||
{
|
||||
atan2(1, 0.100000001);
|
||||
atan2(0.100000001, 1);
|
||||
atan2(1, 0.1);
|
||||
atan2(0.1, 1);
|
||||
}
|
||||
}
|
||||
|
|
|
@ -1,6 +1,8 @@
|
|||
SKIP: FAILED
|
||||
|
||||
bug/chromium/1367602-1.wgsl:2:23 error: array size (65536) must be less than 65536
|
||||
var v = array<bool, 65536>();
|
||||
^^^^^
|
||||
[numthreads(1, 1, 1)]
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
|
||||
void f() {
|
||||
bool v[65535] = (bool[65535])0;
|
||||
}
|
||||
|
|
|
@ -48,7 +48,7 @@ fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
|
|||
if ((index >= config.numLights)) {
|
||||
return;
|
||||
}
|
||||
lightsBuffer.lights[index].position.y = ((lightsBuffer.lights[index].position.y - 0.100000001) + (0.001 * (f32(index) - (64.0 * floor((f32(index) / 64.0))))));
|
||||
lightsBuffer.lights[index].position.y = ((lightsBuffer.lights[index].position.y - 0.1) + (0.001 * (f32(index) - (64.0 * floor((f32(index) / 64.0))))));
|
||||
if ((lightsBuffer.lights[index].position.y < uniforms.min.y)) {
|
||||
lightsBuffer.lights[index].position.y = uniforms.max.y;
|
||||
}
|
||||
|
|
|
@ -1,5 +1,5 @@
|
|||
@compute @workgroup_size(1)
|
||||
fn compute_main() {
|
||||
let a = 1.230000019;
|
||||
let a = 1.23;
|
||||
var b = max(a, 1.17549435e-38);
|
||||
}
|
||||
|
|
|
@ -1,6 +1,6 @@
|
|||
@vertex
|
||||
fn main() -> @builtin(position) vec4<f32> {
|
||||
var light : vec3<f32> = vec3<f32>(1.200000048, 1.0, 2.0);
|
||||
var light : vec3<f32> = vec3<f32>(1.2, 1.0, 2.0);
|
||||
var negative_light : vec3<f32> = -(light);
|
||||
return vec4<f32>();
|
||||
}
|
||||
|
|
|
@ -7,7 +7,7 @@ struct Output {
|
|||
|
||||
@vertex
|
||||
fn main(@builtin(vertex_index) VertexIndex : u32, @builtin(instance_index) InstanceIndex : u32) -> Output {
|
||||
let zv : array<vec2<f32>, 4> = array<vec2<f32>, 4>(vec2<f32>(0.200000003, 0.200000003), vec2<f32>(0.300000012, 0.300000012), vec2<f32>(-0.100000001, -0.100000001), vec2<f32>(1.100000024, 1.100000024));
|
||||
let zv : array<vec2<f32>, 4> = array<vec2<f32>, 4>(vec2<f32>(0.2, 0.2), vec2<f32>(0.3, 0.3), vec2<f32>(-0.1, -0.1), vec2<f32>(1.1, 1.1));
|
||||
let z : f32 = zv[InstanceIndex].x;
|
||||
var output : Output;
|
||||
output.Position = vec4<f32>(0.5, 0.5, z, 1.0);
|
||||
|
|
|
@ -1,6 +1,6 @@
|
|||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let res = frexp(1.230000019);
|
||||
let res = frexp(1.23);
|
||||
let exp : i32 = res.exp;
|
||||
let sig : f32 = res.sig;
|
||||
}
|
||||
|
|
|
@ -1,6 +1,6 @@
|
|||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let res = modf(1.230000019);
|
||||
let res = modf(1.23);
|
||||
let fract : f32 = res.fract;
|
||||
let whole : f32 = res.whole;
|
||||
}
|
||||
|
|
|
@ -2,5 +2,5 @@ enable f16;
|
|||
|
||||
@fragment
|
||||
fn main() -> @location(0) vec4<f32> {
|
||||
return vec4<f32>(0.100000001, 0.200000003, 0.300000012, 0.400000006);
|
||||
return vec4<f32>(0.1, 0.2, 0.3, 0.4);
|
||||
}
|
||||
|
|
|
@ -4,5 +4,5 @@ enable f16;
|
|||
|
||||
@fragment
|
||||
fn main() -> @location(0) vec4<f32> {
|
||||
return vec4<f32>(0.100000001, 0.200000003, 0.300000012, 0.400000006);
|
||||
return vec4<f32>(0.1, 0.2, 0.3, 0.4);
|
||||
}
|
||||
|
|
|
@ -75,7 +75,7 @@ fn comp_main(@builtin(global_invocation_id) gl_GlobalInvocationID : vec3<u32>) {
|
|||
cVel = (cVel / vec2<f32>(f32(cVelCount), f32(cVelCount)));
|
||||
}
|
||||
vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale));
|
||||
vVel = (normalize(vVel) * clamp(length(vVel), 0.0, 0.100000001));
|
||||
vVel = (normalize(vVel) * clamp(length(vVel), 0.0, 0.1));
|
||||
vPos = (vPos + (vVel * params.deltaT));
|
||||
if ((vPos.x < -1.0)) {
|
||||
vPos.x = 1.0;
|
||||
|
|
|
@ -5,5 +5,5 @@ fn bar() {
|
|||
fn main() -> @location(0) vec4<f32> {
|
||||
var a : vec2<f32> = vec2<f32>();
|
||||
bar();
|
||||
return vec4<f32>(0.400000006, 0.400000006, 0.800000012, 1.0);
|
||||
return vec4<f32>(0.4, 0.4, 0.8, 1.0);
|
||||
}
|
||||
|
|
|
@ -9,7 +9,7 @@ struct Interface {
|
|||
|
||||
@vertex
|
||||
fn vert_main() -> Interface {
|
||||
return Interface(0.400000006, 0.600000024, vec4<f32>());
|
||||
return Interface(0.4, 0.6, vec4<f32>());
|
||||
}
|
||||
|
||||
@fragment
|
||||
|
|
Loading…
Reference in New Issue