diff --git a/src/tint/ast/module_clone_test.cc b/src/tint/ast/module_clone_test.cc
index 544e6bf..a79ef0e 100644
--- a/src/tint/ast/module_clone_test.cc
+++ b/src/tint/ast/module_clone_test.cc
@@ -52,7 +52,7 @@
 @group(4) @binding(0) var g6 : texture_external;
 
 var<private> g7 : vec3<f32>;
-@group(0) @binding(1) var<storage, write> g8 : S0;
+@group(0) @binding(1) var<storage, read_write> g8 : S0;
 @group(1) @binding(1) var<storage, read> g9 : S0;
 @group(2) @binding(1) var<storage, read_write> g10 : S0;
 
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
index 737cb07..f201047 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
@@ -275,8 +275,8 @@
     // (`read` for uniform storage class).
     std::string shader = R"(
 struct S {
-  a: i32;
-};
+  a: i32
+}
 
 var<private> a: S;
 @group(1) @binding(1) var<uniform> b: S;
@@ -310,8 +310,8 @@
     // Can't replace `ptr_b` with `a` since the latter is not a pointer.
     std::string shader = R"(
 struct S {
-  a: i32;
-};
+  a: i32
+}
 
 var<private> a: S;
 @group(1) @binding(1) var<uniform> b: S;
@@ -346,12 +346,12 @@
     // storage class.
     std::string shader = R"(
 struct S {
-  a: i32;
-};
+  a: i32
+}
 
 var<private> a: S;
 @group(1) @binding(1) var<uniform> b: S;
-@group(1) @binding(2) var<storage, write> c: S;
+@group(1) @binding(2) var<storage, read_write> c: S;
 fn f() {
   let ptr_b = &b;
   *&a = *ptr_b;
@@ -382,8 +382,8 @@
     // Can't replace `b` with `e` since the latter is not a reference.
     std::string shader = R"(
 struct S {
-  a: i32;
-};
+  a: i32
+}
 
 var<private> a: S;
 let e = 3;
@@ -418,11 +418,11 @@
     // Can't replace `b` with `e` since the latter has a wrong access mode.
     std::string shader = R"(
 struct S {
-  a: i32;
-};
+  a: i32
+}
 
 var<private> a: S;
-@group(0) @binding(0) var<storage, write> e: S;
+@group(0) @binding(0) var<storage, read_write> e: S;
 @group(1) @binding(1) var<uniform> b: S;
 fn f() {
   *&a = *&b;
diff --git a/src/tint/number.cc b/src/tint/number.cc
index 17b005b..6371442 100644
--- a/src/tint/number.cc
+++ b/src/tint/number.cc
@@ -204,4 +204,108 @@
     return value;
 }
 
