diff --git a/src/tint/number.cc b/src/tint/number.cc
index 3ead5b0..17b005b 100644
--- a/src/tint/number.cc
+++ b/src/tint/number.cc
@@ -15,9 +15,12 @@
 #include "src/tint/number.h"
 
 #include <algorithm>
+#include <cmath>
 #include <cstring>
 #include <ostream>
 
+#include "src/tint/debug.h"
+
 namespace tint {
 
 std::ostream& operator<<(std::ostream& out, ConversionFailure failure) {
@@ -38,18 +41,165 @@
         return -std::numeric_limits<f16::type>::infinity();
     }
     // Below value must be within the finite range of a f16.
+    // Assert we use binary32 (i.e. float) as underlying type, which has 4 bytes.
+    static_assert(std::is_same<f16::type, float>());
+    const uint32_t sign_mask = 0x80000000u;      // Mask for the sign bit
+    const uint32_t exponent_mask = 0x7f800000u;  // Mask for 8 exponent bits
+
     uint32_t u32;
     memcpy(&u32, &value, 4);
-    if ((u32 & 0x7fffffffu) == 0) {  // ~sign
+
+    if ((u32 & ~sign_mask) == 0) {
         return value;                // +/- zero
     }
-    if ((u32 & 0x7f800000) == 0x7f800000) {  // exponent all 1's
+    if ((u32 & exponent_mask) == exponent_mask) {  // exponent all 1's
         return value;                        // inf or nan
     }
-    // f32 bits :  1 sign,  8 exponent, 23 mantissa
-    // f16 bits :  1 sign,  5 exponent, 10 mantissa
-    // Mask the value to preserve the sign, exponent and most-significant 10 mantissa bits.
-    u32 = u32 & 0xffffe000u;
+
+    // We are now going to quantize a f32 number into subnormal f16 and store the result value back
+    // into a f32 variable. Notice that all subnormal f16 values are just normal f32 values. Below
+    // will show that we can do this quantization by just masking out 13 or more lowest mantissa
+    // bits of the original f32 number.
+    //
+    // Note:
+    // f32 has 1 sign bit, 8 exponent bits for biased exponent (i.e. unbiased exponent + 127), and
+    // 23 mantissa bits. Binary form: s_eeeeeeee_mmmmmmmmmmmmmmmmmmmmmmm
+    // f16 has 1 sign bit, 5 exponent bits for biased exponent (i.e. unbiased exponent + 15), and
+    // 10 mantissa bits. Binary form: s_eeeee_mmmmmmmmmm
+    // The largest finite f16 number has a biased exponent of 11110 in binary, or 30 decimal, and so
+    // a unbiased exponent of 30 - 15 = 15.
+    // The smallest finite f16 number has a biased exponent of 00001 in binary, or 1 decimal, and so
+    // a unbiased exponent of 1 - 15 = -14.
+    //
+    // We may follow the argument below:
+    // 1. All normal or subnormal f16 values, range from 0x1.p-24 to 0x1.ffcp15, are exactly
+    // representable by normal f32 number.
+    //   1.1. We can denote the set of all f32 number that are exact representation of finite f16
+    //   values by `R`.
+    //   1.2. We can do the quantization by mapping a normal f32 value v (in the f16 finite range)
+    //   to a certain f32 number v' in the set R, which is the largest (by the meaning of absolute
+    //   value) one among all values in R that are no larger than v.
+    // 2. We can decide whether a given normal f32 number v is in the set R, by looking at its
+    // mantissa bits and biased exponent `e`. Recall that biased exponent e is unbiased exponent +
+    // 127, and in the range of 1 to 254 for normal f32 number.
+    //   2.1. If e >= 143, i.e. abs(v) >= 2^16 > f16::kHighest = 0x1.ffcp15, v is larger than any
+    //   finite f16 value and can not be in set R.
+    //   2.2. If 142 >= e >= 113, or f16::kHighest >= abs(v) >= f16::kSmallest = 2^-14, v falls in
+    //   the range of normal f16 values. In this case, v is in the set R iff the lowest 13 mantissa
+    //   bits are all 0. (See below for proof)
+    //     2.2.1. If we let v' be v with lowest 13 mantissa bits masked to 0, v' will be in set R
+    //     and the largest one in set R that no larger than v. Such v' is the quantized value of v.
+    //   2.3. If 112 >= e >= 103, i.e. 2^-14 > abs(v) >= f16::kSmallestSubnormal = 2^-24, v falls in
+    //   the range of subnormal f16 values. In this case, v is in the set R iff the lowest 126-e
+    //   mantissa bits are all 0. Notice that 126-e is in range 14 to 23, inclusive. (See below for
+    //   proof)
+    //     2.3.1. If we let v' be v with lowest 126-e mantissa bits masked to 0, v' will be in set R
+    //     and the largest on in set R that no larger than v. Such v' is the quantized value of v.
+    //   2.4. If 2^-24 > abs(v) > 0, i.e. 103 > e, v is smaller than any finite f16 value and not
+    //   equal to 0.0, thus can not be in set R.
+    //   2.5. If abs(v) = 0, v is in set R and is just +-0.0.
+    //
+    // Proof for 2.2:
+    // Any normal f16 number, in binary form, s_eeeee_mmmmmmmmmm, has value
+    //      (s==0?1:-1)*(1+uint(mmmmm_mmmmm)*(2^-10))*2^(uint(eeeee)-15)
+    // in which unit(bbbbb) means interprete binary pattern "bbbbb" as unsigned binary number,
+    // and we have 1 <= uint(eeeee) <= 30.
+    // This value is equal to a normal f32 number with binary
+    //      s_EEEEEEEE_mmmmmmmmmm0000000000000
+    // where uint(EEEEEEEE) = uint(eeeee) + 112, so that unbiased exponent keep unchanged
+    //      uint(EEEEEEEE) - 127 = uint(eeeee) - 15
+    // and its value is
+    //         (s==0?1:-1)*(1+uint(mmmmm_mmmmm_00000_00000_000)*(2^-23))*2^(uint(EEEEEEEE)-127)
+    //      == (s==0?1:-1)*(1+uint(mmmmm_mmmmm)*(2^-10))*2^(uint(eeeee)-15)
+    // Notice that uint(EEEEEEEE) is in range [113, 142], showing that it is a normal f32 number.
+    // So we proof that any normal f16 number can be exactly representd by a normal f32 number
+    // with biased exponent in range [113,142] and the lowest 13 mantissa bits 0.
+    // On the other hand, since mantissa bits mmmmmmmmmm are arbitrary, the value of any f32
+    // that has a biased exponent in range [113, 142] and lowest 13 mantissa bits zero is equal
+    // to a normal f16 value. Hence we proof 2.2.
+    //
+    // Proof for 2.3:
+    // Any subnormal f16 number has a binary form of s_00000_mmmmmmmmmm, and its value is
+    // (s==0?1:-1)*uint(mmmmmmmmmm)*(2^-10)*(2^-14) = (s==0?1:-1)*uint(mmmmmmmmmm)*(2^-24).
+    // We discuss on bit pattern of mantissa bits mmmmmmmmmm.
+    //   Case 1: mantissa bits has no leading zero bit, s_00000_1mmmmmmmmm
+    //      In this case the value is
+    //              (s==0?1:-1)*uint(1mmmm_mmmmm)*(2^-10)*(2^-14)
+    //          ==  (s==0?1:-1)*(uint(1_mmmmm_mmmm)*(2^-9))*(2^-15)
+    //          ==  (s==0?1:-1)*(1+uint(mmmmm_mmmm)*(2^-9))*(2^-15)
+    //          ==  (s==0?1:-1)*(1+uint(mmmmm_mmmm0_00000_00000_000)*(2^-23))*(2^-15)
+    //      which is equal to the value of normal f32 number
+    //          s_EEEEEEEE_mmmmm_mmmm0_00000_00000_000
+    //      where uint(EEEEEEEE) = -15 + 127 = 112. Hence we proof that any subnormal f16 number
+    //      with no leading zero mantissa bit can be exactly represented by a f32 number with
+    //      biased exponent 112 and the lowest 14 mantissa bits zero, and the value of any f32
+    //      number with biased exponent 112 and the lowest 14 mantissa bits zero are equal to a
+    //      subnormal f16 number with no leading zero mantissa bit.
+    //   Case 2: mantissa bits has 1 leading zero bit, s_00000_01mmmmmmmm
+    //      In this case the value is
+    //              (s==0?1:-1)*uint(01mmm_mmmmm)*(2^-10)*(2^-14)
+    //          ==  (s==0?1:-1)*(uint(01_mmmmm_mmm)*(2^-8))*(2^-16)
+    //          ==  (s==0?1:-1)*(1+uint(mmmmm_mmm)*(2^-8))*(2^-16)
+    //          ==  (s==0?1:-1)*(1+uint(mmmmm_mmm00_00000_00000_000)*(2^-23))*(2^-16)
+    //      which is equal to the value of normal f32 number
+    //          s_EEEEEEEE_mmmmm_mmm00_00000_00000_000
+    //      where uint(EEEEEEEE) = -16 + 127 = 111. Hence we proof that any subnormal f16 number
+    //      with 1 leading zero mantissa bit can be exactly represented by a f32 number with
+    //      biased exponent 111 and the lowest 15 mantissa bits zero, and the value of any f32
+    //      number with biased exponent 111 and the lowest 15 mantissa bits zero are equal to a
+    //      subnormal f16 number with 1 leading zero mantissa bit.
+    //   Case 3 to case 8: ......
+    //   Case 9: mantissa bits has 8 leading zero bit, s_00000_000000001m
+    //      In this case the value is
+    //              (s==0?1:-1)*uint(00000_0001m)*(2^-10)*(2^-14)
+    //          ==  (s==0?1:-1)*(uint(000000001_m)*(2^-1))*(2^-23)
+    //          ==  (s==0?1:-1)*(1+uint(m)*(2^-1))*(2^-23)
+    //          ==  (s==0?1:-1)*(1+uint(m0000_00000_00000_00000_000)*(2^-23))*(2^-23)
+    //      which is equal to the value of normal f32 number
+    //          s_EEEEEEEE_m0000_00000_00000_00000_000
+    //      where uint(EEEEEEEE) = -23 + 127 = 104. Hence we proof that any subnormal f16 number
+    //      with 8 leading zero mantissa bit can be exactly represented by a f32 number with
+    //      biased exponent 104 and the lowest 22 mantissa bits zero, and the value of any f32
+    //      number with biased exponent 104 and the lowest 22 mantissa bits zero are equal to a
+    //      subnormal f16 number with 8 leading zero mantissa bit.
+    //   Case 10: mantissa bits has 9 leading zero bit, s_00000_0000000001
+    //      In this case the value is just +-2^-24 = +-0x1.0p-24,
+    //      the f32 number has biased exponent 103 and all 23 mantissa bits zero.
+    //   Case 11: mantissa bits has 10 leading zero bit, s_00000_0000000000, just 0.0
+    // Concluding all these case, we proof that any subnormal f16 number with N leading zero
+    // mantissa bit can be exactly represented by a f32 number with biased exponent 112-N and the
+    // lowest 14+N mantissa bits zero, and the value of any f32 number with biased exponent 112-N (=
+    // e) and the lowest 14+N (= 126-e) mantissa bits zero are equal to a subnormal f16 number with
+    // N leading zero mantissa bit. N is in range [0, 9], so the f32 number's biased exponent e is
+    // in range [103, 112], or unbiased exponent in [-24, -15].
+
+    float abs_value = std::fabs(value);
+    if (abs_value >= kSmallest) {
+        // Value falls in the normal f16 range, quantize it to a normal f16 value by masking out
+        // lowest 13 mantissa bits.
+        u32 = u32 & ~((1u << 13) - 1);
+    } else if (abs_value >= kSmallestSubnormal) {
+        // Value should be quantized to a subnormal f16 value.
+
+        // Get the biased exponent `e` of f32 value, e.g. value 127 representing exponent 2^0.
+        uint32_t biased_exponent_original = (u32 & exponent_mask) >> 23;
+        // Since we ensure that kSmallest = 0x1f-14 > abs(value) >= kSmallestSubnormal = 0x1f-24,
+        // value will have a unbiased exponent in range -24 to -15 (inclusive), and the
+        // corresponding biased exponent in f32 is in range 103 to 112 (inclusive).
+        TINT_ASSERT(Semantic,
+                    (103 <= biased_exponent_original) && (biased_exponent_original <= 112));
+
+        // As we have proved, masking out the lowest 126-e mantissa bits of input value will result
+        // in a valid subnormal f16 value, which is exactly the required quantization result.
+        uint32_t discard_bits = 126 - biased_exponent_original;  // In range 14 to 23 (inclusive)
+        TINT_ASSERT(Semantic, (14 <= discard_bits) && (discard_bits <= 23));
+        uint32_t discard_mask = (1u << discard_bits) - 1;
+        u32 = u32 & ~discard_mask;
+    } else {
+        // value is too small that it can't even be represented as subnormal f16 number. Quantize
+        // to zero.
+        return value > 0 ? 0.0 : -0.0;
+    }
     memcpy(&value, &u32, 4);
     return value;
 }
diff --git a/src/tint/number.h b/src/tint/number.h
index b4c5ca4..7a0b13b 100644
--- a/src/tint/number.h
+++ b/src/tint/number.h
@@ -86,6 +86,10 @@
     static constexpr type kSmallest =
         std::is_integral_v<type> ? 0 : std::numeric_limits<type>::min();
 
+    /// Smallest positive subnormal value of this type, 0 for integral type.
+    static constexpr type kSmallestSubnormal =
+        std::is_integral_v<type> ? 0 : std::numeric_limits<type>::denorm_min();
+
     /// Constructor. The value is zero-initialized.
     Number() = default;
 
@@ -201,7 +205,12 @@
     static constexpr type kLowest = -65504.0f;
 
     /// Smallest positive normal value of this type.
-    static constexpr type kSmallest = 0.00006103515625f;  // 2⁻¹⁴
+    /// binary16 0_00001_0000000000, value is 2⁻¹⁴.
+    static constexpr type kSmallest = 0x1p-14f;
+
+    /// Smallest positive subnormal value of this type.
+    /// binary16 0_00000_0000000001, value is 2⁻¹⁴ * 2⁻¹⁰ = 2⁻²⁴.
+    static constexpr type kSmallestSubnormal = 0x1p-24f;
 
     /// Constructor. The value is zero-initialized.
     Number() = default;
diff --git a/src/tint/number_test.cc b/src/tint/number_test.cc
index 34b4d39..52ba4ae 100644
--- a/src/tint/number_test.cc
+++ b/src/tint/number_test.cc
@@ -52,7 +52,7 @@
 // Smallest positive normal float16 value.
 constexpr double kSmallestF16 = 0x1p-14;
 
-// Highest subnormal value for a float32.
+// Highest subnormal value for a float16.
 constexpr double kHighestF16Subnormal = 0x0.ffcp-14;
 
 constexpr double kLowestF32 = -kHighestF32;
@@ -141,6 +141,67 @@
     EXPECT_EQ(f16(inf), inf);
     EXPECT_EQ(f16(-inf), -inf);
     EXPECT_TRUE(std::isnan(f16(nan)));
+
+    // 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;
+
+    // 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);
+
+    // 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);
 }
 
 using BinaryCheckedCase = std::tuple<std::optional<AInt>, AInt, AInt>;
diff --git a/src/tint/reader/wgsl/lexer.cc b/src/tint/reader/wgsl/lexer.cc
index 58e3c85..7943894 100644
--- a/src/tint/reader/wgsl/lexer.cc
+++ b/src/tint/reader/wgsl/lexer.cc
@@ -343,12 +343,16 @@
     }
 
     bool has_f_suffix = false;
