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