+uint16_t f16::BitsRepresentation() const {
+    constexpr uint16_t f16_nan = 0x7e00u;
+    constexpr uint16_t f16_pos_inf = 0x7c00u;
+    constexpr uint16_t f16_neg_inf = 0xfc00u;
+
+    // Assert we use binary32 (i.e. float) as underlying type, which has 4 bytes.
+    static_assert(std::is_same<f16::type, float>());
+
+    // The stored value in f16 object must be already quantized, so it should be either NaN, +/-
+    // Inf, or exactly representable by normal or subnormal f16.
+
+    if (std::isnan(value)) {
+        return f16_nan;
+    }
+
+    if (std::isinf(value)) {
+        return value > 0 ? f16_pos_inf : f16_neg_inf;
+    }
+
+    // Now quantized_value must be a finite f16 exactly-representable value.
+    // The following table shows exponent cases for all finite f16 exactly-representable value.
+    // ---------------------------------------------------------------------------
+    // |  Value category  |  Unbiased exp  |  F16 biased exp  |  F32 biased exp  |
+    // |------------------|----------------|------------------|------------------|
+    // |     +/- zero     |        \       |         0        |         0        |
+    // |  Subnormal f16   |   [-24, -15]   |         0        |    [103, 112]    |
+    // |    Normal f16    |   [-14, 15]    |      [1, 30]     |    [113, 142]    |
+    // ---------------------------------------------------------------------------
+
+    constexpr uint32_t max_f32_biased_exp_for_f16_normal_number = 142;
+    constexpr uint32_t min_f32_biased_exp_for_f16_normal_number = 113;
+    constexpr uint32_t max_f32_biased_exp_for_f16_subnormal_number = 112;
+    constexpr uint32_t min_f32_biased_exp_for_f16_subnormal_number = 103;
+
+    constexpr uint32_t f32_sign_mask = 0x80000000u;
+    constexpr uint32_t f32_exp_mask = 0x7f800000u;
+    constexpr uint32_t f32_mantissa_mask = 0x007fffffu;
+    constexpr uint32_t f32_mantissa_bis_number = 23;
+    constexpr uint32_t f32_exp_bias = 127;
+
+    constexpr uint16_t f16_sign_mask = 0x8000u;
+    constexpr uint16_t f16_exp_mask = 0x7c00u;
+    constexpr uint16_t f16_mantissa_mask = 0x03ffu;
+    constexpr uint32_t f16_mantissa_bis_number = 10;
+    constexpr uint32_t f16_exp_bias = 15;
+
+    uint32_t f32_bit_pattern;
+    memcpy(&f32_bit_pattern, &value, 4);
+    uint32_t f32_biased_exponent = (f32_bit_pattern & f32_exp_mask) >> f32_mantissa_bis_number;
+    uint32_t f32_mantissa = f32_bit_pattern & f32_mantissa_mask;
+
+    uint16_t f16_sign_part = static_cast<uint16_t>((f32_bit_pattern & f32_sign_mask) >> 16);
+    TINT_ASSERT(Semantic, (f16_sign_part & ~f16_sign_mask) == 0);
+
+    if ((f32_bit_pattern & ~f32_sign_mask) == 0) {
+        // +/- zero
+        return f16_sign_part;
+    }
+
+    if ((min_f32_biased_exp_for_f16_normal_number <= f32_biased_exponent) &&
+        (f32_biased_exponent <= max_f32_biased_exp_for_f16_normal_number)) {
+        // Normal f16
+        uint32_t f16_biased_exponent = f32_biased_exponent - f32_exp_bias + f16_exp_bias;
+        uint16_t f16_exp_part =
+            static_cast<uint16_t>(f16_biased_exponent << f16_mantissa_bis_number);
+        uint16_t f16_mantissa_part = static_cast<uint16_t>(
+            f32_mantissa >> (f32_mantissa_bis_number - f16_mantissa_bis_number));
+
+        TINT_ASSERT(Semantic, (f16_exp_part & ~f16_exp_mask) == 0);
+        TINT_ASSERT(Semantic, (f16_mantissa_part & ~f16_mantissa_mask) == 0);
+
+        return f16_sign_part | f16_exp_part | f16_mantissa_part;
+    }
+
+    if ((min_f32_biased_exp_for_f16_subnormal_number <= f32_biased_exponent) &&
+        (f32_biased_exponent <= max_f32_biased_exp_for_f16_subnormal_number)) {
+        // Subnormal f16
+        // The resulting exp bits are always 0, and the mantissa bits should be handled specially.
+        uint16_t f16_exp_part = 0;
+        // The resulting subnormal f16 will have only 1 valid mantissa bit if the unbiased exponent
+        // of value is of the minimum, i.e. -24; and have all 10 mantissa bits valid if the unbiased
+        // exponent of value is of the maximum, i.e. -15.
+        uint32_t f16_valid_mantissa_bits =
+            f32_biased_exponent - min_f32_biased_exp_for_f16_subnormal_number + 1;
+        // The resulting f16 mantissa part comes from right-shifting the f32 mantissa bits with
+        // leading 1 added.
+        uint16_t f16_mantissa_part =
+            static_cast<uint16_t>((f32_mantissa | (f32_mantissa_mask + 1)) >>
+                                  (f32_mantissa_bis_number + 1 - f16_valid_mantissa_bits));
+
+        TINT_ASSERT(Semantic, (1 <= f16_valid_mantissa_bits) &&
+                                  (f16_valid_mantissa_bits <= f16_mantissa_bis_number));
+        TINT_ASSERT(Semantic, (f16_mantissa_part & ~((1u << f16_valid_mantissa_bits) - 1)) == 0);
+        TINT_ASSERT(Semantic, (f16_mantissa_part != 0));
+
+        return f16_sign_part | f16_exp_part | f16_mantissa_part;
+    }
+
+    // Neither zero, subnormal f16 or normal f16, shall never hit.
+    tint::diag::List diag;
+    TINT_UNREACHABLE(Semantic, diag);
+    return f16_nan;
+}
+
 }  // namespace tint
diff --git a/src/tint/number.h b/src/tint/number.h
index e130019..32cca59 100644
--- a/src/tint/number.h
+++ b/src/tint/number.h
@@ -186,6 +186,13 @@
         return *this;
     }
 
+    /// Get the binary16 bit pattern in type uint16_t of this value.
+    /// @returns the binary16 bit pattern, in type uint16_t, of the stored quantized f16 value. If
+    /// the value is NaN, the returned value will be 0x7e00u. If the value is positive infinity, the
+    /// returned value will be 0x7c00u. If the input value is negative infinity, the returned value
+    /// will be 0xfc00u.
+    uint16_t BitsRepresentation() const;
+
     /// @param value the input float32 value
     /// @returns the float32 value quantized to the smaller float16 value, through truncation of the
     /// mantissa bits (no rounding). If the float32 value is too large (positive or negative) to be
diff --git a/src/tint/number_test.cc b/src/tint/number_test.cc
index 81acc04..eeb31ed 100644
--- a/src/tint/number_test.cc
+++ b/src/tint/number_test.cc
@@ -217,83 +217,164 @@
     EXPECT_EQ(CheckedConvert<f16>(AFloat(-kHighestF16Subnormal)), f16(-kHighestF16Subnormal));
 }
 