+    bool has_h_suffix = false;
     if (end < length() && matches(end, "f")) {
         end++;
         has_f_suffix = true;
+    } else if (end < length() && matches(end, "h")) {
+        end++;
+        has_h_suffix = true;
     }
 
-    if (!has_point && !has_exponent && !has_f_suffix) {
+    if (!has_point && !has_exponent && !has_f_suffix && !has_h_suffix) {
         // If it only has digits then it's an integer.
         return {};
     }
@@ -369,6 +373,14 @@
         }
     }
 
+    if (has_h_suffix) {
+        if (auto f = CheckedConvert<f16>(AFloat(value))) {
+            return {Token::Type::kFloatLiteral_H, source, static_cast<double>(f.Get())};
+        } else {
+            return {Token::Type::kError, source, "value cannot be represented as 'f16'"};
+        }
+    }
+
     if (value == HUGE_VAL || -value == HUGE_VAL) {
         return {Token::Type::kError, source, "value cannot be represented as 'abstract-float'"};
     } else {
@@ -547,6 +559,7 @@
     int64_t exponent_sign = 1;
     // If the 'p' part is present, the rest of the exponent must exist.
     bool has_f_suffix = false;
+    bool has_h_suffix = false;
     if (has_exponent) {
         // Parse the rest of the exponent.
         // (+|-)?
@@ -574,12 +587,15 @@
             end++;
         }
 
-        // Parse optional 'f' suffix.  For a hex float, it can only exist
+        // Parse optional 'f' or 'h' suffix.  For a hex float, it can only exist
         // when the exponent is present. Otherwise it will look like
         // one of the mantissa digits.
         if (end < length() && matches(end, "f")) {
             has_f_suffix = true;
             end++;
+        } else if (end < length() && matches(end, "h")) {
+            has_h_suffix = true;
+            end++;
         }
 
         if (!has_exponent_digits) {
@@ -648,7 +664,7 @@
     }
 
     if (signed_exponent >= kExponentMax || (signed_exponent == kExponentMax && mantissa != 0)) {
-        std::string type = has_f_suffix ? "f32" : "abstract-float";
+        std::string type = has_f_suffix ? "f32" : (has_h_suffix ? "f16" : "abstract-float");
         return {Token::Type::kError, source, "value cannot be represented as '" + type + "'"};
     }
 
@@ -667,14 +683,106 @@
             result_f64 > static_cast<double>(f32::kHighest)) {
             return {Token::Type::kError, source, "value cannot be represented as 'f32'"};
         }
