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>
diff --git a/src/tint/writer/float_to_string.cc b/src/tint/writer/float_to_string.cc
index 7494ace..3b4260e 100644
--- a/src/tint/writer/float_to_string.cc
+++ b/src/tint/writer/float_to_string.cc
@@ -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 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 @@
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 @@
// 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 @@
// 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 @@
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
diff --git a/src/tint/writer/float_to_string.h b/src/tint/writer/float_to_string.h
index 4a5afd4..387c0a8 100644
--- a/src/tint/writer/float_to_string.h
+++ b/src/tint/writer/float_to_string.h
@@ -27,11 +27,24 @@
/// @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_
diff --git a/src/tint/writer/float_to_string_test.cc b/src/tint/writer/float_to_string_test.cc
index 8e5f244..901334e 100644
--- a/src/tint/writer/float_to_string_test.cc
+++ b/src/tint/writer/float_to_string_test.cc
@@ -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 @@
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 @@
"-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
diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc
index 56eabc1..7795455 100644
--- a/src/tint/writer/wgsl/generator_impl.cc
+++ b/src/tint/writer/wgsl/generator_impl.cc
@@ -263,7 +263,11 @@
// 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) { //
diff --git a/test/tint/bug/chromium/1350147.wgsl.expected.wgsl b/test/tint/bug/chromium/1350147.wgsl.expected.wgsl
index 03763db..463497f 100644
--- a/test/tint/bug/chromium/1350147.wgsl.expected.wgsl
+++ b/test/tint/bug/chromium/1350147.wgsl.expected.wgsl
@@ -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);
}
}
diff --git a/test/tint/bug/chromium/1367602-1.wgsl.expected.dxc.hlsl b/test/tint/bug/chromium/1367602-1.wgsl.expected.dxc.hlsl
index 35c457d..cd8659f 100644
--- a/test/tint/bug/chromium/1367602-1.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/chromium/1367602-1.wgsl.expected.dxc.hlsl
@@ -1,6 +1,8 @@
-SKIP: FAILED
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
-bug/chromium/1367602-1.wgsl:2:23 error: array size (65536) must be less than 65536
- var v = array<bool, 65536>();
- ^^^^^
-
+void f() {
+ bool v[65535] = (bool[65535])0;
+}
diff --git a/test/tint/bug/tint/1121.wgsl.expected.wgsl b/test/tint/bug/tint/1121.wgsl.expected.wgsl
index bea6cab..6156c87 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/1121.wgsl.expected.wgsl
@@ -48,7 +48,7 @@
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;
}
diff --git a/test/tint/bug/tint/1332.wgsl.expected.wgsl b/test/tint/bug/tint/1332.wgsl.expected.wgsl
index f95540e..27e8e44 100644
--- a/test/tint/bug/tint/1332.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/1332.wgsl.expected.wgsl
@@ -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);
}
diff --git a/test/tint/bug/tint/292.wgsl.expected.wgsl b/test/tint/bug/tint/292.wgsl.expected.wgsl
index b89b327..994483a 100644
--- a/test/tint/bug/tint/292.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/292.wgsl.expected.wgsl
@@ -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>();
}
diff --git a/test/tint/bug/tint/824.wgsl.expected.wgsl b/test/tint/bug/tint/824.wgsl.expected.wgsl
index 9ac4af3..364a771 100644
--- a/test/tint/bug/tint/824.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/824.wgsl.expected.wgsl
@@ -7,7 +7,7 @@
@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);
diff --git a/test/tint/builtins/frexp.wgsl.expected.wgsl b/test/tint/builtins/frexp.wgsl.expected.wgsl
index a36f01a..8d76498 100644
--- a/test/tint/builtins/frexp.wgsl.expected.wgsl
+++ b/test/tint/builtins/frexp.wgsl.expected.wgsl
@@ -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;
}
diff --git a/test/tint/builtins/modf.wgsl.expected.wgsl b/test/tint/builtins/modf.wgsl.expected.wgsl
index a0f85b8..c9d61eb 100644
--- a/test/tint/builtins/modf.wgsl.expected.wgsl
+++ b/test/tint/builtins/modf.wgsl.expected.wgsl
@@ -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;
}
diff --git a/test/tint/extensions/parsing/basic.wgsl.expected.wgsl b/test/tint/extensions/parsing/basic.wgsl.expected.wgsl
index f8bd65a..0cbd662 100644
--- a/test/tint/extensions/parsing/basic.wgsl.expected.wgsl
+++ b/test/tint/extensions/parsing/basic.wgsl.expected.wgsl
@@ -2,5 +2,5 @@
@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);
}
diff --git a/test/tint/extensions/parsing/duplicated_extensions.wgsl.expected.wgsl b/test/tint/extensions/parsing/duplicated_extensions.wgsl.expected.wgsl
index b20c55f..c0c6259 100644
--- a/test/tint/extensions/parsing/duplicated_extensions.wgsl.expected.wgsl
+++ b/test/tint/extensions/parsing/duplicated_extensions.wgsl.expected.wgsl
@@ -4,5 +4,5 @@
@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);
}
diff --git a/test/tint/samples/compute_boids.wgsl.expected.wgsl b/test/tint/samples/compute_boids.wgsl.expected.wgsl
index b41ba8c..253d6f0 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.wgsl
+++ b/test/tint/samples/compute_boids.wgsl.expected.wgsl
@@ -75,7 +75,7 @@
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;
diff --git a/test/tint/samples/simple.wgsl.expected.wgsl b/test/tint/samples/simple.wgsl.expected.wgsl
index 0c37a65..cf8998d 100644
--- a/test/tint/samples/simple.wgsl.expected.wgsl
+++ b/test/tint/samples/simple.wgsl.expected.wgsl
@@ -5,5 +5,5 @@
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);
}
diff --git a/test/tint/shader_io/shared_struct_different_stages.wgsl.expected.wgsl b/test/tint/shader_io/shared_struct_different_stages.wgsl.expected.wgsl
index 4c3967d..ddd459f 100644
--- a/test/tint/shader_io/shared_struct_different_stages.wgsl.expected.wgsl
+++ b/test/tint/shader_io/shared_struct_different_stages.wgsl.expected.wgsl
@@ -9,7 +9,7 @@
@vertex
fn vert_main() -> Interface {
- return Interface(0.400000006, 0.600000024, vec4<f32>());
+ return Interface(0.4, 0.6, vec4<f32>());
}
@fragment