-TEST(NumberTest, QuantizeF16) {
-    constexpr float nan = std::numeric_limits<float>::quiet_NaN();
-    constexpr float inf = std::numeric_limits<float>::infinity();
+// Test cases for f16 subnormal quantization and BitsRepresentation.
+// The ULP is based on float rather than double or f16, since F16::Quantize and
+// F16::BitsRepresentation take float as input.
+constexpr float lowestPositiveNormalF16 = 0x1p-14;
+constexpr float lowestPositiveNormalF16PlusULP = 0x1.000002p-14;
+constexpr float lowestPositiveNormalF16MinusULP = 0x1.fffffep-15;
+constexpr float highestPositiveSubnormalF16 = 0x0.ffcp-14;
+constexpr float highestPositiveSubnormalF16PlusULP = 0x1.ff8002p-15;
+constexpr float highestPositiveSubnormalF16MinusULP = 0x1.ff7ffep-15;
+constexpr float lowestPositiveSubnormalF16 = 0x1.p-24;
+constexpr float lowestPositiveSubnormalF16PlusULP = 0x1.000002p-24;
+constexpr float lowestPositiveSubnormalF16MinusULP = 0x1.fffffep-25;
 
-    EXPECT_EQ(f16(0.0), 0.0f);
-    EXPECT_EQ(f16(1.0), 1.0f);
-    EXPECT_EQ(f16(0.00006106496), 0.000061035156f);
-    EXPECT_EQ(f16(1.0004883), 1.0f);
-    EXPECT_EQ(f16(-8196), -8192.f);
-    EXPECT_EQ(f16(65504.003), inf);
-    EXPECT_EQ(f16(-65504.003), -inf);
-    EXPECT_EQ(f16(inf), inf);
-    EXPECT_EQ(f16(-inf), -inf);
-    EXPECT_TRUE(std::isnan(f16(nan)));
+constexpr uint16_t lowestPositiveNormalF16Bits = 0x0400u;
+constexpr uint16_t highestPositiveSubnormalF16Bits = 0x03ffu;
+constexpr uint16_t lowestPositiveSubnormalF16Bits = 0x0001u;
 
-    // Test for subnormal quantization.
-    // The ULP is based on float rather than double or f16, since F16::Quantize take float as input.
-    constexpr float lowestPositiveNormalF16 = 0x1p-14;
-    constexpr float lowestPositiveNormalF16PlusULP = 0x1.000002p-14;
-    constexpr float lowestPositiveNormalF16MinusULP = 0x1.fffffep-15;
-    constexpr float highestPositiveSubnormalF16 = 0x0.ffcp-14;
-    constexpr float highestPositiveSubnormalF16PlusULP = 0x1.ff8002p-15;
-    constexpr float highestPositiveSubnormalF16MinusULP = 0x1.ff7ffep-15;
-    constexpr float lowestPositiveSubnormalF16 = 0x1.p-24;
-    constexpr float lowestPositiveSubnormalF16PlusULP = 0x1.000002p-24;
-    constexpr float lowestPositiveSubnormalF16MinusULP = 0x1.fffffep-25;
+constexpr float highestNegativeNormalF16 = -lowestPositiveNormalF16;
+constexpr float highestNegativeNormalF16PlusULP = -lowestPositiveNormalF16MinusULP;
+constexpr float highestNegativeNormalF16MinusULP = -lowestPositiveNormalF16PlusULP;
+constexpr float lowestNegativeSubnormalF16 = -highestPositiveSubnormalF16;
+constexpr float lowestNegativeSubnormalF16PlusULP = -highestPositiveSubnormalF16MinusULP;
+constexpr float lowestNegativeSubnormalF16MinusULP = -highestPositiveSubnormalF16PlusULP;
+constexpr float highestNegativeSubnormalF16 = -lowestPositiveSubnormalF16;
+constexpr float highestNegativeSubnormalF16PlusULP = -lowestPositiveSubnormalF16MinusULP;
+constexpr float highestNegativeSubnormalF16MinusULP = -lowestPositiveSubnormalF16PlusULP;
 
-    constexpr float highestNegativeNormalF16 = -lowestPositiveNormalF16;
-    constexpr float highestNegativeNormalF16PlusULP = -lowestPositiveNormalF16MinusULP;
-    constexpr float highestNegativeNormalF16MinusULP = -lowestPositiveNormalF16PlusULP;
-    constexpr float lowestNegativeSubnormalF16 = -highestPositiveSubnormalF16;
-    constexpr float lowestNegativeSubnormalF16PlusULP = -highestPositiveSubnormalF16MinusULP;
-    constexpr float lowestNegativeSubnormalF16MinusULP = -highestPositiveSubnormalF16PlusULP;
-    constexpr float highestNegativeSubnormalF16 = -lowestPositiveSubnormalF16;
-    constexpr float highestNegativeSubnormalF16PlusULP = -lowestPositiveSubnormalF16MinusULP;
-    constexpr float highestNegativeSubnormalF16MinusULP = -lowestPositiveSubnormalF16PlusULP;
+constexpr uint16_t highestNegativeNormalF16Bits = 0x8400u;
+constexpr uint16_t lowestNegativeSubnormalF16Bits = 0x83ffu;
+constexpr uint16_t highestNegativeSubnormalF16Bits = 0x8001u;
 
-    // Value larger than or equal to lowest positive normal f16 will be quantized to normal f16.
-    EXPECT_EQ(f16(lowestPositiveNormalF16PlusULP), lowestPositiveNormalF16);
-    EXPECT_EQ(f16(lowestPositiveNormalF16), lowestPositiveNormalF16);
-    // Positive value smaller than lowest positive normal f16 but not smaller than lowest positive
-    // subnormal f16 will be quantized to subnormal f16 or zero.
-    EXPECT_EQ(f16(lowestPositiveNormalF16MinusULP), highestPositiveSubnormalF16);
-    EXPECT_EQ(f16(highestPositiveSubnormalF16PlusULP), highestPositiveSubnormalF16);
-    EXPECT_EQ(f16(highestPositiveSubnormalF16), highestPositiveSubnormalF16);
-    EXPECT_EQ(f16(highestPositiveSubnormalF16MinusULP), 0x0.ff8p-14);
-    EXPECT_EQ(f16(lowestPositiveSubnormalF16PlusULP), lowestPositiveSubnormalF16);
-    EXPECT_EQ(f16(lowestPositiveSubnormalF16), lowestPositiveSubnormalF16);
-    // Positive value smaller than lowest positive subnormal f16 will be quantized to zero.
-    EXPECT_EQ(f16(lowestPositiveSubnormalF16MinusULP), 0.0);
-    // Test the mantissa discarding, the least significant mantissa bit is 0x1p-24 = 0x0.004p-14.
-    EXPECT_EQ(f16(0x0.064p-14), 0x0.064p-14);
-    EXPECT_EQ(f16(0x0.067fecp-14), 0x0.064p-14);
-    EXPECT_EQ(f16(0x0.063ffep-14), 0x0.060p-14);
-    EXPECT_EQ(f16(0x0.008p-14), 0x0.008p-14);
-    EXPECT_EQ(f16(0x0.00bffep-14), 0x0.008p-14);
-    EXPECT_EQ(f16(0x0.007ffep-14), 0x0.004p-14);
+constexpr float f32_nan = std::numeric_limits<float>::quiet_NaN();
+constexpr float f32_inf = std::numeric_limits<float>::infinity();
 
-    // Vice versa for negative cases.
-    EXPECT_EQ(f16(highestNegativeNormalF16MinusULP), highestNegativeNormalF16);
-    EXPECT_EQ(f16(highestNegativeNormalF16), highestNegativeNormalF16);
-    EXPECT_EQ(f16(highestNegativeNormalF16PlusULP), lowestNegativeSubnormalF16);
-    EXPECT_EQ(f16(lowestNegativeSubnormalF16MinusULP), lowestNegativeSubnormalF16);
-    EXPECT_EQ(f16(lowestNegativeSubnormalF16), lowestNegativeSubnormalF16);
-    EXPECT_EQ(f16(lowestNegativeSubnormalF16PlusULP), -0x0.ff8p-14);
-    EXPECT_EQ(f16(highestNegativeSubnormalF16MinusULP), highestNegativeSubnormalF16);
-    EXPECT_EQ(f16(highestNegativeSubnormalF16), highestNegativeSubnormalF16);
-    EXPECT_EQ(f16(highestNegativeSubnormalF16PlusULP), -0.0);
-    // Test the mantissa discarding.
-    EXPECT_EQ(f16(-0x0.064p-14), -0x0.064p-14);
-    EXPECT_EQ(f16(-0x0.067fecp-14), -0x0.064p-14);
-    EXPECT_EQ(f16(-0x0.063ffep-14), -0x0.060p-14);
-    EXPECT_EQ(f16(-0x0.008p-14), -0x0.008p-14);
-    EXPECT_EQ(f16(-0x0.00bffep-14), -0x0.008p-14);
-    EXPECT_EQ(f16(-0x0.007ffep-14), -0x0.004p-14);
+struct F16TestCase {
+    float input_value;
+    float quantized_value;
+    uint16_t f16_bit_pattern;
+};
+
+using NumberF16Test = testing::TestWithParam<F16TestCase>;
+
+TEST_P(NumberF16Test, QuantizeF16) {
+    float input_value = GetParam().input_value;
+    float quantized_value = GetParam().quantized_value;
+
+    std::stringstream ss;
+    ss << "input value = " << input_value << ", expected quantized value = " << quantized_value;
+    SCOPED_TRACE(ss.str());
+
+    if (std::isnan(quantized_value)) {
+        EXPECT_TRUE(std::isnan(f16(input_value)));
+    } else {
+        EXPECT_EQ(f16(input_value), quantized_value);
+    }
 }
 
+TEST_P(NumberF16Test, BitsRepresentation) {
+    float input_value = GetParam().input_value;
+    uint16_t representation = GetParam().f16_bit_pattern;
+
+    std::stringstream ss;
+    ss << "input value = " << input_value
+       << ", expected binary16 bits representation = " << std::hex << std::showbase
+       << representation;
+    SCOPED_TRACE(ss.str());
+
+    EXPECT_EQ(f16(input_value).BitsRepresentation(), representation);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    NumberF16Test,
+    NumberF16Test,
+    testing::ValuesIn(std::vector<F16TestCase>{
+        // NaN, Inf
+        {f32_inf, f32_inf, 0x7c00u},
+        {-f32_inf, -f32_inf, 0xfc00u},
+        {f32_nan, f32_nan, 0x7e00u},
+        {-f32_nan, -f32_nan, 0x7e00u},
+        // +/- zero
+        {+0.0f, 0.0f, 0x0000u},
+        {-0.0f, -0.0f, 0x8000u},
+        // Value in normal f16 range
+        {1.0f, 1.0f, 0x3c00u},
+        {-1.0f, -1.0f, 0xbc00u},
+        //   0.00006106496 quantized to 0.000061035156 = 0x1p-14
+        {0.00006106496f, 0.000061035156f, 0x0400u},
+        {-0.00006106496f, -0.000061035156f, 0x8400u},
+        //   1.0004883 quantized to 1.0 = 0x1p0
+        {1.0004883f, 1.0f, 0x3c00u},
+        {-1.0004883f, -1.0f, 0xbc00u},
+        //   8196.0 quantized to 8192.0 = 0x1p13
+        {8196.0f, 8192.f, 0x7000u},
+        {-8196.0f, -8192.f, 0xf000u},
+        // Value in subnormal f16 range
+        {0x0.034p-14f, 0x0.034p-14f, 0x000du},
+        {-0x0.034p-14f, -0x0.034p-14f, 0x800du},
+        {0x0.068p-14f, 0x0.068p-14f, 0x001au},
+        {-0x0.068p-14f, -0x0.068p-14f, 0x801au},
+        //   0x0.06b7p-14 quantized to 0x0.068p-14
+        {0x0.06b7p-14f, 0x0.068p-14f, 0x001au},
+        {-0x0.06b7p-14f, -0x0.068p-14, 0x801au},
+        // Value out of f16 range
+        {65504.003f, f32_inf, 0x7c00u},
+        {-65504.003f, -f32_inf, 0xfc00u},
+        {0x1.234p56f, f32_inf, 0x7c00u},
+        {-0x4.321p65f, -f32_inf, 0xfc00u},
+
+        // Test for subnormal quantization.
+        // Value larger than or equal to lowest positive normal f16 will be quantized to normal f16.
+        {lowestPositiveNormalF16PlusULP, lowestPositiveNormalF16, lowestPositiveNormalF16Bits},
+        {lowestPositiveNormalF16, lowestPositiveNormalF16, lowestPositiveNormalF16Bits},
+        // Positive value smaller than lowest positive normal f16 but not smaller than lowest
+        // positive
+        // subnormal f16 will be quantized to subnormal f16 or zero.
+        {lowestPositiveNormalF16MinusULP, highestPositiveSubnormalF16,
+         highestPositiveSubnormalF16Bits},
+        {highestPositiveSubnormalF16PlusULP, highestPositiveSubnormalF16,
+         highestPositiveSubnormalF16Bits},
+        {highestPositiveSubnormalF16, highestPositiveSubnormalF16, highestPositiveSubnormalF16Bits},
+        {highestPositiveSubnormalF16MinusULP, 0x0.ff8p-14, 0x03feu},
+        {lowestPositiveSubnormalF16PlusULP, lowestPositiveSubnormalF16,
+         lowestPositiveSubnormalF16Bits},
+        {lowestPositiveSubnormalF16, lowestPositiveSubnormalF16, lowestPositiveSubnormalF16Bits},
+        // Positive value smaller than lowest positive subnormal f16 will be quantized to zero.
+        {lowestPositiveSubnormalF16MinusULP, 0.0, 0x0000u},
+        // Test the mantissa discarding, the least significant mantissa bit is 0x1p-24 =
+        // 0x0.004p-14.
+        {0x0.064p-14f, 0x0.064p-14, 0x0019u},
+        {0x0.067fecp-14f, 0x0.064p-14, 0x0019u},
+        {0x0.063ffep-14f, 0x0.060p-14, 0x0018u},
+        {0x0.008p-14f, 0x0.008p-14, 0x0002u},
+        {0x0.00bffep-14f, 0x0.008p-14, 0x0002u},
+        {0x0.007ffep-14f, 0x0.004p-14, 0x0001u},
+
+        // Vice versa for negative cases.
+        {highestNegativeNormalF16MinusULP, highestNegativeNormalF16, highestNegativeNormalF16Bits},
+        {highestNegativeNormalF16, highestNegativeNormalF16, highestNegativeNormalF16Bits},
+        {highestNegativeNormalF16PlusULP, lowestNegativeSubnormalF16,
+         lowestNegativeSubnormalF16Bits},
+        {lowestNegativeSubnormalF16MinusULP, lowestNegativeSubnormalF16,
+         lowestNegativeSubnormalF16Bits},
+        {lowestNegativeSubnormalF16, lowestNegativeSubnormalF16, lowestNegativeSubnormalF16Bits},
+        {lowestNegativeSubnormalF16PlusULP, -0x0.ff8p-14, 0x83feu},
+        {highestNegativeSubnormalF16MinusULP, highestNegativeSubnormalF16,
+         highestNegativeSubnormalF16Bits},
+        {highestNegativeSubnormalF16, highestNegativeSubnormalF16, highestNegativeSubnormalF16Bits},
+        {highestNegativeSubnormalF16PlusULP, -0.0, 0x8000u},
+        // Test the mantissa discarding.
+        {-0x0.064p-14f, -0x0.064p-14, 0x8019u},
+        {-0x0.067fecp-14f, -0x0.064p-14, 0x8019u},
+        {-0x0.063ffep-14f, -0x0.060p-14, 0x8018u},
+        {-0x0.008p-14f, -0x0.008p-14, 0x8002u},
+        {-0x0.00bffep-14f, -0x0.008p-14, 0x8002u},
+        {-0x0.007ffep-14f, -0x0.004p-14, 0x8001u},
+        /////////////////////////////////////
+    }));
+
 using BinaryCheckedCase = std::tuple<std::optional<AInt>, AInt, AInt>;
 
 #undef OVERFLOW  // corecrt_math.h :(
diff --git a/src/tint/reader/spirv/function_cfg_test.cc b/src/tint/reader/spirv/function_cfg_test.cc
index edeea55..efeeee8 100644
--- a/src/tint/reader/spirv/function_cfg_test.cc
+++ b/src/tint/reader/spirv/function_cfg_test.cc
@@ -3188,6 +3188,9 @@
     EXPECT_EQ(fe.GetBlockInfo(40)->construct, constructs[1].get());
     EXPECT_EQ(fe.GetBlockInfo(50)->construct, constructs[1].get());
     EXPECT_EQ(fe.GetBlockInfo(99)->construct, constructs[0].get());
+    
+    // SPIR-V 1.6 Rev 2 made this invalid SPIR-V.
+    p->DeliberatelyInvalidSpirv();
 }
 
 TEST_F(SpvParserCFGTest, LabelControlFlowConstructs_MergeBlockIsAlsoSingleBlockLoop) {
@@ -4726,6 +4729,12 @@
     ASSERT_NE(bi50, nullptr);
     EXPECT_EQ(bi50->succ_edge.count(20), 1u);
     EXPECT_EQ(bi50->succ_edge[20], EdgeKind::kBack);
+
+    // SPIR-V 1.6 Rev 2 made this invalid SPIR-V.
+    // The continue target also has the LoopMerge in it, but the continue
+    // target 20 is not structurally post-dominated by the back-edge block 50.
+    // Don't dump this as an end-to-end test.
+    p->DeliberatelyInvalidSpirv();
 }
 
 TEST_F(SpvParserCFGTest, ClassifyCFGEdges_PrematureExitFromContinueConstruct) {
@@ -8640,6 +8649,9 @@
 return;
 )";
     ASSERT_EQ(expect, got);