-        // Check the value can be exactly represented (low 29 mantissa bits must be 0)
-        if (result_u64 & 0x1fffffff) {
+        // Check the value can be exactly represented, i.e. only high 23 mantissa bits are valid for
+        // normal f32 values, and less for subnormal f32 values. The rest low mantissa bits must be
+        // 0.
+        int valid_mantissa_bits = 0;
+        double abs_result_f64 = std::fabs(result_f64);
+        if (abs_result_f64 >= static_cast<double>(f32::kSmallest)) {
+            // The result shall be a normal f32 value.
+            valid_mantissa_bits = 23;
+        } else if (abs_result_f64 >= static_cast<double>(f32::kSmallestSubnormal)) {
+            // The result shall be a subnormal f32 value, represented as double.
+            // The smallest positive normal f32 is f32::kSmallest = 2^-126 = 0x1.0p-126, and the
+            //   smallest positive subnormal f32 is f32::kSmallestSubnormal = 2^-149. Thus, the
+            //   value v in range 2^-126 > v >= 2^-149 must be represented as a subnormal f32
+            //   number, but is still normal double (f64) number, and has a exponent in range -127
+            //   to -149, inclusive.
+            // A value v, if 2^-126 > v >= 2^-127, its binary32 representation will have binary form
+            //   s_00000000_1xxxxxxxxxxxxxxxxxxxxxx, having mantissa of 1 leading 1 bit and 22
+            //   arbitrary bits. Since this value is represented as normal double number, the
+            //   leading 1 bit is omitted, only the highest 22 mantissia bits can be arbitrary, and
+            //   the rest lowest 40 mantissa bits of f64 number must be zero.
+            // 2^-127 > v >= 2^-128, binary32 s_00000000_01xxxxxxxxxxxxxxxxxxxxx, having mantissa of
+            //   1 leading 0 bit, 1 leading 1 bit, and 21 arbitrary bits. The f64 representation
+            //   omits the leading 0 and 1 bits, and only the highest 21 mantissia bits can be
+            //   arbitrary.
+            // 2^-128 > v >= 2^-129, binary32 s_00000000_001xxxxxxxxxxxxxxxxxxxx, 20 arbitrary bits.
+            // ...
+            // 2^-147 > v >= 2^-148, binary32 s_00000000_0000000000000000000001x, 1 arbitrary bit.
+            // 2^-148 > v >= 2^-149, binary32 s_00000000_00000000000000000000001, 0 arbitrary bit.
+
+            // signed_exponent must be in range -149 + 1023 = 874 to -127 + 1023 = 896, inclusive
+            TINT_ASSERT(Reader, (874 <= signed_exponent) && (signed_exponent <= 896));
+            int unbiased_exponent =
+                static_cast<int>(signed_exponent) - static_cast<int>(kExponentBias);
+            TINT_ASSERT(Reader, (-149 <= unbiased_exponent) && (unbiased_exponent <= -127));
+            valid_mantissa_bits = unbiased_exponent + 149;  // 0 for -149, and 22 for -127
+        } else if (abs_result_f64 != 0.0) {
+            // The result is smaller than the smallest subnormal f32 value, but not equal to zero.
+            // Such value will never be exactly represented by f32.
             return {Token::Type::kError, source, "value cannot be exactly represented as 'f32'"};
         }
+        // Check the low 52-valid_mantissa_bits mantissa bits must be 0.
+        TINT_ASSERT(Reader, (0 <= valid_mantissa_bits) && (valid_mantissa_bits <= 23));
+        if (result_u64 & ((uint64_t(1) << (52 - valid_mantissa_bits)) - 1)) {
+            return {Token::Type::kError, source, "value cannot be exactly represented as 'f32'"};
+        }
+        return {Token::Type::kFloatLiteral_F, source, result_f64};
+    } else if (has_h_suffix) {
+        // Check value fits in f16
+        if (result_f64 < static_cast<double>(f16::kLowest) ||
+            result_f64 > static_cast<double>(f16::kHighest)) {
+            return {Token::Type::kError, source, "value cannot be represented as 'f16'"};
+        }
+        // Check the value can be exactly represented, i.e. only high 10 mantissa bits are valid for
+        // normal f16 values, and less for subnormal f16 values. The rest low mantissa bits must be
+        // 0.
+        int valid_mantissa_bits = 0;
+        double abs_result_f64 = std::fabs(result_f64);
+        if (abs_result_f64 >= static_cast<double>(f16::kSmallest)) {
+            // The result shall be a normal f16 value.
+            valid_mantissa_bits = 10;
+        } else if (abs_result_f64 >= static_cast<double>(f16::kSmallestSubnormal)) {
+            // The result shall be a subnormal f16 value, represented as double.
+            // The smallest positive normal f16 is f16::kSmallest = 2^-14 = 0x1.0p-14, and the
+            //   smallest positive subnormal f16 is f16::kSmallestSubnormal = 2^-24. Thus, the value
+            //   v in range 2^-14 > v >= 2^-24 must be represented as a subnormal f16 number, but
+            //   is still normal double (f64) number, and has a exponent in range -15 to -24,
+            //   inclusive.
+            // A value v, if 2^-14 > v >= 2^-15, its binary16 representation will have binary form
+            //   s_00000_1xxxxxxxxx, having mantissa of 1 leading 1 bit and 9 arbitrary bits. Since
+            //   this value is represented as normal double number, the leading 1 bit is omitted,
+            //   only the highest 9 mantissia bits can be arbitrary, and the rest lowest 43 mantissa
+            //   bits of f64 number must be zero.
+            // 2^-15 > v >= 2^-16, binary16 s_00000_01xxxxxxxx, having mantissa of 1 leading 0 bit,
+            //   1 leading 1 bit, and 8 arbitrary bits. The f64 representation omits the leading 0
+            //   and 1 bits, and only the highest 8 mantissia bits can be arbitrary.
+            // 2^-16 > v >= 2^-17, binary16 s_00000_001xxxxxxx, 7 arbitrary bits.
+            // ...
+            // 2^-22 > v >= 2^-23, binary16 s_00000_000000001x, 1 arbitrary bits.
+            // 2^-23 > v >= 2^-24, binary16 s_00000_0000000001, 0 arbitrary bits.
+
+            // signed_exponent must be in range -24 + 1023 = 999 to -15 + 1023 = 1008, inclusive
+            TINT_ASSERT(Reader, (999 <= signed_exponent) && (signed_exponent <= 1008));
+            int unbiased_exponent =
+                static_cast<int>(signed_exponent) - static_cast<int>(kExponentBias);
+            TINT_ASSERT(Reader, (-24 <= unbiased_exponent) && (unbiased_exponent <= -15));
+            valid_mantissa_bits = unbiased_exponent + 24;  // 0 for -24, and 9 for -15
+        } else if (abs_result_f64 != 0.0) {
+            // The result is smaller than the smallest subnormal f16 value, but not equal to zero.
+            // Such value will never be exactly represented by f16.
+            return {Token::Type::kError, source, "value cannot be exactly represented as 'f16'"};
+        }
+        // Check the low 52-valid_mantissa_bits mantissa bits must be 0.
+        TINT_ASSERT(Reader, (0 <= valid_mantissa_bits) && (valid_mantissa_bits <= 10));
+        if (result_u64 & ((uint64_t(1) << (52 - valid_mantissa_bits)) - 1)) {
+            return {Token::Type::kError, source, "value cannot be exactly represented as 'f16'"};
+        }
+        return {Token::Type::kFloatLiteral_H, source, result_f64};
     }
 
-    return {has_f_suffix ? Token::Type::kFloatLiteral_F : Token::Type::kFloatLiteral, source,
-            result_f64};
+    return {Token::Type::kFloatLiteral, source, result_f64};
 }
 
 Token Lexer::build_token_from_int_if_possible(Source source, size_t start, int32_t base) {
diff --git a/src/tint/reader/wgsl/lexer_test.cc b/src/tint/reader/wgsl/lexer_test.cc
index 16ae46d..d8decba 100644
--- a/src/tint/reader/wgsl/lexer_test.cc
+++ b/src/tint/reader/wgsl/lexer_test.cc
@@ -19,6 +19,7 @@
 #include <vector>
 
 #include "gtest/gtest.h"
+#include "src/tint/number.h"
 
 namespace tint::reader::wgsl {
 namespace {
@@ -320,6 +321,8 @@
     auto t = l.next();
     if (std::string(params.input).back() == 'f') {
         EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral_F));
+    } else if (std::string(params.input).back() == 'h') {
+        EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral_H));
     } else {
         EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral));
     }
@@ -340,6 +343,11 @@
                              FloatData{"1f", 1.0},
                              FloatData{"-0f", 0.0},
                              FloatData{"-1f", -1.0},
+                             // No decimal, with 'h' suffix
+                             FloatData{"0h", 0.0},
+                             FloatData{"1h", 1.0},
+                             FloatData{"-0h", 0.0},
+                             FloatData{"-1h", -1.0},
 
                              // Zero, with decimal.
                              FloatData{"0.0", 0.0},
@@ -354,7 +362,14 @@
                              FloatData{".0f", 0.0},
                              FloatData{"-0.0f", 0.0},
                              FloatData{"-0.f", 0.0},
-                             FloatData{"-.0", 0.0},
+                             FloatData{"-.0f", 0.0},
+                             // Zero, with decimal and 'h' suffix
+                             FloatData{"0.0h", 0.0},
+                             FloatData{"0.h", 0.0},
+                             FloatData{".0h", 0.0},
+                             FloatData{"-0.0h", 0.0},
+                             FloatData{"-0.h", 0.0},
+                             FloatData{"-.0h", 0.0},
 
                              // Non-zero with decimal
                              FloatData{"5.7", 5.7},
@@ -370,6 +385,13 @@
                              FloatData{"-5.7f", static_cast<double>(-5.7f)},
                              FloatData{"-5.f", static_cast<double>(-5.f)},
                              FloatData{"-.7f", static_cast<double>(-.7f)},
+                             // Non-zero with decimal and 'h' suffix
+                             FloatData{"5.7h", static_cast<double>(f16::Quantize(5.7f))},
+                             FloatData{"5.h", static_cast<double>(f16::Quantize(5.f))},
+                             FloatData{".7h", static_cast<double>(f16::Quantize(.7f))},
+                             FloatData{"-5.7h", static_cast<double>(f16::Quantize(-5.7f))},
+                             FloatData{"-5.h", static_cast<double>(f16::Quantize(-5.f))},
+                             FloatData{"-.7h", static_cast<double>(f16::Quantize(-.7f))},
 
                              // No decimal, with exponent
                              FloatData{"1e5", 1e5},
@@ -381,6 +403,11 @@
                              FloatData{"1E5f", static_cast<double>(1e5f)},
                              FloatData{"1e-5f", static_cast<double>(1e-5f)},
                              FloatData{"1E-5f", static_cast<double>(1e-5f)},
+                             // No decimal, with exponent and 'h' suffix
+                             FloatData{"6e4h", static_cast<double>(f16::Quantize(6e4f))},
+                             FloatData{"6E4h", static_cast<double>(f16::Quantize(6e4f))},
+                             FloatData{"1e-5h", static_cast<double>(f16::Quantize(1e-5f))},
+                             FloatData{"1E-5h", static_cast<double>(f16::Quantize(1e-5f))},
                              // With decimal and exponents
                              FloatData{"0.2e+12", 0.2e12},
                              FloatData{"1.2e-5", 1.2e-5},
@@ -393,9 +420,16 @@
                              FloatData{"2.57e23f", static_cast<double>(2.57e23f)},
                              FloatData{"2.5e+0f", static_cast<double>(2.5f)},
                              FloatData{"2.5e-0f", static_cast<double>(2.5f)},
+                             // With decimal and exponents and 'h' suffix
+                             FloatData{"0.2e+5h", static_cast<double>(f16::Quantize(0.2e5f))},
+                             FloatData{"1.2e-5h", static_cast<double>(f16::Quantize(1.2e-5f))},
+                             FloatData{"6.55e4h", static_cast<double>(f16::Quantize(6.55e4f))},
+                             FloatData{"2.5e+0h", static_cast<double>(f16::Quantize(2.5f))},
+                             FloatData{"2.5e-0h", static_cast<double>(f16::Quantize(2.5f))},
                              // Quantization
-                             FloatData{"3.141592653589793", 3.141592653589793},   // no quantization
-                             FloatData{"3.141592653589793f", 3.1415927410125732}  // f32 quantized
+                             FloatData{"3.141592653589793", 3.141592653589793},  // no quantization
+                             FloatData{"3.141592653589793f", 3.1415927410125732},  // f32 quantized
+                             FloatData{"3.141592653589793h", 3.140625}             // f16 quantized
                              ));
 
 using FloatTest_Invalid = testing::TestWithParam<const char*>;
@@ -404,7 +438,8 @@
     Lexer l(&file);
 
     auto t = l.next();
-    EXPECT_FALSE(t.Is(Token::Type::kFloatLiteral));
+    EXPECT_FALSE(t.Is(Token::Type::kFloatLiteral) || t.Is(Token::Type::kFloatLiteral_F) ||
+                 t.Is(Token::Type::kFloatLiteral_H));
 }
 INSTANTIATE_TEST_SUITE_P(LexerTest,
                          FloatTest_Invalid,
@@ -423,9 +458,8 @@
                                          // Overflow
                                          "2.5e+256f",
                                          "-2.5e+127f",
-                                         // Magnitude smaller than smallest positive f32.
-                                         "2.5e-300f",
-                                         "-2.5e-300f",
+                                         "6.5520e+4h",
+                                         "-6.5e+12h",
                                          // Decimal exponent must immediately
                                          // follow the 'e'.
                                          "2.5e 12",
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index a28b798..ddd9e96 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -3018,6 +3018,10 @@
         return create<ast::FloatLiteralExpression>(t.source(), t.to_f64(),
                                                    ast::FloatLiteralExpression::Suffix::kF);
     }
+    if (match(Token::Type::kFloatLiteral_H)) {
+        return create<ast::FloatLiteralExpression>(t.source(), t.to_f64(),
+                                                   ast::FloatLiteralExpression::Suffix::kH);
+    }
     if (match(Token::Type::kTrue)) {
         return create<ast::BoolLiteralExpression>(t.source(), true);
     }
@@ -3251,6 +3255,7 @@
         });
     }
 
+    // TODO(crbug.com/tint/1503): Remove when deprecation period is over.
     if (t == kStageAttribute) {
         return expect_paren_block("stage attribute", [&]() -> Result {
             auto stage = expect_pipeline_stage();
@@ -3258,26 +3263,22 @@
                 return Failure::kErrored;
             }
 
-            // TODO(crbug.com/tint/1503): Enable this once all the Dawn and CTS
-            // tests are updated to use the new format so we can avoid spamming
-            // the log files.
-            if ((false)) {
-                std::string warning = "stage should use @";
-                switch (stage.value) {
-                    case ast::PipelineStage::kVertex:
-                        warning += "vertex";
-                        break;
-                    case ast::PipelineStage::kFragment:
-                        warning += "fragment";
-                        break;
-                    case ast::PipelineStage::kCompute:
-                        warning += "compute";
-                        break;
-                    case ast::PipelineStage::kNone:
-                        break;
-                }
-                deprecated(t.source(), warning);
+            std::string warning = "remove stage and use @";
+            switch (stage.value) {
+                case ast::PipelineStage::kVertex:
+                    warning += "vertex";
+                    break;
+                case ast::PipelineStage::kFragment:
+                    warning += "fragment";
+                    break;
+                case ast::PipelineStage::kCompute:
+                    warning += "compute";
+                    break;
+                case ast::PipelineStage::kNone:
+                    break;
             }
+            deprecated(t.source(), warning);
+
             return create<ast::StageAttribute>(t.source(), stage.value);
         });
     }
diff --git a/src/tint/reader/wgsl/parser_impl_const_literal_test.cc b/src/tint/reader/wgsl/parser_impl_const_literal_test.cc
index b07ec1b..16861e4 100644
--- a/src/tint/reader/wgsl/parser_impl_const_literal_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_const_literal_test.cc
@@ -151,13 +151,18 @@
     ASSERT_NE(c.value, nullptr);
     auto* literal = c->As<ast::FloatLiteralExpression>();
     ASSERT_NE(literal, nullptr);