+
+    // Continue target does not structurally dominate the backedge block.
+    p->DeliberatelyInvalidSpirv();
 }
 
 TEST_F(SpvParserCFGTest, EmitBody_Loop_Never) {
@@ -12608,6 +12620,11 @@
     const Construct* c = fe.GetBlockInfo(20)->construct;
     EXPECT_EQ(c->kind, Construct::kContinue);
     EXPECT_EQ(fe.SiblingLoopConstruct(c), nullptr);
+
+    // SPIR-V 1.6 Rev 2 made this invalid SPIR-V.
+    // Continue target is not structurally post dominated by the backedge block.
+    // Don't dump this as an end-to-end test.
+    p->DeliberatelyInvalidSpirv();
 }
 
 TEST_F(SpvParserCFGTest, SiblingLoopConstruct_HasSiblingLoop) {
diff --git a/src/tint/reader/spirv/function_misc_test.cc b/src/tint/reader/spirv/function_misc_test.cc
index 3f9c398..5054c67 100644
--- a/src/tint/reader/spirv/function_misc_test.cc
+++ b/src/tint/reader/spirv/function_misc_test.cc
@@ -290,51 +290,6 @@
                              {4, "", "vector component index is larger than 3: 4"},
                              {99999, "", "vector component index is larger than 3: 99999"}}));
 