-    EXPECT_DOUBLE_EQ(literal->value, params.expected)
+    // Use EXPECT_EQ instead of EXPECT_DOUBLE_EQ here, because EXPECT_DOUBLE_EQ use AlmostEquals(),
+    // which allows an error up to 4 ULPs.
+    EXPECT_EQ(literal->value, params.expected)
         << "\n"
         << "got:      " << std::hexfloat << literal->value << "\n"
         << "expected: " << std::hexfloat << params.expected;
     if (params.input.back() == 'f') {
         EXPECT_EQ(c->As<ast::FloatLiteralExpression>()->suffix,
                   ast::FloatLiteralExpression::Suffix::kF);
+    } else if (params.input.back() == 'h') {
+        EXPECT_EQ(c->As<ast::FloatLiteralExpression>()->suffix,
+                  ast::FloatLiteralExpression::Suffix::kH);
     } else {
         EXPECT_EQ(c->As<ast::FloatLiteralExpression>()->suffix,
                   ast::FloatLiteralExpression::Suffix::kNone);
@@ -181,6 +186,7 @@
 
                              {"234.e12", 234.e12},
                              {"234.e12f", static_cast<double>(234.e12f)},
+                             {"234.e2h", static_cast<double>(f16::Quantize(234.e2))},
 
                              // Tiny cases
                              {"1e-5000", 0.0},
@@ -189,6 +195,12 @@
                              {"-1e-5000f", 0.0},
                              {"1e-50f", 0.0},
                              {"-1e-50f", 0.0},
+                             {"1e-5000h", 0.0},
+                             {"-1e-5000h", 0.0},
+                             {"1e-50h", 0.0},
+                             {"-1e-50h", 0.0},
+                             {"1e-8h", 0.0},  // The smallest positive subnormal f16 is 5.96e-8
+                             {"-1e-8h", 0.0},
 
                              // Nearly overflow
                              {"1.e308", 1.e308},
@@ -209,6 +221,16 @@
                              {"-3.5e37f", static_cast<double>(-3.5e37f)},
                              {"3.403e37f", static_cast<double>(3.403e37f)},
                              {"-3.403e37f", static_cast<double>(-3.403e37f)},
+
+                             // Nearly overflow
+                             {"6e4h", 6e4},
+                             {"-6e4h", -6e4},
+                             {"8.0e3h", 8.0e3},
+                             {"-8.0e3h", -8.0e3},
+                             {"3.5e3h", 3.5e3},
+                             {"-3.5e3h", -3.5e3},
+                             {"3.403e3h", 3.402e3},    // Quantized
+                             {"-3.403e3h", -3.402e3},  // Quantized
                          }));
 
 const double NegInf = MakeDouble(1, 0x7FF, 0);
@@ -229,6 +251,10 @@
         {"-0x1p-1", -0x1p-1},
         {"-0x1p-2", -0x1p-2},
         {"-0x1.8p-1", -0x1.8p-1},
+        {"0x0.4p+1", 0x0.4p+1},
+        {"0x0.02p+3", 0x0.02p+3},
+        {"0x4.4p+1", 0x4.4p+1},
+        {"0x8c.02p+3", 0x8c.02p+3},
 
         // Large numbers
         {"0x1p+9", 0x1p+9},
@@ -257,6 +283,11 @@
         {"-0x1p-124f", -0x1p-124},
         {"-0x1p-125f", -0x1p-125},
 
+        {"0x1p-12h", 0x1p-12},
+        {"0x1p-13h", 0x1p-13},
+        {"-0x1p-12h", -0x1p-12},
+        {"-0x1p-13h", -0x1p-13},
+
         // Lowest non-denorm
         {"0x1p-1022", 0x1p-1022},
         {"-0x1p-1022", -0x1p-1022},
@@ -264,9 +295,14 @@
         {"0x1p-126f", 0x1p-126},
         {"-0x1p-126f", -0x1p-126},
 
+        {"0x1p-14h", 0x1p-14},
+        {"-0x1p-14h", -0x1p-14},
+
         // Denormalized values
         {"0x1p-1023", 0x1p-1023},
+        {"0x0.8p-1022", 0x0.8p-1022},
         {"0x1p-1024", 0x1p-1024},
+        {"0x0.2p-1021", 0x0.2p-1021},
         {"0x1p-1025", 0x1p-1025},
         {"0x1p-1026", 0x1p-1026},
         {"-0x1p-1023", -0x1p-1023},
@@ -277,7 +313,9 @@
         {"0x1.8p-1024", 0x1.8p-1024},
 
         {"0x1p-127f", 0x1p-127},
+        {"0x0.8p-126f", 0x0.8p-126},
         {"0x1p-128f", 0x1p-128},
+        {"0x0.2p-125f", 0x0.2p-125},
         {"0x1p-129f", 0x1p-129},
         {"0x1p-130f", 0x1p-130},
         {"-0x1p-127f", -0x1p-127},
@@ -287,13 +325,28 @@
         {"0x1.8p-127f", 0x1.8p-127},
         {"0x1.8p-128f", 0x1.8p-128},
 
+        {"0x1p-15h", 0x1p-15},
+        {"0x0.8p-14h", 0x0.8p-14},
+        {"0x1p-16h", 0x1p-16},
+        {"0x0.2p-13h", 0x0.2p-13},
+        {"0x1p-17h", 0x1p-17},
+        {"0x1p-18h", 0x1p-18},
+        {"-0x1p-15h", -0x1p-15},
+        {"-0x1p-16h", -0x1p-16},
+        {"-0x1p-17h", -0x1p-17},
+        {"-0x1p-18h", -0x1p-18},
+        {"0x1.8p-15h", 0x1.8p-15},
+        {"0x1.8p-16h", 0x1.8p-16},
+
         // F64 extremities
-        {"0x1p-1074", 0x1p-1074},                              // +SmallestDenormal
-        {"0x1p-1073", 0x1p-1073},                              // +BiggerDenormal
-        {"0x1.ffffffffffffp-1027", 0x1.ffffffffffffp-1027},    // +LargestDenormal
-        {"-0x1p-1074", -0x1p-1074},                            // -SmallestDenormal
-        {"-0x1p-1073", -0x1p-1073},                            // -BiggerDenormal
-        {"-0x1.ffffffffffffp-1027", -0x1.ffffffffffffp-1027},  // -LargestDenormal
+        {"0x1p-1074", 0x1p-1074},                                // +SmallestDenormal
+        {"0x1p-1073", 0x1p-1073},                                // +BiggerDenormal
+        {"0x1.ffffffffffffep-1023", 0x1.ffffffffffffep-1023},    // +LargestDenormal
+        {"0x0.fffffffffffffp-1022", 0x0.fffffffffffffp-1022},    // +LargestDenormal
+        {"-0x1p-1074", -0x1p-1074},                              // -SmallestDenormal
+        {"-0x1p-1073", -0x1p-1073},                              // -BiggerDenormal
+        {"-0x1.ffffffffffffep-1023", -0x1.ffffffffffffep-1023},  // -LargestDenormal
+        {"-0x0.fffffffffffffp-1022", -0x0.fffffffffffffp-1022},  // -LargestDenormal
 
         {"0x0.cafebeeff000dp-1022", 0x0.cafebeeff000dp-1022},    // +Subnormal
         {"-0x0.cafebeeff000dp-1022", -0x0.cafebeeff000dp-1022},  // -Subnormal
@@ -301,21 +354,29 @@
         {"-0x1.2bfaf8p-1052", -0x1.2bfaf8p-1052},                // +Subnormal
         {"0x1.55554p-1055", 0x1.55554p-1055},                    // +Subnormal
         {"-0x1.55554p-1055", -0x1.55554p-1055},                  // -Subnormal
+        {"0x1.fffffffffffp-1027", 0x1.fffffffffffp-1027},  // +Subnormal, = 0x0.0fffffffffff8p-1022
+        {"-0x1.fffffffffffp-1027", -0x1.fffffffffffp-1027},  // -Subnormal
 
         // F32 extremities
-        {"0x1p-149", 0x1p-149},                  // +SmallestDenormal
-        {"0x1p-148", 0x1p-148},                  // +BiggerDenormal
-        {"0x1.fffffcp-127", 0x1.fffffcp-127},    // +LargestDenormal
-        {"-0x1p-149", -0x1p-149},                // -SmallestDenormal
-        {"-0x1p-148", -0x1p-148},                // -BiggerDenormal
-        {"-0x1.fffffcp-127", -0x1.fffffcp-127},  // -LargestDenormal
+        {"0x1p-149f", 0x1p-149},                  // +SmallestDenormal
+        {"0x1p-148f", 0x1p-148},                  // +BiggerDenormal
+        {"0x1.fffffcp-127f", 0x1.fffffcp-127},    // +LargestDenormal
+        {"0x0.fffffep-126f", 0x0.fffffep-126},    // +LargestDenormal
+        {"0x1.0p-126f", 0x1.0p-126},              // +SmallestNormal
+        {"0x8.0p-129f", 0x8.0p-129},              // +SmallestNormal
+        {"-0x1p-149f", -0x1p-149},                // -SmallestDenormal
+        {"-0x1p-148f", -0x1p-148},                // -BiggerDenormal
+        {"-0x1.fffffcp-127f", -0x1.fffffcp-127},  // -LargestDenormal
+        {"-0x0.fffffep-126f", -0x0.fffffep-126},  // -LargestDenormal
+        {"-0x1.0p-126f", -0x1.0p-126},            // -SmallestNormal
+        {"-0x8.0p-129f", -0x8.0p-129},            // -SmallestNormal
 
-        {"0x0.cafebp-129", 0x0.cafebp-129},      // +Subnormal
-        {"-0x0.cafebp-129", -0x0.cafebp-129},    // -Subnormal
-        {"0x1.2bfaf8p-127", 0x1.2bfaf8p-127},    // +Subnormal
-        {"-0x1.2bfaf8p-127", -0x1.2bfaf8p-127},  // -Subnormal
-        {"0x1.55554p-130", 0x1.55554p-130},      // +Subnormal
-        {"-0x1.55554p-130", -0x1.55554p-130},    // -Subnormal
+        {"0x0.cafebp-129f", 0x0.cafebp-129},      // +Subnormal
+        {"-0x0.cafebp-129f", -0x0.cafebp-129},    // -Subnormal
+        {"0x1.2bfaf8p-127f", 0x1.2bfaf8p-127},    // +Subnormal
+        {"-0x1.2bfaf8p-127f", -0x1.2bfaf8p-127},  // -Subnormal
+        {"0x1.55554p-130f", 0x1.55554p-130},      // +Subnormal
+        {"-0x1.55554p-130f", -0x1.55554p-130},    // -Subnormal
 
         // F32 exactly representable
         {"0x1.000002p+0f", 0x1.000002p+0},
@@ -324,10 +385,47 @@
         {"0x8.00003p+0f", 0x8.00003p+0},
         {"0x2.123p+0f", 0x2.123p+0},
         {"0x2.cafefp+0f", 0x2.cafefp+0},
+        {"0x0.0000fep-126f", 0x0.0000fep-126},    // Subnormal
+        {"-0x0.0000fep-126f", -0x0.0000fep-126},  // Subnormal
+        {"0x3.f8p-144f", 0x3.f8p-144},            // Subnormal
+        {"-0x3.f8p-144f", -0x3.f8p-144},          // Subnormal
+
+        // F16 extremities
+        {"0x1p-24h", 0x1p-24},            // +SmallestDenormal
+        {"0x1p-23h", 0x1p-23},            // +BiggerDenormal
+        {"0x1.ff8p-15h", 0x1.ff8p-15},    // +LargestDenormal
+        {"0x0.ffcp-14h", 0x0.ffcp-14},    // +LargestDenormal
+        {"0x1.0p-14h", 0x1.0p-14},        // +SmallestNormal
+        {"0x8.0p-17h", 0x8.0p-17},        // +SmallestNormal
+        {"-0x1p-24h", -0x1p-24},          // -SmallestDenormal
+        {"-0x1p-23h", -0x1p-23},          // -BiggerDenormal
+        {"-0x1.ff8p-15h", -0x1.ff8p-15},  // -LargestDenormal
+        {"-0x0.ffcp-14h", -0x0.ffcp-14},  // -LargestDenormal
+        {"-0x1.0p-14h", -0x1.0p-14},      // -SmallestNormal
+        {"-0x8.0p-17h", -0x8.0p-17},      // -SmallestNormal
+
+        {"0x0.a8p-19h", 0x0.a8p-19},    // +Subnormal
+        {"-0x0.a8p-19h", -0x0.a8p-19},  // -Subnormal
+        {"0x1.7ap-17h", 0x1.7ap-17},    // +Subnormal
+        {"-0x1.7ap-17h", -0x1.7ap-17},  // -Subnormal
+        {"0x1.dp-20h", 0x1.dp-20},      // +Subnormal
+        {"-0x1.dp-20h", -0x1.dp-20},    // -Subnormal
+
+        // F16 exactly representable
+        {"0x1.004p+0h", 0x1.004p+0},
+        {"0x8.02p+0h", 0x8.02p+0},
+        {"0x8.fep+0h", 0x8.fep+0},
+        {"0x8.06p+0h", 0x8.06p+0},
+        {"0x2.128p+0h", 0x2.128p+0},
+        {"0x2.ca8p+0h", 0x2.ca8p+0},
+        {"0x0.0fcp-14h", 0x0.0fcp-14},    // Subnormal
+        {"-0x0.0fcp-14h", -0x0.0fcp-14},  // Subnormal
+        {"0x3.f00p-20h", 0x3.f00p-20},    // Subnormal
+        {"-0x3.f00p-20h", -0x3.f00p-20},  // Subnormal
 
         // Underflow -> Zero
-        {"0x1p-1074", 0.0},  // Exponent underflows
-        {"-0x1p-1074", 0.0},
+        {"0x1p-1075", 0.0},  // Exponent underflows
+        {"-0x1p-1075", 0.0},
         {"0x1p-5000", 0.0},
         {"-0x1p-5000", 0.0},
         {"0x0.00000000000000000000001p-1022", 0.0},  // Fraction causes underflow
@@ -399,6 +497,16 @@
         {"-0x.8p2f", -2.0},
         {"-0x1.8p-1f", -0.75},
         {"-0x2p-2f", -0.5},  // No binary point