-TEST_F(SpvParserTest, ValueFromBlockNotInBlockOrder) {
-    // crbug.com/tint/804
-    const auto assembly = Preamble() + CommonTypes() + R"(
-     %float_42 = OpConstant %float 42.0
-     %cond = OpUndef %bool
-
-     %100 = OpFunction %void None %voidfn
-     %10 = OpLabel
-     OpBranch %30
-
-     ; unreachable
-     %20 = OpLabel
-     %499 = OpFAdd %float %float_42 %float_42
-     %500 = OpFAdd %float %499 %float_42
-     OpBranch %25
-
-     %25 = OpLabel
-     OpBranch %80
-
-
-     %30 = OpLabel
-     OpLoopMerge %90 %80 None
-     OpBranchConditional %cond %90 %40
-
-     %40 = OpLabel
-     OpBranch %90
-
-     %80 = OpLabel ; unreachable continue target
-                ; but "dominated" by %20 and %25
-     %81 = OpFMul %float %500 %float_42 ; %500 is defined in %20
-     OpBranch %30 ; backedge
-
-     %90 = OpLabel
-     OpReturn
-     OpFunctionEnd
-)";
-    auto p = parser(test::Assemble(assembly));
-    ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
-    auto fe = p->function_emitter(100);
-    EXPECT_TRUE(fe.EmitBody()) << p->error();
-    auto ast_body = fe.ast_body();
-    const auto got = test::ToString(p->program(), ast_body);
-    EXPECT_THAT(got, HasSubstr("let x_81 : f32 = (0.0f * 42.0f);"));
-}
-
 // TODO(dneto): OpSizeof : requires Kernel (OpenCL)
 
 }  // namespace