+
+        // Examples with a binary exponent and a 'h' suffix.
+        {"0x1.p0h", 1.0},
+        {"0x.8p2h", 2.0},
+        {"0x1.8p-1h", 0.75},
+        {"0x2p-2h", 0.5},  // No binary point
+        {"-0x1.p0h", -1.0},
+        {"-0x.8p2h", -2.0},
+        {"-0x1.8p-1h", -0.75},
+        {"-0x2p-2h", -0.5},  // No binary point
     };
 }
 INSTANTIATE_TEST_SUITE_P(ParserImplFloatLiteralTest_HexFloat,
@@ -542,6 +650,23 @@
                      })));
 
 INSTANTIATE_TEST_SUITE_P(
+    HexNaNF16,
+    ParserImplInvalidLiteralTest,
+    testing::Combine(testing::Values("1:1: value cannot be represented as 'f16'"),
+                     testing::ValuesIn(std::vector<const char*>{
+                         "0x1.8p+16h",
+                         "0x1.004p+16h",
+                         "0x1.018p+16h",
+                         "0x1.1ep+16h",
+                         "0x1.ffcp+16h",
+                         "-0x1.8p+16h",
+                         "-0x1.004p+16h",
+                         "-0x1.018p+16h",
+                         "-0x1.1ep+16h",
+                         "-0x1.ffcp+16h",
+                     })));
+
+INSTANTIATE_TEST_SUITE_P(
     HexOverflowAFloat,
     ParserImplInvalidLiteralTest,
     testing::Combine(testing::Values("1:1: value cannot be represented as 'abstract-float'"),
@@ -578,17 +703,94 @@
                      })));
 
 INSTANTIATE_TEST_SUITE_P(
+    HexOverflowF16,
+    ParserImplInvalidLiteralTest,
+    testing::Combine(testing::Values("1:1: value cannot be represented as 'f16'"),
+                     testing::ValuesIn(std::vector<const char*>{
+                         "0x1p+16h",
+                         "-0x1p+16h",
+                         "0x1.1p+16h",
+                         "-0x1.1p+16h",
+                         "0x1p+17h",
+                         "-0x1p+17h",
+                         "0x32p+15h",
+                         "-0x32p+15h",
+                         "0x32p+500h",
+                         "-0x32p+500h",
+                     })));
+
+INSTANTIATE_TEST_SUITE_P(
     HexNotExactlyRepresentableF32,
     ParserImplInvalidLiteralTest,
     testing::Combine(testing::Values("1:1: value cannot be exactly represented as 'f32'"),
                      testing::ValuesIn(std::vector<const char*>{
-                         "0x1.000001p+0f",    // Quantizes to 0x1.0p+0
-                         "0x8.0000f8p+0f",    // Quantizes to 0x8.0000fp+0
-                         "0x8.000038p+0f",    // Quantizes to 0x8.00003p+0
-                         "0x2.cafef00dp+0f",  // Quantizes to 0x2.cafefp+0
+                         "0x1.000001p+0f",            // Quantizes to 0x1.0p+0
+                         "0x1.0000008p+0f",           // Quantizes to 0x1.0p+0
+                         "0x1.0000000000001p+0f",     // Quantizes to 0x1.0p+0
+                         "0x8.0000f8p+0f",            // Quantizes to 0x8.0000fp+0
+                         "0x8.000038p+0f",            // Quantizes to 0x8.00003p+0
+                         "0x2.cafef00dp+0f",          // Quantizes to 0x2.cafefp+0
+                         "0x0.0000ffp-126f",          // Subnormal, quantizes to 0x0.0000fep-126
+                         "0x3.fcp-144f",              // Subnormal, quantizes to 0x3.f8p-144
+                         "-0x0.0000ffp-126f",         // Subnormal, quantizes to -0x0.0000fep-126
+                         "-0x3.fcp-144f",             // Subnormal, quantizes to -0x3.f8p-144
+                         "0x0.ffffffp-126f",          // Subnormal, quantizes to 0x0.fffffep-144
+                         "0x0.fffffe0000001p-126f",   // Subnormal, quantizes to 0x0.fffffep-144
+                         "-0x0.ffffffp-126f",         // Subnormal, quantizes to -0x0.fffffep-144
+                         "-0x0.fffffe0000001p-126f",  // Subnormal, quantizes to -0x0.fffffep-144
+                         "0x1.8p-149f",               // Subnormal, quantizes to 0x1.0p-149f
+                         "0x1.4p-149f",               // Subnormal, quantizes to 0x1.0p-149f
+                         "0x1.000002p-149f",          // Subnormal, quantizes to 0x1.0p-149f
+                         "0x1.0000000000001p-149f",   // Subnormal, quantizes to 0x1.0p-149f
+                         "-0x1.8p-149f",              // Subnormal, quantizes to -0x1.0p-149f
+                         "-0x1.4p-149f",              // Subnormal, quantizes to -0x1.0p-149f
+                         "-0x1.000002p-149f",         // Subnormal, quantizes to -0x1.0p-149f
+                         "-0x1.0000000000001p-149f",  // Subnormal, quantizes to -0x1.0p-149f
+                         "0x1.0p-150f",   // Smaller than the smallest subnormal, quantizes to 0.0
+                         "0x1.8p-150f",   // Smaller than the smallest subnormal, quantizes to 0.0
+                         "-0x1.0p-150f",  // Smaller than the smallest subnormal, quantizes to -0.0
+                         "-0x1.8p-150f",  // Smaller than the smallest subnormal, quantizes to -0.0
                      })));
 
 INSTANTIATE_TEST_SUITE_P(
+    HexNotExactlyRepresentableF16,
+    ParserImplInvalidLiteralTest,
+    testing::Combine(
+        testing::Values("1:1: value cannot be exactly represented as 'f16'"),
+        testing::ValuesIn(std::vector<const char*>{
+            "0x1.002p+0h",             // Quantizes to 0x1.0p+0, has 11 mantissa bits rather than 10
+            "0x1.001p+0h",             // Quantizes to 0x1.0p+0, has 12 mantissa bits rather than 10
+            "0x1.0000000000001p+0h",   // Quantizes to 0x1.0p+0, has 52 mantissa bits rather than 10
+            "0x8.0fp+0h",              // Quantizes to 0x8.0ep+0
+            "0x8.31p+0h",              // Quantizes to 0x8.30p+0
+            "0x2.ca80dp+0h",           // Quantizes to 0x2.ca8p+0
+            "0x4.ba8p+0h",             // Quantizes to 0x4.bap+0
+            "0x4.011p+0h",             // Quantizes to 0x4.01p+0
+            "0x0.0fep-14h",            // Subnormal, quantizes to 0x0.0fcp-14
+            "0x3.f8p-20h",             // Subnormal, quantizes to 0x3.f0p-20
+            "-0x0.0fep-14h",           // Subnormal, quantizes to -0x0.0fcp-14
+            "-0x3.f8p-20h",            // Subnormal, quantizes to -0x3.f0p-20
+            "0x0.ffep-14h",            // Subnormal, quantizes to 0x0.ffcp-14
+            "0x0.ffe0000000001p-14h",  // Subnormal, quantizes to 0x0.ffcp-14
+            "0x0.fffffffffffffp-14h",  // Subnormal, quantizes to 0x0.ffcp-14
+            "-0x0.ffep-14h",           // Subnormal, quantizes to -0x0.ffcp-14
+            "-0x0.ffe0000000001p-14h",  // Subnormal, quantizes to -0x0.ffcp-14
+            "-0x0.fffffffffffffp-14h",  // Subnormal, quantizes to -0x0.ffcp-14
+            "0x1.8p-24h",               // Subnormal, quantizes to 0x1.0p-24f
+            "0x1.4p-24h",               // Subnormal, quantizes to 0x1.0p-24f
+            "0x1.004p-24h",             // Subnormal, quantizes to 0x1.0p-24f
+            "0x1.0000000000001p-24h",   // Subnormal, quantizes to 0x1.0p-24f
+            "-0x1.8p-24h",              // Subnormal, quantizes to -0x1.0p-24f
+            "-0x1.4p-24h",              // Subnormal, quantizes to -0x1.0p-24f
+            "-0x1.004p-24h",            // Subnormal, quantizes to -0x1.0p-24f
+            "-0x1.0000000000001p-24h",  // Subnormal, quantizes to -0x1.0p-24f
+            "0x1.0p-25h",               // Smaller than the smallest subnormal, quantizes to 0.0
+            "0x1.8p-25h",               // Smaller than the smallest subnormal, quantizes to 0.0
+            "-0x1.0p-25h",              // Smaller than the smallest subnormal, quantizes to -0.0
+            "-0x1.8p-25h",              // Smaller than the smallest subnormal, quantizes to -0.0
+        })));
+
+INSTANTIATE_TEST_SUITE_P(
     DecOverflowAFloat,
     ParserImplInvalidLiteralTest,
     testing::Combine(testing::Values("1:1: value cannot be represented as 'abstract-float'"),
@@ -622,6 +824,25 @@
                          "-1.2e+256f",
                      })));
 