diff --git a/src/tint/reader/spirv/function_var_test.cc b/src/tint/reader/spirv/function_var_test.cc
index 6371234..ff28299 100644
--- a/src/tint/reader/spirv/function_var_test.cc
+++ b/src/tint/reader/spirv/function_var_test.cc
@@ -1418,65 +1418,6 @@
     EXPECT_EQ(expect, got);
 }
 
-TEST_F(SpvParserFunctionVarTest, EmitStatement_Phi_ValueFromBlockNotInBlockOrderIgnored) {
-    // From crbug.com/tint/804
-    const auto assembly = Preamble() + R"(
-     %float_42 = OpConstant %float 42.0
-     %cond = OpUndef %bool
-
-     %100 = OpFunction %void None %voidfn
-     %10 = OpLabel
-     OpBranch %30
-
-     ; unreachable
-     %20 = OpLabel
-     %499 = OpFAdd %float %float_42 %float_42
-     %500 = OpFAdd %float %499 %float_42
-     OpBranch %25
-
-     %25 = OpLabel
-     OpBranch %80
-
-
-     %30 = OpLabel
-     OpLoopMerge %90 %80 None
-     OpBranchConditional %cond %90 %40
-
-     %40 = OpLabel
-     OpBranch %90
-
-     %80 = OpLabel ; unreachable continue target
-                ; but "dominated" by %20 and %25
-     %81 = OpPhi %float %500 %25
-     OpBranch %30 ; backedge
-
-     %90 = OpLabel
-     OpReturn
-     OpFunctionEnd
-)";
-    auto p = parser(test::Assemble(assembly));
-    ASSERT_TRUE(p->BuildAndParseInternalModule()) << p->error() << assembly;
-    auto fe = p->function_emitter(100);
-    EXPECT_TRUE(fe.EmitBody()) << p->error();
-
-    const auto* expected = R"(loop {
-  if (false) {
-    break;
-  }
-  break;
-
-  continuing {
-    var x_81_phi_1 : f32;
-    let x_81 : f32 = x_81_phi_1;
-  }
-}
-return;
-)";
-    auto ast_body = fe.ast_body();
-    const auto got = test::ToString(p->program(), ast_body);
-    EXPECT_EQ(got, expected);
-}
-
 TEST_F(SpvParserFunctionVarTest, EmitStatement_Hoist_CompositeInsert) {
     // From crbug.com/tint/804
     const auto assembly = Preamble() + R"(
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index deaa17c..94468c9 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -1348,12 +1348,6 @@
                 members.push_back(member.value);
             }
 
-            // TODO(crbug.com/tint/1475): Remove support for semicolons.
-            if (auto sc = peek(); sc.Is(Token::Type::kSemicolon)) {
-                deprecated(sc.source(), "struct members should be separated with commas");
-                next();
-                continue;
-            }
             if (!match(Token::Type::kComma)) {
                 break;
             }
diff --git a/src/tint/reader/wgsl/parser_impl_struct_decl_test.cc b/src/tint/reader/wgsl/parser_impl_struct_decl_test.cc
index 106ffd5..1c6d6bc 100644
--- a/src/tint/reader/wgsl/parser_impl_struct_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_struct_decl_test.cc
@@ -104,23 +104,5 @@
     EXPECT_EQ(p->error(), "1:10: expected '{' for struct declaration");
 }
 
-// TODO(crbug.com/tint/1475): Remove this.
-TEST_F(ParserImplTest, DEPRECATED_StructDecl_Parses_WithSemicolons) {
-    auto p = parser(R"(
-struct S {
-  a : i32;
-  b : f32;
-})");
-    auto s = p->struct_decl();
-    EXPECT_FALSE(p->has_error());
-    EXPECT_FALSE(s.errored);
-    EXPECT_TRUE(s.matched);
-    ASSERT_NE(s.value, nullptr);
-    ASSERT_EQ(s->name, p->builder().Symbols().Register("S"));
-    ASSERT_EQ(s->members.size(), 2u);
-    EXPECT_EQ(s->members[0]->symbol, p->builder().Symbols().Register("a"));
-    EXPECT_EQ(s->members[1]->symbol, p->builder().Symbols().Register("b"));
-}
-
 }  // namespace
 }  // namespace tint::reader::wgsl
diff --git a/src/tint/resolver/storage_class_validation_test.cc b/src/tint/resolver/storage_class_validation_test.cc
index 0e2be29..3b98fdb 100644
--- a/src/tint/resolver/storage_class_validation_test.cc
+++ b/src/tint/resolver/storage_class_validation_test.cc
@@ -180,6 +180,33 @@
         R"(56:78 error: only variables in <storage> storage class may declare an access mode)");
 }
 
+TEST_F(ResolverStorageClassValidationTest, Storage_ReadAccessMode) {
+    // @group(0) @binding(0) var<storage, read> a : i32;
+    GlobalVar(Source{{56, 78}}, "a", ty.i32(), ast::StorageClass::kStorage, ast::Access::kRead,
+              GroupAndBinding(0, 0));
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, Storage_ReadWriteAccessMode) {
+    // @group(0) @binding(0) var<storage, read_write> a : i32;
+    GlobalVar(Source{{56, 78}}, "a", ty.i32(), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+              GroupAndBinding(0, 0));
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, Storage_WriteAccessMode) {
+    // @group(0) @binding(0) var<storage, read_write> a : i32;
+    GlobalVar(Source{{56, 78}}, "a", ty.i32(), ast::StorageClass::kStorage, ast::Access::kWrite,
+              GroupAndBinding(0, 0));
+
+    ASSERT_FALSE(r()->Resolve());
+
+    EXPECT_EQ(r()->error(),
+              R"(56:78 error: access mode 'write' is not valid for the 'storage' address space)");
+}
+
 TEST_F(ResolverStorageClassValidationTest, StorageBufferNoError_Basic) {
     // struct S { x : i32 };
     // var<storage, read> g : S;
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index a2e7799..ee98d03 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -554,11 +554,20 @@
             // https://gpuweb.github.io/gpuweb/wgsl/#variable-declaration
             // The access mode always has a default, and except for variables in the
             // storage storage class, must not be written.
-            if (global->StorageClass() != ast::StorageClass::kStorage &&
-                var->declared_access != ast::Access::kUndefined) {
-                AddError("only variables in <storage> storage class may declare an access mode",
-                         var->source);
-                return false;
+            if (var->declared_access != ast::Access::kUndefined) {
+                if (global->StorageClass() == ast::StorageClass::kStorage) {
+                    // The access mode for the storage address space can only be 'read' or
+                    // 'read_write'.
+                    if (var->declared_access == ast::Access::kWrite) {
+                        AddError("access mode 'write' is not valid for the 'storage' address space",
+                                 decl->source);
+                        return false;
+                    }
+                } else {
+                    AddError("only variables in <storage> storage class may declare an access mode",
+                             decl->source);
+                    return false;
+                }
             }
 
             if (!AtomicVariable(global, atomic_composite_info)) {
diff --git a/src/tint/transform/binding_remapper_test.cc b/src/tint/transform/binding_remapper_test.cc
index 3274886..564a3a5 100644
--- a/src/tint/transform/binding_remapper_test.cc
+++ b/src/tint/transform/binding_remapper_test.cc
@@ -137,9 +137,9 @@
   a : f32,
 };
 
-@group(2) @binding(1) var<storage, read> a : S;
+@group(2) @binding(1) var<storage, read_write> a : S;
 
-@group(3) @binding(2) var<storage, write> b : S;
+@group(3) @binding(2) var<storage, read_write> b : S;
 
 @group(4) @binding(3) var<storage, read> c : S;
 
@@ -153,9 +153,9 @@
   a : f32,
 }
 
-@group(2) @binding(1) var<storage, write> a : S;
+@group(2) @binding(1) var<storage, read_write> a : S;
 
-@group(3) @binding(2) var<storage, write> b : S;
+@group(3) @binding(2) var<storage, read_write> b : S;
 
 @group(4) @binding(3) var<storage, read> c : S;
 
@@ -168,7 +168,7 @@
     data.Add<BindingRemapper::Remappings>(
         BindingRemapper::BindingPoints{},
         BindingRemapper::AccessControls{
-            {{2, 1}, ast::Access::kWrite},  // Modify access control
+            {{2, 1}, ast::Access::kReadWrite},  // Modify access control
             // Keep @group(3) @binding(2) as is
             {{4, 3}, ast::Access::kRead},  // Add access control
         });
@@ -197,9 +197,9 @@
   a : f32,
 }
 