+INSTANTIATE_TEST_SUITE_P(
+    DecOverflowF16,
+    ParserImplInvalidLiteralTest,
+    testing::Combine(testing::Values("1:1: value cannot be represented as 'f16'"),
+                     testing::ValuesIn(std::vector<const char*>{
+                         "1.0e5h",
+                         "-1.0e5h",
+                         "7.0e4h",
+                         "-7.0e4h",
+                         "6.6e4h",
+                         "-6.6e4h",
+                         "6.56e4h",
+                         "-6.56e4h",
+                         "6.554e4h",
+                         "-6.554e4h",
+                         "1.2e+32h",
+                         "-1.2e+32h",
+                     })));
+
 TEST_F(ParserImplTest, ConstLiteral_FloatHighest) {
     const auto highest = std::numeric_limits<float>::max();
     const auto expected_highest = 340282346638528859811704183484516925440.0f;
@@ -636,8 +857,7 @@
     EXPECT_FALSE(p->has_error()) << p->error();
     ASSERT_NE(c.value, nullptr);
     ASSERT_TRUE(c->Is<ast::FloatLiteralExpression>());
-    EXPECT_DOUBLE_EQ(c->As<ast::FloatLiteralExpression>()->value,
-                     std::numeric_limits<float>::max());
+    EXPECT_EQ(c->As<ast::FloatLiteralExpression>()->value, std::numeric_limits<float>::max());
     EXPECT_EQ(c->As<ast::FloatLiteralExpression>()->suffix,
               ast::FloatLiteralExpression::Suffix::kNone);
     EXPECT_EQ(c->source.range, (Source::Range{{1u, 1u}, {1u, 42u}}));
@@ -660,8 +880,7 @@
     EXPECT_FALSE(p->has_error()) << p->error();
     ASSERT_NE(c.value, nullptr);
     ASSERT_TRUE(c->Is<ast::FloatLiteralExpression>());
-    EXPECT_DOUBLE_EQ(c->As<ast::FloatLiteralExpression>()->value,
-                     std::numeric_limits<float>::lowest());
+    EXPECT_EQ(c->As<ast::FloatLiteralExpression>()->value, std::numeric_limits<float>::lowest());
     EXPECT_EQ(c->As<ast::FloatLiteralExpression>()->suffix,
               ast::FloatLiteralExpression::Suffix::kNone);
     EXPECT_EQ(c->source.range, (Source::Range{{1u, 1u}, {1u, 43u}}));
diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
index ab97d94..972e118 100644
--- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
@@ -316,8 +316,13 @@
 }
 
 TEST_F(ParserImplErrorTest, FunctionDeclStageMissingRParen) {
-    EXPECT("@stage(vertex fn f() {}",
-           R"(test.wgsl:1:15 error: expected ')' for stage attribute
+    EXPECT(
+        "@stage(vertex fn f() {}",
+        R"(test.wgsl:1:2 warning: use of deprecated language feature: remove stage and use @vertex
+@stage(vertex fn f() {}
+ ^^^^^
+
+test.wgsl:1:15 error: expected ')' for stage attribute
 @stage(vertex fn f() {}
               ^^
 )");
diff --git a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
index 44f3d78..a32e140 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
@@ -263,7 +263,8 @@
     EXPECT_TRUE(attr.errored);
     EXPECT_EQ(attr.value, nullptr);
     EXPECT_TRUE(p->has_error());
-    EXPECT_EQ(p->error(), "1:14: expected ')' for stage attribute");
+    EXPECT_EQ(p->error(), R"(1:1: use of deprecated language feature: remove stage and use @compute
+1:14: expected ')' for stage attribute)");
 }
 
 TEST_F(ParserImplTest, Attribute_Compute) {
diff --git a/src/tint/reader/wgsl/token.cc b/src/tint/reader/wgsl/token.cc
index 4680eee..dcd72cf 100644
--- a/src/tint/reader/wgsl/token.cc
+++ b/src/tint/reader/wgsl/token.cc
@@ -29,6 +29,8 @@
             return "abstract float literal";
         case Token::Type::kFloatLiteral_F:
             return "'f'-suffixed float literal";
+        case Token::Type::kFloatLiteral_H:
+            return "'h'-suffixed float literal";
         case Token::Type::kIntLiteral:
             return "abstract integer literal";
         case Token::Type::kIntLiteral_I:
@@ -311,6 +313,8 @@
             return std::to_string(std::get<double>(value_));
         case Type::kFloatLiteral_F:
             return std::to_string(std::get<double>(value_)) + "f";
+        case Type::kFloatLiteral_H:
+            return std::to_string(std::get<double>(value_)) + "h";
         case Type::kIntLiteral:
             return std::to_string(std::get<int64_t>(value_));
         case Type::kIntLiteral_I:
diff --git a/src/tint/reader/wgsl/token.h b/src/tint/reader/wgsl/token.h
index 0a68f9b0..9587f36 100644
--- a/src/tint/reader/wgsl/token.h
+++ b/src/tint/reader/wgsl/token.h
@@ -42,6 +42,8 @@
         kFloatLiteral,
         /// A float literal with an 'f' suffix
         kFloatLiteral_F,
+        /// A float literal with an 'h' suffix
+        kFloatLiteral_H,
         /// An integer literal with no suffix
         kIntLiteral,
         /// An integer literal with an 'i' suffix
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index 90ae4b0..b1f5a95 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -1658,10 +1658,15 @@
             return nullptr;
         },
         [&](const ast::FloatLiteralExpression* f) -> sem::Type* {
-            if (f->suffix == ast::FloatLiteralExpression::Suffix::kNone) {
-                return builder_->create<sem::AbstractFloat>();
+            switch (f->suffix) {
+                case ast::FloatLiteralExpression::Suffix::kNone:
+                    return builder_->create<sem::AbstractFloat>();
+                case ast::FloatLiteralExpression::Suffix::kF:
+                    return builder_->create<sem::F32>();
+                case ast::FloatLiteralExpression::Suffix::kH:
+                    return builder_->create<sem::F16>();
             }
-            return builder_->create<sem::F32>();
+            return nullptr;
         },
         [&](const ast::BoolLiteralExpression*) { return builder_->create<sem::Bool>(); },
         [&](Default) { return nullptr; });
@@ -1672,6 +1677,11 @@
         return nullptr;
     }
 
+    if ((ty->Is<sem::F16>()) && (!enabled_extensions_.contains(tint::ast::Extension::kF16))) {
+        AddError("f16 literal used without 'f16' extension enabled", literal->source);
+        return nullptr;
+    }
+
     auto val = EvaluateConstantValue(literal, ty);
     if (!val) {
         return nullptr;
diff --git a/src/tint/resolver/resolver_test.cc b/src/tint/resolver/resolver_test.cc
index 16725ba..38f4d49 100644
--- a/src/tint/resolver/resolver_test.cc
+++ b/src/tint/resolver/resolver_test.cc
@@ -2134,5 +2134,22 @@
                                         std::to_string(kMaxExpressionDepth)));
 }
 
+TEST_F(ResolverTest, Literal_F16WithoutExtension) {
+    // fn test() {_ = 1.23h;}
+    WrapInFunction(Ignore(Expr(f16(1.23f))));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_THAT(r()->error(), HasSubstr("error: f16 literal used without 'f16' extension enabled"));
+}
+
+TEST_F(ResolverTest, Literal_F16WithExtension) {
+    // enable f16;
+    // fn test() {_ = 1.23h;}
+    Enable(ast::Extension::kF16);
+    WrapInFunction(Ignore(Expr(f16(1.23f))));
+
+    EXPECT_TRUE(r()->Resolve());
+}
+
 }  // namespace
 }  // namespace tint::resolver
diff --git a/src/tint/resolver/uniformity.cc b/src/tint/resolver/uniformity.cc
index 97612a4..2bbb487 100644
--- a/src/tint/resolver/uniformity.cc
+++ b/src/tint/resolver/uniformity.cc
@@ -305,7 +305,7 @@
     /// @param ast the optional AST node that this node corresponds to
     /// @returns the new node
     Node* CreateNode(std::string tag, const ast::Node* ast = nullptr) {
-        return current_function_->CreateNode(tag, ast);
+        return current_function_->CreateNode(std::move(tag), ast);
     }
 
     /// Process a function.
@@ -1248,6 +1248,10 @@
                     if (func_info->parameters[i].pointer_may_become_non_uniform) {
                         ptr_result->AddEdge(current_function_->may_be_non_uniform);
                     } else {
+                        // Add edge to the call to catch when it's called in non-uniform control
+                        // flow.
+                        ptr_result->AddEdge(call_node);
+
                         // Add edges from the resulting pointer value to any other arguments that
                         // feed it.
                         for (auto* source : func_info->parameters[i].pointer_param_output_sources) {
diff --git a/src/tint/resolver/uniformity_test.cc b/src/tint/resolver/uniformity_test.cc
index 7ed4a6b..d661099 100644
--- a/src/tint/resolver/uniformity_test.cc
+++ b/src/tint/resolver/uniformity_test.cc
@@ -4624,6 +4624,145 @@
 )");
 }
 
+TEST_F(UniformityAnalysisTest, PointerParamModifiedInNonUniformControlFlow) {
+    std::string src = R"(
+@binding(0) @group(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo(p : ptr<function, i32>) {
+  *p = 42;
+}
+
+@compute @workgroup_size(64)
+fn main() {
+  var a : i32;
+  if (non_uniform_global == 0) {
+    foo(&a);
+  }
+
+  if (a == 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, false);
+    EXPECT_EQ(error_,
+              R"(test:16:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+    workgroupBarrier();
+    ^^^^^^^^^^^^^^^^
+
+test:11:3 note: control flow depends on non-uniform value
+  if (non_uniform_global == 0) {
+  ^^
+
+test:11:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
+  if (non_uniform_global == 0) {
+      ^^^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamAssumedModifiedInNonUniformControlFlow) {
+    std::string src = R"(
+@binding(0) @group(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo(p : ptr<function, i32>) {
+  // Do not modify 'p', uniformity analysis presently assumes it will be.
+}
+
+@compute @workgroup_size(64)
+fn main() {
+  var a : i32;
+  if (non_uniform_global == 0) {
+    foo(&a);
+  }
+
+  if (a == 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, false);
+    EXPECT_EQ(error_,
+              R"(test:16:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+    workgroupBarrier();
+    ^^^^^^^^^^^^^^^^
+
+test:11:3 note: control flow depends on non-uniform value
+  if (non_uniform_global == 0) {
+  ^^
+
+test:11:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
+  if (non_uniform_global == 0) {
+      ^^^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamModifiedInNonUniformControlFlow_NestedCall) {
+    std::string src = R"(
+@binding(0) @group(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo2(p : ptr<function, i32>) {
+  *p = 42;
+}
+
+fn foo(p : ptr<function, i32>) {
+  foo2(p);
+}
+
+@compute @workgroup_size(64)
+fn main() {
+  var a : i32;
+  if (non_uniform_global == 0) {
+    foo(&a);
+  }
+
+  if (a == 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, false);
+    EXPECT_EQ(error_,
+              R"(test:20:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+    workgroupBarrier();
+    ^^^^^^^^^^^^^^^^
+
+test:15:3 note: control flow depends on non-uniform value
+  if (non_uniform_global == 0) {
+  ^^
+
+test:15:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
+  if (non_uniform_global == 0) {
+      ^^^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamModifiedInUniformControlFlow) {
+    std::string src = R"(
+@binding(0) @group(0) var<uniform> uniform_global : i32;
+
+fn foo(p : ptr<function, i32>) {
+  *p = 42;
+}
+
+@compute @workgroup_size(64)
+fn main() {
+  var a : i32;
+  if (uniform_global == 0) {
+    foo(&a);
+  }
+
+  if (a == 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, true);
+}
+
 TEST_F(UniformityAnalysisTest, NonUniformPointerParameterBecomesUniform_AfterUse) {
     std::string src = R"(
 @group(0) @binding(0) var<storage, read_write> non_uniform : i32;