-@group(4) @binding(5) var<storage, write> a : S;
+@group(4) @binding(5) var<storage, read_write> a : S;
 
-@group(6) @binding(7) var<storage, write> b : S;
+@group(6) @binding(7) var<storage, read_write> b : S;
 
 @compute @workgroup_size(1)
 fn f() {
@@ -213,8 +213,8 @@
             {{3, 2}, {6, 7}},
         },
         BindingRemapper::AccessControls{
-            {{2, 1}, ast::Access::kWrite},
-            {{3, 2}, ast::Access::kWrite},
+            {{2, 1}, ast::Access::kReadWrite},
+            {{3, 2}, ast::Access::kReadWrite},
         });
     auto got = Run<BindingRemapper>(src, data);
 
diff --git a/src/tint/transform/num_workgroups_from_uniform_test.cc b/src/tint/transform/num_workgroups_from_uniform_test.cc
index 8562c01..093081c 100644
--- a/src/tint/transform/num_workgroups_from_uniform_test.cc
+++ b/src/tint/transform/num_workgroups_from_uniform_test.cc
@@ -568,7 +568,7 @@
 @group(3) @binding(0) var g5 : texture_depth_cube_array;
 @group(4) @binding(0) var g6 : texture_external;
 
-@group(0) @binding(1) var<storage, write> g8 : S0;
+@group(0) @binding(1) var<storage, read_write> g8 : S0;
 @group(1) @binding(3) var<storage, read> g9 : S0;
 @group(3) @binding(2) var<storage, read_write> g10 : S0;
 
@@ -634,7 +634,7 @@
 
 @group(4) @binding(0) var g6 : texture_external;
 
-@group(0) @binding(1) var<storage, write> g8 : S0;
+@group(0) @binding(1) var<storage, read_write> g8 : S0;
 
 @group(1) @binding(3) var<storage, read> g9 : S0;
 
diff --git a/src/tint/writer/glsl/generator_impl_function_test.cc b/src/tint/writer/glsl/generator_impl_function_test.cc
index a70e238..450041b 100644
--- a/src/tint/writer/glsl/generator_impl_function_test.cc
+++ b/src/tint/writer/glsl/generator_impl_function_test.cc
@@ -549,7 +549,7 @@
                                     Member("b", ty.f32()),
                                 });
 
-    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kWrite,
+    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
               ast::AttributeList{
                   create<ast::BindingAttribute>(0u),
                   create<ast::GroupAttribute>(1u),
diff --git a/src/tint/writer/hlsl/generator_impl_function_test.cc b/src/tint/writer/hlsl/generator_impl_function_test.cc
index 85647a5..89d74c5 100644
--- a/src/tint/writer/hlsl/generator_impl_function_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_function_test.cc
@@ -503,7 +503,7 @@
                                     Member("b", ty.f32()),
                                 });
 
-    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kWrite,
+    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
               ast::AttributeList{
                   create<ast::BindingAttribute>(0u),
                   create<ast::GroupAttribute>(1u),
diff --git a/src/tint/writer/wgsl/generator_impl_variable_test.cc b/src/tint/writer/wgsl/generator_impl_variable_test.cc
index a02aed8..a058bc7 100644
--- a/src/tint/writer/wgsl/generator_impl_variable_test.cc
+++ b/src/tint/writer/wgsl/generator_impl_variable_test.cc
@@ -56,21 +56,6 @@
     EXPECT_EQ(out.str(), R"(@binding(0) @group(0) var<storage, read> a : S;)");
 }
 
-TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Write) {
-    auto* s = Structure("S", {Member("a", ty.i32())});
-    auto* v = GlobalVar("a", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kWrite,
-                        ast::AttributeList{
-                            create<ast::BindingAttribute>(0u),
-                            create<ast::GroupAttribute>(0u),
-                        });
-
-    GeneratorImpl& gen = Build();
-
-    std::stringstream out;
-    ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
-    EXPECT_EQ(out.str(), R"(@binding(0) @group(0) var<storage, write> a : S;)");
-}
-
 TEST_F(WgslGeneratorImplTest, EmitVariable_Access_ReadWrite) {
     auto* s = Structure("S", {Member("a", ty.i32())});
     auto* v = GlobalVar("a", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,