diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index 1960c60..ada52f9 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -27,18 +27,29 @@
 
 #include "src/tint/lang/core/ir/binary/decode.h"
 
+#include <cmath>
+#include <cstdint>
+#include <string>
 #include <utility>
 
 #include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/control_instruction.h"
 #include "src/tint/lang/core/ir/module.h"
 #include "src/tint/lang/core/type/depth_multisampled_texture.h"
 #include "src/tint/lang/core/type/depth_texture.h"
 #include "src/tint/lang/core/type/external_texture.h"
+#include "src/tint/lang/core/type/invalid.h"
 #include "src/tint/lang/core/type/multisampled_texture.h"
 #include "src/tint/lang/core/type/sampled_texture.h"
 #include "src/tint/lang/core/type/storage_texture.h"
+#include "src/tint/lang/core/type/vector.h"
+#include "src/tint/utils/containers/hashset.h"
 #include "src/tint/utils/containers/transform.h"
+#include "src/tint/utils/diagnostic/diagnostic.h"
 #include "src/tint/utils/macros/compiler.h"
+#include "src/tint/utils/result/result.h"
+#include "src/tint/utils/text/string.h"
+#include "src/tint/utils/text/text_style.h"
 
 TINT_BEGIN_DISABLE_PROTOBUF_WARNINGS();
 #include "src/tint/lang/core/ir/binary/ir.pb.h"
@@ -66,6 +77,9 @@
     Vector<ir::BreakIf*, 32> break_ifs_{};
     Vector<ir::Continue*, 32> continues_{};
 
+    diag::List diags_{};
+    Hashset<std::string, 4> struct_names_{};
+
     Result<Module> Decode() {
         {
             const size_t n = static_cast<size_t>(mod_in_.types().size());
@@ -85,7 +99,7 @@
             const size_t n = static_cast<size_t>(mod_in_.blocks().size());
             blocks_.Reserve(n);
             for (size_t i = 0; i < n; i++) {
-                auto id = static_cast<uint32_t>(i + 1);
+                auto id = static_cast<uint32_t>(i);
                 if (id == mod_in_.root_block()) {
                     blocks_.Push(mod_out_.root_block);
                 } else {
@@ -115,28 +129,93 @@
             PopulateBlock(blocks_[i], mod_in_.blocks()[static_cast<int>(i)]);
         }
 
-        for (auto* exit : exit_ifs_) {
-            InferControlInstruction(exit, &ExitIf::SetIf);
-        }
-        for (auto* exit : exit_switches_) {
-            InferControlInstruction(exit, &ExitSwitch::SetSwitch);
-        }
-        for (auto* exit : exit_loops_) {
-            InferControlInstruction(exit, &ExitLoop::SetLoop);
-        }
-        for (auto* break_ifs : break_ifs_) {
-            InferControlInstruction(break_ifs, &BreakIf::SetLoop);
-        }
-        for (auto* next_iters : next_iterations_) {
-            InferControlInstruction(next_iters, &NextIteration::SetLoop);
-        }
-        for (auto* cont : continues_) {
-            InferControlInstruction(cont, &Continue::SetLoop);
+        if (diags_.ContainsErrors()) {
+            // Note: Its not safe to call InferControlInstruction() with a broken IR.
+            return Failure{std::move(diags_)};
         }
 
+        if (CheckBlocks()) {
+            for (auto* exit : exit_ifs_) {
+                InferControlInstruction(exit, &ExitIf::SetIf);
+            }
+            for (auto* exit : exit_switches_) {
+                InferControlInstruction(exit, &ExitSwitch::SetSwitch);
+            }
+            for (auto* exit : exit_loops_) {
+                InferControlInstruction(exit, &ExitLoop::SetLoop);
+            }
+            for (auto* break_ifs : break_ifs_) {
+                InferControlInstruction(break_ifs, &BreakIf::SetLoop);
+            }
+            for (auto* next_iters : next_iterations_) {
+                InferControlInstruction(next_iters, &NextIteration::SetLoop);
+            }
+            for (auto* cont : continues_) {
+                InferControlInstruction(cont, &Continue::SetLoop);
+            }
+        }
+
+        if (diags_.ContainsErrors()) {
+            return Failure{std::move(diags_)};
+        }
         return std::move(mod_out_);
     }
 
+    /// Adds a new error to the diagnostics and returns a reference to it
+    diag::Diagnostic& Error() { return diags_.AddError(Source{}); }
+
+    /// Errors if @p number is not finite.
+    /// @returns @p number if finite, otherwise 0.
+    template <typename T>
+    Number<T> CheckFinite(Number<T> number) {
+        if (TINT_UNLIKELY(!std::isfinite(number.value))) {
+            Error() << "value must be finite";
+            return Number<T>{};
+        }
+        return number;
+    }
+
+    /// @returns true if all blocks are reachable, acyclic nesting depth is less than or equal to
+    /// kMaxBlockDepth.
+    bool CheckBlocks() {
+        const size_t kMaxBlockDepth = 128;
+        Vector<std::pair<const ir::Block*, size_t>, 32> pending;
+        pending.Push(std::make_pair(mod_out_.root_block, 0));
+        for (auto& fn : mod_out_.functions) {
+            pending.Push(std::make_pair(fn->Block(), 0));
+        }
+        Hashset<const ir::Block*, 32> seen;
+        while (!pending.IsEmpty()) {
+            const auto block_depth = pending.Pop();
+            const auto* block = block_depth.first;
+            const size_t depth = block_depth.second;
+            if (!seen.Add(block)) {
+                Error() << "cyclic nesting of blocks";
+                return false;
+            }
+            if (depth > kMaxBlockDepth) {
+                Error() << "block nesting exceeds " << kMaxBlockDepth;
+                return false;
+            }
+            for (auto* inst = block->Instructions(); inst; inst = inst->next) {
+                if (auto* ctrl = inst->As<ir::ControlInstruction>()) {
+                    ctrl->ForeachBlock([&](const ir::Block* child) {
+                        pending.Push(std::make_pair(child, depth + 1));
+                    });
+                }
+            }
+        }
+
+        for (auto* block : blocks_) {
+            if (!seen.Contains(block)) {
+                Error() << "unreachable block";
+                return false;
+            }
+        }
+
+        return true;
+    }
+
     template <typename EXIT, typename CTRL_INST>
     void InferControlInstruction(EXIT* exit, void (EXIT::*set)(CTRL_INST*)) {
         for (auto* block = exit->Block(); block;) {
@@ -174,7 +253,10 @@
 
         Vector<FunctionParam*, 8> params_out;
         for (auto param_in : fn_in.parameters()) {
-            params_out.Push(ValueAs<ir::FunctionParam>(param_in));
+            auto* param_out = ValueAs<FunctionParam>(param_in);
+            if (TINT_LIKELY(param_out)) {
+                params_out.Push(param_out);
+            }
         }
         if (fn_in.has_return_location()) {
             fn_out->SetReturnLocation(Location(fn_in.return_location()));
@@ -189,7 +271,13 @@
         fn_out->SetBlock(Block(fn_in.block()));
     }
 
-    ir::Function* Function(uint32_t id) { return id > 0 ? mod_out_.functions[id - 1] : nullptr; }
+    ir::Function* Function(uint32_t id) {
+        if (TINT_UNLIKELY(id >= mod_out_.functions.Length())) {
+            Error() << "function id " << id << " out of range";
+            return nullptr;
+        }
+        return mod_out_.functions[id];
+    }
 
     Function::PipelineStage PipelineStage(pb::PipelineStage stage) {
         switch (stage) {
@@ -215,28 +303,37 @@
     }
 
     void PopulateBlock(ir::Block* block_out, const pb::Block& block_in) {
-        if (block_in.is_multi_in()) {
+        if (auto* mib = block_out->As<ir::MultiInBlock>()) {
             Vector<ir::BlockParam*, 8> params;
-            for (auto param : block_in.parameters()) {
-                params.Push(ValueAs<BlockParam>(param));
+            for (auto param_in : block_in.parameters()) {
+                auto* param_out = ValueAs<BlockParam>(param_in);
+                if (TINT_LIKELY(param_out)) {
+                    params.Push(param_out);
+                }
             }
-            block_out->As<ir::MultiInBlock>()->SetParams(std::move(params));
+            mib->SetParams(std::move(params));
         }
         for (auto& inst : block_in.instructions()) {
             block_out->Append(Instruction(inst));
         }
     }
 
-    ir::Block* Block(uint32_t id) { return id > 0 ? blocks_[id - 1] : nullptr; }
+    ir::Block* Block(uint32_t id) {
+        if (TINT_UNLIKELY(id >= blocks_.Length())) {
+            Error() << "block id " << id << " out of range";
+            return b.Block();
+        }
+        return blocks_[id];
+    }
 
     template <typename T>
     T* BlockAs(uint32_t id) {
         auto* block = Block(id);
-        if (auto cast = block->As<T>(); TINT_LIKELY(cast)) {
+        if (auto cast = As<T>(block); TINT_LIKELY(cast)) {
             return cast;
         }
-        TINT_ICE() << "block " << id << " is " << (block ? block->TypeInfo().name : "<null>")
-                   << " expected " << TypeInfo::Of<T>().name;
+        Error() << "block " << id << " is " << (block ? block->TypeInfo().name : "<null>")
+                << " expected " << TypeInfo::Of<T>().name;
         return nullptr;
     }
 
@@ -330,6 +427,11 @@
             case pb::Instruction::KindCase::KIND_NOT_SET:
                 break;
         }
+        if (!inst_out) {
+            Error() << "invalid Instruction.kind: " << std::to_string(inst_in.kind_case());
+            return b.Let(mod_out_.Types().invalid());
+        }
+
         TINT_ASSERT(inst_out);
 
         Vector<ir::Value*, 4> operands;
@@ -345,8 +447,15 @@
         inst_out->SetResults(std::move(results));
 
         if (inst_in.has_break_if()) {
-            static_cast<BreakIf*>(inst_out)->SetNumNextIterValues(
-                inst_in.break_if().num_next_iter_values());
+            auto num_next_iter_values = inst_in.break_if().num_next_iter_values();
+            bool is_valid =
+                inst_out->Operands().Length() >= num_next_iter_values + BreakIf::kArgsOperandOffset;
+            if (TINT_LIKELY(is_valid)) {
+                static_cast<BreakIf*>(inst_out)->SetNumNextIterValues(
+                    inst_in.break_if().num_next_iter_values());
+            } else {
+                Error() << "invalid value for num_next_iter_values()";
+            }
         }
 
         return inst_out;
@@ -416,12 +525,8 @@
 
     ir::If* CreateInstructionIf(const pb::InstructionIf& if_in) {
         auto* if_out = mod_out_.allocators.instructions.Create<ir::If>();
-        if (if_in.has_true_()) {
-            if_out->SetTrue(Block(if_in.true_()));
-        }
-        if (if_in.has_false_()) {
-            if_out->SetFalse(Block(if_in.false_()));
-        }
+        if_out->SetTrue(if_in.has_true_() ? Block(if_in.true_()) : b.Block());
+        if_out->SetFalse(if_in.has_false_() ? Block(if_in.false_()) : b.Block());
         return if_out;
     }
 
@@ -440,16 +545,16 @@
 
     ir::Loop* CreateInstructionLoop(const pb::InstructionLoop& loop_in) {
         auto* loop_out = mod_out_.allocators.instructions.Create<ir::Loop>();
-        if (loop_in.has_initalizer()) {
-            loop_out->SetInitializer(Block(loop_in.initalizer()));
+        if (loop_in.has_initializer()) {
+            loop_out->SetInitializer(Block(loop_in.initializer()));
         } else {
-            loop_out->SetInitializer(mod_out_.blocks.Create());
+            loop_out->SetInitializer(b.Block());
         }
         loop_out->SetBody(BlockAs<ir::MultiInBlock>(loop_in.body()));
         if (loop_in.has_continuing()) {
             loop_out->SetContinuing(BlockAs<ir::MultiInBlock>(loop_in.continuing()));
         } else {
-            loop_out->SetContinuing(mod_out_.blocks.Create<ir::MultiInBlock>());
+            loop_out->SetContinuing(b.MultiInBlock());
         }
         return loop_out;
     }
@@ -491,7 +596,7 @@
             case_out.block->SetParent(switch_out);
             for (auto selector_in : case_in.selectors()) {
                 ir::Switch::CaseSelector selector_out{};
-                selector_out.val = b.Constant(ConstantValue(selector_in));
+                selector_out.val = Constant(selector_in);
                 case_out.selectors.Push(std::move(selector_out));
             }
             if (case_in.is_default()) {
@@ -562,7 +667,9 @@
             case pb::Type::KindCase::KIND_NOT_SET:
                 break;
         }
-        TINT_ICE() << type_in.kind_case();
+
+        Error() << "invalid Type.kind: " << std::to_string(type_in.kind_case());
+        return mod_out_.Types().invalid();
     }
 
     const type::Type* CreateTypeBasic(pb::TypeBasic basic_in) {
@@ -584,15 +691,28 @@
             case pb::TypeBasic::TypeBasic_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
-        TINT_ICE() << "invalid TypeBasic: " << basic_in;
+
+        Error() << "invalid TypeBasic: " << std::to_string(basic_in);
+        return mod_out_.Types().invalid();
     }
 
-    const type::Vector* CreateTypeVector(const pb::TypeVector& vector_in) {
+    const type::Type* CreateTypeVector(const pb::TypeVector& vector_in) {
+        const auto width = vector_in.width();
+        if (TINT_UNLIKELY(width < 2 || width > 4)) {
+            Error() << "invalid vector width";
+            return mod_out_.Types().invalid();
+        }
         auto* el_ty = Type(vector_in.element_type());
         return mod_out_.Types().vec(el_ty, vector_in.width());
     }
 
-    const type::Matrix* CreateTypeMatrix(const pb::TypeMatrix& matrix_in) {
+    const type::Type* CreateTypeMatrix(const pb::TypeMatrix& matrix_in) {
+        const auto rows = matrix_in.num_rows();
+        const auto cols = matrix_in.num_columns();
+        if (TINT_UNLIKELY(rows < 2 || rows > 4 || cols < 2 || cols > 4)) {
+            Error() << "invalid matrix dimensions";
+            return mod_out_.Types().invalid();
+        }
         auto* el_ty = Type(matrix_in.element_type());
         auto* column_ty = mod_out_.Types().vec(el_ty, matrix_in.num_rows());
         return mod_out_.Types().mat(column_ty, matrix_in.num_columns());
@@ -605,15 +725,38 @@
         return mod_out_.Types().ptr(address_space, store_ty, access);
     }
 
-    const type::Struct* CreateTypeStruct(const pb::TypeStruct& struct_in) {
+    const type::Type* CreateTypeStruct(const pb::TypeStruct& struct_in) {
+        auto struct_name = struct_in.name();
+        if (TINT_UNLIKELY(struct_name.empty())) {
+            Error() << "struct must have a name";
+            return mod_out_.Types().invalid();
+        }
+        if (!struct_names_.Add(struct_name)) {
+            Error() << "duplicate struct name: " << style::Type(struct_name);
+            return mod_out_.Types().invalid();
+        }
+
         Vector<const core::type::StructMember*, 8> members_out;
         uint32_t offset = 0;
         for (auto& member_in : struct_in.member()) {
-            auto symbol = mod_out_.symbols.Register(member_in.name());
+            auto member_name = member_in.name();
+            if (TINT_UNLIKELY(member_name.empty())) {
+                Error() << "struct member must have a name";
+                return mod_out_.Types().invalid();
+            }
+            auto symbol = mod_out_.symbols.Register(member_name);
             auto* type = Type(member_in.type());
             auto index = static_cast<uint32_t>(members_out.Length());
             auto align = member_in.align();
             auto size = member_in.size();
+            if (TINT_UNLIKELY(align == 0)) {
+                Error() << "struct member must have non-zero alignment";
+                align = 1;
+            }
+            if (TINT_UNLIKELY(size == 0)) {
+                Error() << "struct member must have non-zero size";
+                size = 1;
+            }
             core::type::StructMemberAttributes attributes_out{};
             if (member_in.has_attributes()) {
                 auto& attributes_in = member_in.attributes();
@@ -641,7 +784,11 @@
             offset += size;
             members_out.Push(member_out);
         }
-        auto name = mod_out_.symbols.Register(struct_in.name());
+        if (TINT_UNLIKELY(members_out.IsEmpty())) {
+            Error() << "struct requires at least one member";
+            return mod_out_.Types().invalid();
+        }
+        auto name = mod_out_.symbols.Register(struct_name);
         return mod_out_.Types().Struct(name, std::move(members_out));
     }
 
@@ -649,16 +796,29 @@
         return mod_out_.Types().atomic(Type(atomic_in.type()));
     }
 
-    const type::Array* CreateTypeArray(const pb::TypeArray& array_in) {
+    const type::Type* CreateTypeArray(const pb::TypeArray& array_in) {
         auto* element = Type(array_in.element());
-        uint32_t stride = static_cast<uint32_t>(array_in.stride());
-        uint32_t count = static_cast<uint32_t>(array_in.count());
+        uint32_t stride = array_in.stride();
+        uint32_t count = array_in.count();
+        if (element->Align() == 0 || element->Size() == 0) {
+            Error() << "cannot create an array of an unsized type";
+            return mod_out_.Types().invalid();
+        }
+        uint32_t implicit_stride = tint::RoundUp(element->Align(), element->Size());
+        if (stride < implicit_stride) {
+            Error() << "array element stride is smaller than the implicit stride";
+            return mod_out_.Types().invalid();
+        }
         return count > 0 ? mod_out_.Types().array(element, count, stride)
                          : mod_out_.Types().runtime_array(element, stride);
     }
 
-    const type::DepthTexture* CreateTypeDepthTexture(const pb::TypeDepthTexture& texture_in) {
+    const type::Type* CreateTypeDepthTexture(const pb::TypeDepthTexture& texture_in) {
         auto dimension = TextureDimension(texture_in.dimension());
+        if (!type::DepthTexture::IsValidDimension(dimension)) {
+            Error() << "invalid DepthTexture dimension";
+            return mod_out_.Types().invalid();
+        }
         return mod_out_.Types().Get<type::DepthTexture>(dimension);
     }
 
@@ -675,9 +835,13 @@
         return mod_out_.Types().Get<type::MultisampledTexture>(dimension, sub_type);
     }
 
-    const type::DepthMultisampledTexture* CreateTypeDepthMultisampledTexture(
+    const type::Type* CreateTypeDepthMultisampledTexture(
         const pb::TypeDepthMultisampledTexture& texture_in) {
         auto dimension = TextureDimension(texture_in.dimension());
+        if (!type::DepthMultisampledTexture::IsValidDimension(dimension)) {
+            Error() << "invalid DepthMultisampledTexture dimension";
+            return mod_out_.Types().invalid();
+        }
         return mod_out_.Types().Get<type::DepthMultisampledTexture>(dimension);
     }
 
@@ -699,7 +863,13 @@
         return mod_out_.Types().Get<type::Sampler>(kind);
     }
 
-    const type::Type* Type(size_t id) { return id > 0 ? types_[id - 1] : nullptr; }
+    const type::Type* Type(size_t id) {
+        if (TINT_UNLIKELY(id >= types_.Length())) {
+            Error() << "type id " << id << " out of range";
+            return mod_out_.Types().invalid();
+        }
+        return types_[id];
+    }
 
     ////////////////////////////////////////////////////////////////////////////
     // Values
@@ -720,14 +890,15 @@
                 value_out = BlockParameter(value_in.block_parameter());
                 break;
             case pb::Value::KindCase::kConstant:
-                value_out = b.Constant(ConstantValue(value_in.constant()));
+                value_out = Constant(value_in.constant());
                 break;
             case pb::Value::KindCase::KIND_NOT_SET:
                 break;
         }
 
         if (!value_out) {
-            TINT_ICE() << "invalid TypeDecl.kind: " << value_in.kind_case();
+            Error() << "invalid value kind: " << std::to_string(value_in.kind_case());
+            return b.InvalidConstant();
         }
 
         return value_out;
@@ -736,7 +907,7 @@
     ir::InstructionResult* InstructionResult(const pb::InstructionResult& res_in) {
         auto* type = Type(res_in.type());
         auto* res_out = b.InstructionResult(type);
-        if (res_in.has_name()) {
+        if (!res_in.name().empty()) {
             mod_out_.SetName(res_out, res_in.name());
         }
         return res_out;
@@ -745,7 +916,7 @@
     ir::FunctionParam* FunctionParameter(const pb::FunctionParameter& param_in) {
         auto* type = Type(param_in.type());
         auto* param_out = b.FunctionParam(type);
-        if (param_in.has_name()) {
+        if (!param_in.name().empty()) {
             mod_out_.SetName(param_out, param_in.name());
         }
 
@@ -772,22 +943,30 @@
     ir::BlockParam* BlockParameter(const pb::BlockParameter& param_in) {
         auto* type = Type(param_in.type());
         auto* param_out = b.BlockParam(type);
-        if (param_in.has_name()) {
+        if (!param_in.name().empty()) {
             mod_out_.SetName(param_out, param_in.name());
         }
         return param_out;
     }
 
-    ir::Value* Value(uint32_t id) { return id > 0 ? values_[id - 1] : nullptr; }
+    ir::Constant* Constant(uint32_t value_id) { return b.Constant(ConstantValue(value_id)); }
+
+    ir::Value* Value(uint32_t id) {
+        if (TINT_UNLIKELY(id > values_.Length())) {
+            Error() << "value id " << id << " out of range";
+            return nullptr;
+        }
+        return id > 0 ? values_[id - 1] : nullptr;
+    }
 
     template <typename T>
     T* ValueAs(uint32_t id) {
         auto* value = Value(id);
-        if (auto cast = value->As<T>(); TINT_LIKELY(cast)) {
+        if (auto cast = As<T>(value); TINT_LIKELY(cast)) {
             return cast;
         }
-        TINT_ICE() << "Value " << id << " is " << (value ? value->TypeInfo().name : "<null>")
-                   << " expected " << TypeInfo::Of<T>().name;
+        Error() << "value " << id << " is " << (value ? value->TypeInfo().name : "<null>")
+                << " expected " << TypeInfo::Of<T>().name;
         return nullptr;
     }
 
@@ -805,7 +984,8 @@
             case pb::ConstantValue::KindCase::KIND_NOT_SET:
                 break;
         }
-        TINT_ICE() << "invalid ConstantValue.kind: " << value_in.kind_case();
+        Error() << "invalid ConstantValue.kind: " << std::to_string(value_in.kind_case());
+        return b.InvalidConstant()->Value();
     }
 
     const core::constant::Value* CreateConstantScalar(const pb::ConstantValueScalar& value_in) {
@@ -817,33 +997,69 @@
             case pb::ConstantValueScalar::KindCase::kU32:
                 return b.ConstantValue(u32(value_in.u32()));
             case pb::ConstantValueScalar::KindCase::kF32:
-                return b.ConstantValue(f32(value_in.f32()));
+                return b.ConstantValue(CheckFinite(f32(value_in.f32())));
             case pb::ConstantValueScalar::KindCase::kF16:
-                return b.ConstantValue(f16(value_in.f16()));
+                return b.ConstantValue(CheckFinite(f16(value_in.f16())));
             case pb::ConstantValueScalar::KindCase::KIND_NOT_SET:
                 break;
         }
-        TINT_ICE() << "invalid ConstantValueScalar.kind: " << value_in.kind_case();
+        Error() << "invalid ConstantValueScalar.kind: " << std::to_string(value_in.kind_case());
+        return b.InvalidConstant()->Value();
     }
 
     const core::constant::Value* CreateConstantComposite(
         const pb::ConstantValueComposite& composite_in) {
         auto* type = Type(composite_in.type());
+        auto type_elements = type->Elements();
+        size_t num_values = static_cast<size_t>(composite_in.elements().size());
+        if (TINT_UNLIKELY(type_elements.count == 0)) {
+            Error() << "cannot create a composite of type " << type->FriendlyName();
+            return b.InvalidConstant()->Value();
+        }
+        if (TINT_UNLIKELY(type_elements.count != num_values)) {
+            Error() << "constant composite type " << type->FriendlyName() << " expects "
+                    << type_elements.count << " elements, but " << num_values << " values encoded";
+            return b.InvalidConstant()->Value();
+        }
         Vector<const core::constant::Value*, 8> elements_out;
         for (auto element_id : composite_in.elements()) {
-            elements_out.Push(ConstantValue(element_id));
+            uint32_t i = static_cast<uint32_t>(elements_out.Length());
+            auto* value = ConstantValue(element_id);
+            if (auto* el_type = type->Element(i); TINT_UNLIKELY(value->Type() != el_type)) {
+                Error() << "constant composite element value type " << value->Type()->FriendlyName()
+                        << " does not match element type " << el_type->FriendlyName();
+                return b.InvalidConstant()->Value();
+            }
+            elements_out.Push(value);
         }
         return mod_out_.constant_values.Composite(type, std::move(elements_out));
     }
 
     const core::constant::Value* CreateConstantSplat(const pb::ConstantValueSplat& splat_in) {
         auto* type = Type(splat_in.type());
-        auto* elem = ConstantValue(splat_in.elements());
-        return mod_out_.constant_values.Splat(type, elem);
+        uint32_t num_elements = type->Elements().count;
+        if (TINT_UNLIKELY(num_elements == 0)) {
+            Error() << "cannot create a splat of type " << type->FriendlyName();
+            return b.InvalidConstant()->Value();
+        }
+        auto* value = ConstantValue(splat_in.elements());
+        for (uint32_t i = 0; i < num_elements; i++) {
+            auto* el_type = type->Element(i);
+            if (TINT_UNLIKELY(el_type != value->Type())) {
+                Error() << "constant splat element value type " << value->Type()->FriendlyName()
+                        << " does not match element " << i << " type " << el_type->FriendlyName();
+                return b.InvalidConstant()->Value();
+            }
+        }
+        return mod_out_.constant_values.Splat(type, value);
     }
 
     const core::constant::Value* ConstantValue(uint32_t id) {
-        return id > 0 ? constant_values_[id - 1] : nullptr;
+        if (TINT_UNLIKELY(id >= constant_values_.Length())) {
+            Error() << "constant value id " << id << " out of range";
+            return b.InvalidConstant()->Value();
+        }
+        return constant_values_[id];
     }
 
     ////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/core/ir/binary/encode.cc b/src/tint/lang/core/ir/binary/encode.cc
index d908788..05f06bd 100644
--- a/src/tint/lang/core/ir/binary/encode.cc
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -110,7 +110,7 @@
         }
         Vector<pb::Function*, 8> fns_out;
         for (auto& fn_in : mod_in_.functions) {
-            uint32_t id = static_cast<uint32_t>(fns_out.Length() + 1);
+            uint32_t id = static_cast<uint32_t>(mod_out_.functions().size());
             fns_out.Push(mod_out_.add_functions());
             functions_.Add(fn_in, id);
         }
@@ -153,7 +153,7 @@
         fn_out->set_block(Block(fn_in->Block()));
     }
 
-    uint32_t Function(const ir::Function* fn_in) { return fn_in ? *functions_.Get(fn_in) : 0; }
+    uint32_t Function(const ir::Function* fn_in) { return *functions_.Get(fn_in); }
 
     pb::PipelineStage PipelineStage(Function::PipelineStage stage) {
         switch (stage) {
@@ -173,12 +173,11 @@
     // Blocks
     ////////////////////////////////////////////////////////////////////////////
     uint32_t Block(const ir::Block* block_in) {
-        if (block_in == nullptr) {
-            return 0;
-        }
+        TINT_ASSERT(block_in != nullptr);
+
         return blocks_.GetOrAdd(block_in, [&]() -> uint32_t {
+            auto id = static_cast<uint32_t>(mod_out_.blocks().size());
             auto& block_out = *mod_out_.add_blocks();
-            auto id = static_cast<uint32_t>(blocks_.Count());
             for (auto* inst : *block_in) {
                 Instruction(*block_out.add_instructions(), inst);
             }
@@ -296,7 +295,7 @@
 
     void InstructionLoop(pb::InstructionLoop& loop_out, const ir::Loop* loop_in) {
         if (loop_in->HasInitializer()) {
-            loop_out.set_initalizer(Block(loop_in->Initializer()));
+            loop_out.set_initializer(Block(loop_in->Initializer()));
         }
         loop_out.set_body(Block(loop_in->Body()));
         if (loop_in->HasContinuing()) {
@@ -352,9 +351,7 @@
     // Types
     ////////////////////////////////////////////////////////////////////////////
     uint32_t Type(const core::type::Type* type_in) {
-        if (type_in == nullptr) {
-            return 0;
-        }
+        TINT_ASSERT(type_in != nullptr);
         return types_.GetOrAdd(type_in, [&]() -> uint32_t {
             pb::Type type_out;
             tint::Switch(
@@ -393,7 +390,7 @@
                 TINT_ICE_ON_NO_MATCH);
 
             mod_out_.mutable_types()->Add(std::move(type_out));
-            return static_cast<uint32_t>(mod_out_.types().size());
+            return static_cast<uint32_t>(mod_out_.types().size() - 1);
         });
     }
 
@@ -564,9 +561,7 @@
     // ConstantValues
     ////////////////////////////////////////////////////////////////////////////
     uint32_t ConstantValue(const core::constant::Value* constant_in) {
-        if (!constant_in) {
-            return 0;
-        }
+        TINT_ASSERT(constant_in != nullptr);
         return constant_values_.GetOrAdd(constant_in, [&] {
             pb::ConstantValue constant_out;
             tint::Switch(
@@ -595,7 +590,7 @@
                 TINT_ICE_ON_NO_MATCH);
 
             mod_out_.mutable_constant_values()->Add(std::move(constant_out));
-            return static_cast<uint32_t>(mod_out_.constant_values().size());
+            return static_cast<uint32_t>(mod_out_.constant_values().size() - 1);
         });
     }
 
diff --git a/src/tint/lang/core/ir/binary/ir.proto b/src/tint/lang/core/ir/binary/ir.proto
index 31c2821..a5271c3 100644
--- a/src/tint/lang/core/ir/binary/ir.proto
+++ b/src/tint/lang/core/ir/binary/ir.proto
@@ -328,9 +328,9 @@
 }
 
 message InstructionLoop {
-    optional uint32 initalizer = 1;  // Module.blocks
-    optional uint32 body = 2;        // Module.blocks
-    optional uint32 continuing = 3;  // Module.blocks
+    optional uint32 initializer = 1;  // Module.blocks
+    optional uint32 body = 2;         // Module.blocks
+    optional uint32 continuing = 3;   // Module.blocks
 }
 
 message InstructionExitIf {}
diff --git a/src/tint/lang/core/ir/disassembly.cc b/src/tint/lang/core/ir/disassembly.cc
index 1a8a9a6..6909d45 100644
--- a/src/tint/lang/core/ir/disassembly.cc
+++ b/src/tint/lang/core/ir/disassembly.cc
@@ -320,7 +320,7 @@
             out_ << ", ";
         }
         SourceMarker sm(this);
-        out_ << NameOf(p) << ":" << StyleType(p->Type()->FriendlyName());
+        out_ << NameOf(p) << ":" << StyleType(p->Type() ? p->Type()->FriendlyName() : "undef");
         sm.Store(p);
 
         EmitParamAttributes(p);
diff --git a/src/tint/lang/core/ir/function_param.cc b/src/tint/lang/core/ir/function_param.cc
index 69b1dc0..00d72e5 100644
--- a/src/tint/lang/core/ir/function_param.cc
+++ b/src/tint/lang/core/ir/function_param.cc
@@ -35,9 +35,7 @@
 
 namespace tint::core::ir {
 
-FunctionParam::FunctionParam(const core::type::Type* ty) : type_(ty) {
-    TINT_ASSERT(ty != nullptr);
-}
+FunctionParam::FunctionParam(const core::type::Type* ty) : type_(ty) {}
 
 FunctionParam::~FunctionParam() = default;
 
diff --git a/src/tint/lang/core/ir/function_param.h b/src/tint/lang/core/ir/function_param.h
index c9a0cea..6ff2235 100644
--- a/src/tint/lang/core/ir/function_param.h
+++ b/src/tint/lang/core/ir/function_param.h
@@ -73,6 +73,10 @@
     /// @returns the type of the var
     const core::type::Type* Type() const override { return type_; }
 
+    /// Sets the type of the parameter to @p type
+    /// @param type the new type of the parameter
+    void SetType(const core::type::Type* type) { type_ = type; }
+
     /// @copydoc Value::Clone()
     FunctionParam* Clone(CloneContext& ctx) override;
 
diff --git a/src/tint/lang/core/ir/function_param_test.cc b/src/tint/lang/core/ir/function_param_test.cc
index 80b150f..30db305 100644
--- a/src/tint/lang/core/ir/function_param_test.cc
+++ b/src/tint/lang/core/ir/function_param_test.cc
@@ -36,16 +36,6 @@
 using namespace tint::core::number_suffixes;  // NOLINT
 using IR_FunctionParamTest = IRTestHelper;
 
-TEST_F(IR_FunctionParamTest, Fail_NullType) {
-    EXPECT_DEATH_IF_SUPPORTED(
-        {
-            Module mod;
-            Builder b{mod};
-            b.FunctionParam(nullptr);
-        },
-        "");
-}
-
 TEST_F(IR_FunctionParamTest, Fail_SetDuplicateBuiltin) {
     EXPECT_DEATH_IF_SUPPORTED(
         {
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 4a9a866..2912315 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -769,6 +769,11 @@
             return;
         }
 
+        if (!param->Type()) {
+            AddError(param) << "function parameter has nullptr type";
+            return;
+        }
+
         // References not allowed on function signatures even with Capability::kAllowRefTypes
         if (HoldsType<type::Reference>(param->Type())) {
             AddError(param) << "references are not permitted as parameter types";
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 820182a..9d5b0f1 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -239,6 +239,28 @@
 )");
 }
 
+TEST_F(IR_ValidatorTest, Function_ParameterWithNullType) {
+    auto* f = b.Function("my_func", ty.void_());
+    auto* p = b.FunctionParam("my_param", nullptr);
+    f->SetParams({p});
+    f->Block()->Append(b.Return(f));
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:1:17 error: function parameter has nullptr type
+%my_func = func(%my_param:undef):void {
+                ^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func(%my_param:undef):void {
+  $B1: {
+    ret
+  }
+}
+)");
+}
+
 TEST_F(IR_ValidatorTest, CallToFunctionOutsideModule) {
     auto* f = b.Function("f", ty.void_());
     auto* g = b.Function("g", ty.void_());
diff --git a/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access.cc b/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access.cc
index 0b2500a..bacc9a0 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access.cc
@@ -321,10 +321,10 @@
 
 /// BufferAccess describes a single storage or uniform buffer access
 struct BufferAccess {
-    sem::GlobalVariable const* var = nullptr;  // Storage or uniform buffer variable
-    Offset const* offset = nullptr;            // The byte offset on var
-    core::type::Type const* type = nullptr;    // The type of the access
-    operator bool() const { return var; }      // Returns true if valid
+    sem::GlobalVariable const* var = nullptr;       // Storage or uniform buffer variable
+    Offset const* offset = nullptr;                 // The byte offset on var
+    core::type::Type const* type = nullptr;         // The type of the access
+    explicit operator bool() const { return var; }  // Returns true if valid
 };
 
 /// Store describes a single storage or uniform buffer write
diff --git a/src/tint/lang/msl/writer/printer/function_test.cc b/src/tint/lang/msl/writer/printer/function_test.cc
index 0513a48..8ba60c2 100644
--- a/src/tint/lang/msl/writer/printer/function_test.cc
+++ b/src/tint/lang/msl/writer/printer/function_test.cc
@@ -25,6 +25,7 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
+#include "src/tint/lang/core/type/sampled_texture.h"
 #include "src/tint/lang/msl/writer/printer/helper_test.h"
 
 namespace tint::msl::writer {
@@ -41,7 +42,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, EntryPointParameterBindingPoint) {
+TEST_F(MslPrinterTest, EntryPointParameterBufferBindingPoint) {
     auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
     auto* storage = b.FunctionParam("storage", ty.ptr(core::AddressSpace::kStorage, ty.i32()));
     auto* uniform = b.FunctionParam("uniform", ty.ptr(core::AddressSpace::kUniform, ty.i32()));
@@ -57,5 +58,22 @@
 )");
 }
 
+TEST_F(MslPrinterTest, EntryPointParameterHandleBindingPoint) {
+    auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32());
+    auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+    auto* texture = b.FunctionParam("texture", t);
+    auto* sampler = b.FunctionParam("sampler", ty.sampler());
+    texture->SetBindingPoint(0, 1);
+    sampler->SetBindingPoint(0, 2);
+    func->SetParams({texture, sampler});
+    func->Block()->Append(b.Return(func));
+
+    ASSERT_TRUE(Generate()) << err_ << output_;
+    EXPECT_EQ(output_, MetalHeader() + R"(
+fragment void foo(texture2d<float, access::sample> texture [[texture(1)]], sampler sampler [[sampler(2)]]) {
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 7c58cd4..0ae4046 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -329,16 +329,28 @@
                 }
 
                 if (auto binding_point = param->BindingPoint()) {
-                    auto ptr = param->Type()->As<core::type::Pointer>();
                     TINT_ASSERT(binding_point->group == 0);
-                    switch (ptr->AddressSpace()) {
-                        case core::AddressSpace::kStorage:
-                        case core::AddressSpace::kUniform:
-                            out << " [[buffer(" << binding_point->binding << ")]]";
-                            break;
-                        default:
-                            TINT_UNREACHABLE() << "invalid address space with binding point: "
-                                               << ptr->AddressSpace();
+                    if (auto ptr = param->Type()->As<core::type::Pointer>()) {
+                        switch (ptr->AddressSpace()) {
+                            case core::AddressSpace::kStorage:
+                            case core::AddressSpace::kUniform:
+                                out << " [[buffer(" << binding_point->binding << ")]]";
+                                break;
+                            default:
+                                TINT_UNREACHABLE() << "invalid address space with binding point: "
+                                                   << ptr->AddressSpace();
+                        }
+                    } else {
+                        // Handle types are declared by value instead of by pointer.
+                        Switch(
+                            param->Type(),
+                            [&](const core::type::Texture*) {
+                                out << " [[texture(" << binding_point->binding << ")]]";
+                            },
+                            [&](const core::type::Sampler*) {
+                                out << " [[sampler(" << binding_point->binding << ")]]";
+                            },
+                            TINT_ICE_ON_NO_MATCH);
                     }
                 }
             }
@@ -1040,7 +1052,6 @@
         switch (sc) {
             case core::AddressSpace::kFunction:
             case core::AddressSpace::kPrivate:
-            case core::AddressSpace::kHandle:
                 out << "thread";
                 break;
             case core::AddressSpace::kWorkgroup:
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars.cc b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
index 0cb07c1..00a569f 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
@@ -88,14 +88,32 @@
             ProcessFunction(*func);
         }
 
-        // Replace uses of each module-scope variable with pointers extracted from the structure.
+        // Replace uses of each module-scope variable with values extracted from the structure.
         uint32_t index = 0;
         for (auto& var : module_vars) {
-            var->Result(0)->ReplaceAllUsesWith([&](core::ir::Usage use) {  //
-                return GetPointerFromStruct(var, use.instruction, index);
+            Vector<core::ir::Instruction*, 16> to_destroy;
+            auto* ptr = var->Result(0)->Type()->As<core::type::Pointer>();
+            var->Result(0)->ForEachUse([&](core::ir::Usage use) {  //
+                auto* extracted_variable = GetVariableFromStruct(var, use.instruction, index);
+
+                // We drop the pointer from handle variables and store them in the struct by value
+                // instead, so remove any load instructions for the handle address space.
+                if (use.instruction->Is<core::ir::Load>() &&
+                    ptr->AddressSpace() == core::AddressSpace::kHandle) {
+                    use.instruction->Result(0)->ReplaceAllUsesWith(extracted_variable);
+                    to_destroy.Push(use.instruction);
+                    return;
+                }
+
+                use.instruction->SetOperand(use.operand_index, extracted_variable);
             });
             var->Destroy();
             index++;
+
+            // Clean up instructions that need to be removed.
+            for (auto* inst : to_destroy) {
+                inst->Destroy();
+            }
         }
     }
 
@@ -106,6 +124,13 @@
         for (auto* global : *ir.root_block) {
             if (auto* var = global->As<core::ir::Var>()) {
                 auto* type = var->Result(0)->Type();
+
+                // Handle types drop the pointer and are passed around by value.
+                auto* ptr = type->As<core::type::Pointer>();
+                if (ptr->AddressSpace() == core::AddressSpace::kHandle) {
+                    type = ptr->StoreType();
+                }
+
                 auto name = ir.NameOf(var);
                 if (!name) {
                     name = ir.symbols.New();
@@ -181,6 +206,14 @@
                         decl = param;
                         break;
                     }
+                    case core::AddressSpace::kHandle: {
+                        // Handle types become function parameters and drop the pointer.
+                        auto* param = b.FunctionParam(ptr->UnwrapPtr());
+                        param->SetBindingPoint(var->BindingPoint());
+                        func->AppendParam(param);
+                        decl = param;
+                        break;
+                    }
                     default:
                         TINT_UNREACHABLE() << "unhandled address space: " << ptr->AddressSpace();
                 }
@@ -218,18 +251,26 @@
         return param;
     }
 
-    /// Get a pointer from the module-scope variable replacement structure, inserting new access
+    /// Get a variable from the module-scope variable replacement structure, inserting new access
     /// instructions before @p inst.
     /// @param var the variable to get the replacement for
     /// @param inst the instruction that uses the variable
     /// @param index the index of the variable in the structure member list
-    /// @returns the pointer extracted from the structure
-    core::ir::Value* GetPointerFromStruct(core::ir::Var* var,
-                                          core::ir::Instruction* inst,
-                                          uint32_t index) {
+    /// @returns the variable extracted from the structure
+    core::ir::Value* GetVariableFromStruct(core::ir::Var* var,
+                                           core::ir::Instruction* inst,
+                                           uint32_t index) {
         auto* func = ContainingFunction(inst);
         auto* struct_value = function_to_struct_value.GetOr(func, nullptr);
-        auto* access = b.Access(var->Result(0)->Type(), struct_value, u32(index));
+        auto* type = var->Result(0)->Type();
+
+        // Handle types drop the pointer and are passed around by value.
+        auto* ptr = type->As<core::type::Pointer>();
+        if (ptr->AddressSpace() == core::AddressSpace::kHandle) {
+            type = ptr->StoreType();
+        }
+
+        auto* access = b.Access(type, struct_value, u32(index));
         access->InsertBefore(inst);
         return access->Result(0);
     }
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
index 17ffdc1..3bef228 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
@@ -30,6 +30,7 @@
 #include <utility>
 
 #include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/core/type/sampled_texture.h"
 
 using namespace tint::core::fluent_types;     // NOLINT
 using namespace tint::core::number_suffixes;  // NOLINT
@@ -306,6 +307,63 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_ModuleScopeVarsTest, HandleTypes) {
+    auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32());
+    auto* var_t = b.Var("t", ty.ptr<handle>(t));
+    auto* var_s = b.Var("s", ty.ptr<handle>(ty.sampler()));
+    var_t->SetBindingPoint(1, 2);
+    var_s->SetBindingPoint(3, 4);
+    mod.root_block->Append(var_t);
+    mod.root_block->Append(var_s);
+
+    auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+    b.Append(func->Block(), [&] {
+        auto* load_t = b.Load(var_t);
+        auto* load_s = b.Load(var_s);
+        b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, load_t, load_s, b.Splat<vec2<f32>>(0_f));
+        b.Return(func);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 2)
+  %s:ptr<handle, sampler, read> = var @binding_point(3, 4)
+}
+
+%foo = @fragment func():void {
+  $B2: {
+    %4:texture_2d<f32> = load %t
+    %5:sampler = load %s
+    %6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+tint_module_vars_struct = struct @align(1) {
+  t:texture_2d<f32> @offset(0)
+  s:sampler @offset(0)
+}
+
+%foo = @fragment func(%t:texture_2d<f32> [@binding_point(1, 2)], %s:sampler [@binding_point(3, 4)]):void {
+  $B1: {
+    %4:tint_module_vars_struct = construct %t, %s
+    %tint_module_vars:tint_module_vars_struct = let %4
+    %6:texture_2d<f32> = access %tint_module_vars, 0u
+    %7:sampler = access %tint_module_vars, 1u
+    %8:vec4<f32> = textureSample %6, %7, vec2<f32>(0.0f)
+    ret
+  }
+}
+)";
+
+    Run(ModuleScopeVars);
+
+    EXPECT_EQ(expect, str());
+}
+
 TEST_F(MslWriter_ModuleScopeVarsTest, MultipleAddressSpaces) {
     auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
     auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
@@ -607,6 +665,84 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_HandleTypes) {
+    auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32());
+    auto* var_t = b.Var("t", ty.ptr<handle>(t));
+    auto* var_s = b.Var("s", ty.ptr<handle>(ty.sampler()));
+    var_t->SetBindingPoint(1, 2);
+    var_s->SetBindingPoint(3, 4);
+    mod.root_block->Append(var_t);
+    mod.root_block->Append(var_s);
+
+    auto* foo = b.Function("foo", ty.vec4<f32>());
+    auto* param = b.FunctionParam<i32>("param");
+    foo->SetParams({param});
+    b.Append(foo->Block(), [&] {
+        auto* load_t = b.Load(var_t);
+        auto* load_s = b.Load(var_s);
+        auto* result = b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, load_t, load_s,
+                                         b.Splat<vec2<f32>>(0_f));
+        b.Return(foo, result);
+    });
+
+    auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+    b.Append(func->Block(), [&] {
+        b.Call(foo, 42_i);
+        b.Return(func);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 2)
+  %s:ptr<handle, sampler, read> = var @binding_point(3, 4)
+}
+
+%foo = func(%param:i32):vec4<f32> {
+  $B2: {
+    %5:texture_2d<f32> = load %t
+    %6:sampler = load %s
+    %7:vec4<f32> = textureSample %5, %6, vec2<f32>(0.0f)
+    ret %7
+  }
+}
+%main = @fragment func():void {
+  $B3: {
+    %9:vec4<f32> = call %foo, 42i
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+tint_module_vars_struct = struct @align(1) {
+  t:texture_2d<f32> @offset(0)
+  s:sampler @offset(0)
+}
+
+%foo = func(%param:i32, %tint_module_vars:tint_module_vars_struct):vec4<f32> {
+  $B1: {
+    %4:texture_2d<f32> = access %tint_module_vars, 0u
+    %5:sampler = access %tint_module_vars, 1u
+    %6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
+    ret %6
+  }
+}
+%main = @fragment func(%t:texture_2d<f32> [@binding_point(1, 2)], %s:sampler [@binding_point(3, 4)]):void {
+  $B2: {
+    %10:tint_module_vars_struct = construct %t, %s
+    %tint_module_vars_1:tint_module_vars_struct = let %10  # %tint_module_vars_1: 'tint_module_vars'
+    %12:vec4<f32> = call %foo, 42i, %tint_module_vars_1
+    ret
+  }
+}
+)";
+
+    Run(ModuleScopeVars);
+
+    EXPECT_EQ(expect, str());
+}
+
 TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_OutOfOrder) {
     auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
     auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
index 8f755f7..106221a 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
@@ -103,7 +103,7 @@
     TypedExpression& operator=(const TypedExpression&);
 
     /// @returns true if both type and expr are not nullptr
-    operator bool() const { return type && expr; }
+    explicit operator bool() const { return type && expr; }
 
     /// The type
     const Type* type = nullptr;
diff --git a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc
index 0dbbc5b..72efe20 100644
--- a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc
+++ b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc
@@ -119,7 +119,7 @@
         ArrayIndices array_indices;
 
         /// @returns true if the expr is not null (null usually indicates a failure)
-        operator bool() const { return expr != nullptr; }
+        explicit operator bool() const { return expr != nullptr; }
     };
 
     /// Statement holds information about a statement that will zero workgroup
diff --git a/src/tint/lang/wgsl/helpers/append_vector.cc b/src/tint/lang/wgsl/helpers/append_vector.cc
index dab1506..8ea2ebd 100644
--- a/src/tint/lang/wgsl/helpers/append_vector.cc
+++ b/src/tint/lang/wgsl/helpers/append_vector.cc
@@ -45,7 +45,7 @@
 struct VectorConstructorInfo {
     const sem::Call* call = nullptr;
     const sem::ValueConstructor* ctor = nullptr;
-    operator bool() const { return call != nullptr; }
+    explicit operator bool() const { return call != nullptr; }
 };
 VectorConstructorInfo AsVectorConstructor(const sem::ValueExpression* expr) {
     if (auto* call = expr->As<sem::Call>()) {
diff --git a/src/tint/lang/wgsl/reader/parser/parser.cc b/src/tint/lang/wgsl/reader/parser/parser.cc
index 3eddee4..b85555e 100644
--- a/src/tint/lang/wgsl/reader/parser/parser.cc
+++ b/src/tint/lang/wgsl/reader/parser/parser.cc
@@ -45,6 +45,7 @@
 #include "src/tint/lang/wgsl/ast/id_attribute.h"
 #include "src/tint/lang/wgsl/ast/if_statement.h"
 #include "src/tint/lang/wgsl/ast/increment_decrement_statement.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
 #include "src/tint/lang/wgsl/ast/invariant_attribute.h"
 #include "src/tint/lang/wgsl/ast/loop_statement.h"
 #include "src/tint/lang/wgsl/ast/return_statement.h"
@@ -3100,6 +3101,8 @@
             return create<ast::GroupAttribute>(t.source(), args[0]);
         case core::Attribute::kId:
             return create<ast::IdAttribute>(t.source(), args[0]);
+        case core::Attribute::kInputAttachmentIndex:
+            return create<ast::InputAttachmentIndexAttribute>(t.source(), args[0]);
         case core::Attribute::kInterpolate:
             return create<ast::InterpolateAttribute>(t.source(), args[0],
                                                      args.Length() == 2 ? args[1] : nullptr);
diff --git a/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc b/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc
index e7c48e4..c61f47d 100644
--- a/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc
@@ -647,5 +647,48 @@
     EXPECT_EQ(p->error(), "1:7: expected expression for group");
 }
 
+TEST_F(WGSLParserTest, Attribute_InputAttachmentIndex) {
+    auto p = parser("input_attachment_index(4)");
+    auto attr = p->attribute();
+    EXPECT_TRUE(attr.matched);
+    EXPECT_FALSE(attr.errored);
+    ASSERT_NE(attr.value, nullptr);
+    auto* var_attr = attr.value->As<ast::Attribute>();
+    ASSERT_FALSE(p->has_error());
+    ASSERT_NE(var_attr, nullptr);
+    ASSERT_TRUE(var_attr->Is<ast::InputAttachmentIndexAttribute>());
+
+    auto* group = var_attr->As<ast::InputAttachmentIndexAttribute>();
+    ASSERT_TRUE(group->expr->Is<ast::IntLiteralExpression>());
+    auto* expr = group->expr->As<ast::IntLiteralExpression>();
+    EXPECT_EQ(expr->value, 4);
+    EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
+}
+
+TEST_F(WGSLParserTest, Attribute_InputAttachmentIndex_expression) {
+    auto p = parser("input_attachment_index(4 + 5)");
+    auto attr = p->attribute();
+    EXPECT_TRUE(attr.matched);
+    EXPECT_FALSE(attr.errored);
+    ASSERT_NE(attr.value, nullptr);
+    auto* var_attr = attr.value->As<ast::Attribute>();
+    ASSERT_FALSE(p->has_error());
+    ASSERT_NE(var_attr, nullptr);
+    ASSERT_TRUE(var_attr->Is<ast::InputAttachmentIndexAttribute>());
+
+    auto* group = var_attr->As<ast::InputAttachmentIndexAttribute>();
+    ASSERT_TRUE(group->expr->Is<ast::BinaryExpression>());
+    auto* expr = group->expr->As<ast::BinaryExpression>();
+
+    EXPECT_EQ(core::BinaryOp::kAdd, expr->op);
+    auto* v = expr->lhs->As<ast::IntLiteralExpression>();
+    ASSERT_NE(nullptr, v);
+    EXPECT_EQ(v->value, 4u);
+
+    v = expr->rhs->As<ast::IntLiteralExpression>();
+    ASSERT_NE(nullptr, v);
+    EXPECT_EQ(v->value, 5u);
+}
+
 }  // namespace
 }  // namespace tint::wgsl::reader
diff --git a/src/tint/lang/wgsl/resolver/BUILD.bazel b/src/tint/lang/wgsl/resolver/BUILD.bazel
index 9893abb..e5ea130 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.bazel
+++ b/src/tint/lang/wgsl/resolver/BUILD.bazel
@@ -125,6 +125,7 @@
     "host_shareable_validation_test.cc",
     "increment_decrement_validation_test.cc",
     "inferred_type_test.cc",
+    "input_attachments_extension_test.cc",
     "is_host_shareable_test.cc",
     "is_storeable_test.cc",
     "language_features_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/BUILD.cmake b/src/tint/lang/wgsl/resolver/BUILD.cmake
index 429b9c0..326af91 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.cmake
+++ b/src/tint/lang/wgsl/resolver/BUILD.cmake
@@ -123,6 +123,7 @@
   lang/wgsl/resolver/host_shareable_validation_test.cc
   lang/wgsl/resolver/increment_decrement_validation_test.cc
   lang/wgsl/resolver/inferred_type_test.cc
+  lang/wgsl/resolver/input_attachments_extension_test.cc
   lang/wgsl/resolver/is_host_shareable_test.cc
   lang/wgsl/resolver/is_storeable_test.cc
   lang/wgsl/resolver/language_features_test.cc
diff --git a/src/tint/lang/wgsl/resolver/BUILD.gn b/src/tint/lang/wgsl/resolver/BUILD.gn
index 9456dbb..6fb89d0 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.gn
+++ b/src/tint/lang/wgsl/resolver/BUILD.gn
@@ -125,6 +125,7 @@
       "host_shareable_validation_test.cc",
       "increment_decrement_validation_test.cc",
       "inferred_type_test.cc",
+      "input_attachments_extension_test.cc",
       "is_host_shareable_test.cc",
       "is_storeable_test.cc",
       "language_features_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
index e94ebbe..809c8e3 100644
--- a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
@@ -65,6 +65,7 @@
     kDiagnostic,
     kGroup,
     kId,
+    kInputAttachmentIndex,
     kInterpolate,
     kInvariant,
     kLocation,
@@ -93,6 +94,8 @@
             return o << "@group";
         case AttributeKind::kId:
             return o << "@id";
+        case AttributeKind::kInputAttachmentIndex:
+            return o << "@input_attachment_index";
         case AttributeKind::kInterpolate:
             return o << "@interpolate";
         case AttributeKind::kInvariant:
@@ -166,6 +169,10 @@
                 "1:2 error: '@id' is not valid for " + thing,
             },
             TestParams{
+                {AttributeKind::kInputAttachmentIndex},
+                "1:2 error: '@input_attachment_index' is not valid for " + thing,
+            },
+            TestParams{
                 {AttributeKind::kInterpolate},
                 "1:2 error: '@interpolate' is not valid for " + thing,
             },
@@ -230,6 +237,8 @@
             return builder.Group(source, 1_a);
         case AttributeKind::kId:
             return builder.Id(source, 0_a);
+        case AttributeKind::kInputAttachmentIndex:
+            return builder.InputAttachmentIndex(source, 2_a);
         case AttributeKind::kBlendSrc:
             return builder.BlendSrc(source, 0_a);
         case AttributeKind::kInterpolate:
@@ -264,6 +273,9 @@
             case AttributeKind::kBlendSrc:
                 Enable(wgsl::Extension::kChromiumInternalDualSourceBlending);
                 break;
+            case AttributeKind::kInputAttachmentIndex:
+                Enable(wgsl::Extension::kChromiumInternalInputAttachments);
+                break;
             default:
                 break;
         }
@@ -345,6 +357,10 @@
             R"(1:2 error: '@id' is not valid for functions)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for functions)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for functions)",
         },
@@ -393,81 +409,86 @@
 
     CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         NonVoidFunctionAttributeTest,
-                         testing::Values(
-                             TestParams{
-                                 {AttributeKind::kAlign},
-                                 R"(1:2 error: '@align' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBinding},
-                                 R"(1:2 error: '@binding' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBlendSrc},
-                                 R"(1:2 error: '@blend_src' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBuiltinPosition},
-                                 R"(1:2 error: '@builtin' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kColor},
-                                 R"(1:2 error: '@color' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kDiagnostic},
-                                 Pass,
-                             },
-                             TestParams{
-                                 {AttributeKind::kGroup},
-                                 R"(1:2 error: '@group' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kId},
-                                 R"(1:2 error: '@id' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInterpolate},
-                                 R"(1:2 error: '@interpolate' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInvariant},
-                                 R"(1:2 error: '@invariant' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kLocation},
-                                 R"(1:2 error: '@location' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kMustUse},
-                                 Pass,
-                             },
-                             TestParams{
-                                 {AttributeKind::kOffset},
-                                 R"(1:2 error: '@offset' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kSize},
-                                 R"(1:2 error: '@size' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStageCompute},
-                                 R"(9:9 error: missing entry point IO attribute on return type)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
-                                 R"(9:9 error: missing entry point IO attribute on return type)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStride},
-                                 R"(1:2 error: '@stride' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kWorkgroupSize},
-                                 R"(1:2 error: '@workgroup_size' is only valid for compute stages)",
-                             }));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    NonVoidFunctionAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: '@align' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: '@binding' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBlendSrc},
+            R"(1:2 error: '@blend_src' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: '@builtin' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kColor},
+            R"(1:2 error: '@color' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: '@group' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: '@id' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: '@interpolate' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: '@invariant' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: '@location' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: '@offset' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: '@size' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(9:9 error: missing entry point IO attribute on return type)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
+            R"(9:9 error: missing entry point IO attribute on return type)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: '@stride' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: '@workgroup_size' is only valid for compute stages)",
+        }));
 }  // namespace FunctionTests
 
 namespace FunctionInputAndOutputTests {
@@ -520,6 +541,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for non-entry point function parameters)",
         },
@@ -605,6 +630,10 @@
             R"(1:2 error: '@id' is not valid for non-entry point function return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for non-entry point function return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for non-entry point function return types)",
         },
@@ -695,6 +724,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' cannot be used by compute shaders)",
         },
@@ -784,6 +817,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(9:9 error: missing entry point IO attribute on parameter)",
         },
@@ -887,6 +924,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(9:9 error: missing entry point IO attribute on parameter)",
         },
@@ -992,6 +1033,10 @@
             R"(1:2 error: '@id' is not valid for entry point return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' cannot be used by compute shaders)",
         },
@@ -1082,6 +1127,10 @@
             R"(1:2 error: '@id' is not valid for entry point return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(9:9 error: missing entry point IO attribute on return type)",
         },
@@ -1190,6 +1239,10 @@
             R"(1:2 error: '@id' is not valid for entry point return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' can only be used with '@location')",
         },
@@ -1315,6 +1368,10 @@
             R"(1:2 error: '@id' is not valid for 'struct' declarations)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'struct' declarations)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for 'struct' declarations)",
         },
@@ -1399,6 +1456,10 @@
             R"(1:2 error: '@id' is not valid for 'struct' members)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'struct' members)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' can only be used with '@location')",
         },
@@ -1635,82 +1696,87 @@
 
     CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         ArrayAttributeTest,
-                         testing::Values(
-                             TestParams{
-                                 {AttributeKind::kAlign},
-                                 R"(1:2 error: '@align' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBinding},
-                                 R"(1:2 error: '@binding' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBlendSrc},
-                                 R"(1:2 error: '@blend_src' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBuiltinPosition},
-                                 R"(1:2 error: '@builtin' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kDiagnostic},
-                                 R"(1:2 error: '@diagnostic' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kGroup},
-                                 R"(1:2 error: '@group' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kId},
-                                 R"(1:2 error: '@id' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInterpolate},
-                                 R"(1:2 error: '@interpolate' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInvariant},
-                                 R"(1:2 error: '@invariant' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kLocation},
-                                 R"(1:2 error: '@location' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kMustUse},
-                                 R"(1:2 error: '@must_use' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kOffset},
-                                 R"(1:2 error: '@offset' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kSize},
-                                 R"(1:2 error: '@size' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStageCompute},
-                                 R"(1:2 error: '@compute' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStride},
-                                 Pass,
-                             },
-                             TestParams{
-                                 {AttributeKind::kWorkgroupSize},
-                                 R"(1:2 error: '@workgroup_size' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBinding, AttributeKind::kGroup},
-                                 R"(1:2 error: '@binding' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStride, AttributeKind::kStride},
-                                 R"(3:4 error: duplicate stride attribute
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    ArrayAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: '@align' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: '@binding' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBlendSrc},
+            R"(1:2 error: '@blend_src' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: '@builtin' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: '@diagnostic' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: '@group' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: '@id' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: '@interpolate' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: '@invariant' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: '@location' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: '@must_use' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: '@offset' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: '@size' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: '@compute' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: '@workgroup_size' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            R"(1:2 error: '@binding' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kStride, AttributeKind::kStride},
+            R"(3:4 error: duplicate stride attribute
 1:2 note: first attribute declared here)",
-                             }));
+        }));
 
 using VariableAttributeTest = TestWithParams;
 TEST_P(VariableAttributeTest, IsValid) {
@@ -1862,6 +1928,10 @@
             R"(1:2 error: '@id' is not valid for 'const' declaration)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'const' declaration)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for 'const' declaration)",
         },
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph.cc b/src/tint/lang/wgsl/resolver/dependency_graph.cc
index 4838cde..bffa92e 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph.cc
@@ -52,6 +52,7 @@
 #include "src/tint/lang/wgsl/ast/identifier.h"
 #include "src/tint/lang/wgsl/ast/if_statement.h"
 #include "src/tint/lang/wgsl/ast/increment_decrement_statement.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
 #include "src/tint/lang/wgsl/ast/internal_attribute.h"
 #include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
 #include "src/tint/lang/wgsl/ast/invariant_attribute.h"
@@ -379,6 +380,7 @@
             [&](const ast::ColorAttribute* color) { TraverseExpression(color->expr); },
             [&](const ast::GroupAttribute* group) { TraverseExpression(group->expr); },
             [&](const ast::IdAttribute* id) { TraverseExpression(id->expr); },
+            [&](const ast::InputAttachmentIndexAttribute* idx) { TraverseExpression(idx->expr); },
             [&](const ast::BlendSrcAttribute* index) { TraverseExpression(index->expr); },
             [&](const ast::InterpolateAttribute* interpolate) {
                 TraverseExpression(interpolate->type);
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
index e435301..f64306a 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
@@ -1761,6 +1761,7 @@
     GlobalVar(Sym(), ty.sampler(core::type::SamplerKind::kSampler));
 
     GlobalVar(Sym(), ty.i32(), Vector{Binding(V), Group(V)});
+    GlobalVar(Sym(), ty.input_attachment(T), Vector{Binding(V), Group(V), InputAttachmentIndex(V)});
     GlobalVar(Sym(), ty.i32(), Vector{Location(V)});
     Override(Sym(), ty.i32(), Vector{Id(V)});
 
diff --git a/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
new file mode 100644
index 0000000..2b0449d
--- /dev/null
+++ b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
@@ -0,0 +1,166 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/wgsl/resolver/resolver.h"
+#include "src/tint/lang/wgsl/resolver/resolver_helper_test.h"
+
+namespace tint::resolver {
+namespace {
+
+using namespace tint::core::fluent_types;     // NOLINT
+using namespace tint::core::number_suffixes;  // NOLINT
+
+using InputAttachmenExtensionTest = ResolverTest;
+
+// Test that input_attachment cannot be used without extension.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentWithoutExtension) {
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : input_attachment<f32>;
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(error: use of 'input_attachment' requires enabling extension 'chromium_internal_input_attachments')");
+}
+
+// Test that input_attachment cannot be declared locally.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentLocalDecl) {
+    // enable chromium_internal_input_attachments;
+    // @fragment fn f() {
+    //    var input_tex : input_attachment<f32>;
+    // }
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    Func("f", Empty, ty.void_(),
+         Vector{
+             Decl(Var("input_tex", ty.input_attachment(ty.Of<f32>()))),
+         },
+         Vector{Stage(ast::PipelineStage::kFragment)});
+    EXPECT_FALSE(r()->Resolve());
+}
+
+// Test that input_attachment cannot be declared without index.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentWithoutIndex) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()), Vector{Binding(0_u), Group(0_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(error: 'input_attachment' variables require '@input_attachment_index' attribute)");
+}
+
+// Test that Resolver can get input_attachment_index value.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexValue) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    auto* ast_var = GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+                              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem_var = Sem().Get<sem::GlobalVariable>(ast_var);
+    ASSERT_NE(sem_var, nullptr);
+    EXPECT_EQ(sem_var->Attributes().input_attachment_index, 3u);
+}
+
+// Test that @input_attachment_index cannot be used without extension.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexWithoutExtension) {
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : texture_2d<f32>;
+
+    GlobalVar("input_tex", ty.sampled_texture(core::type::TextureDimension::k2d, ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(error: use of '@input_attachment_index' requires enabling extension 'chromium_internal_input_attachments')");
+}
+
+// Test that input_attachment_index's value cannot be float.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexInvalidValueType) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(3.0)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_f)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(error: '@input_attachment_index' must be an 'i32' or 'u32' value)");
+}
+
+// Test that input_attachment_index's value cannot be negative.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexNegative) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(-2)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(core::i32(-2))});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(error: '@input_attachment_index' value must be non-negative)");
+}
+
+// Test that input_attachment_index cannot be used on non input_attachment variable.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexInvalidType) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : texture_2d<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.sampled_texture(core::type::TextureDimension::k2d, ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(error: cannot apply '@input_attachment_index' to declaration of type 'texture_2d<f32>'
+note: '@input_attachment_index' must only be applied to declarations of 'input_attachment' type)");
+}
+
+}  // namespace
+}  // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index 0a17841..fe24ffc 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -65,6 +65,7 @@
 #include "src/tint/lang/wgsl/ast/for_loop_statement.h"
 #include "src/tint/lang/wgsl/ast/id_attribute.h"
 #include "src/tint/lang/wgsl/ast/if_statement.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
 #include "src/tint/lang/wgsl/ast/internal_attribute.h"
 #include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
 #include "src/tint/lang/wgsl/ast/loop_statement.h"
@@ -614,7 +615,7 @@
         bool has_io_address_space = sem->AddressSpace() == core::AddressSpace::kIn ||
                                     sem->AddressSpace() == core::AddressSpace::kOut;
 
-        std::optional<uint32_t> group, binding;
+        std::optional<uint32_t> group, binding, input_attachment_index;
         for (auto* attribute : var->attributes) {
             Mark(attribute);
             enum Status { kSuccess, kErrored, kInvalid };
@@ -636,6 +637,14 @@
                     group = value.Get();
                     return kSuccess;
                 },
+                [&](const ast::InputAttachmentIndexAttribute* attr) {
+                    auto value = InputAttachmentIndexAttribute(attr);
+                    if (value != Success) {
+                        return kErrored;
+                    }
+                    input_attachment_index = value.Get();
+                    return kSuccess;
+                },
                 [&](const ast::LocationAttribute* attr) {
                     if (!has_io_address_space) {
                         return kInvalid;
@@ -708,6 +717,10 @@
             global->Attributes().binding_point = BindingPoint{group.value(), binding.value()};
         }
 
+        if (input_attachment_index) {
+            global->Attributes().input_attachment_index = input_attachment_index;
+        }
+
     } else {
         for (auto* attribute : var->attributes) {
             Mark(attribute);
@@ -3876,6 +3889,31 @@
     return static_cast<uint32_t>(value);
 }
 
+tint::Result<uint32_t> Resolver::InputAttachmentIndexAttribute(
+    const ast::InputAttachmentIndexAttribute* attr) {
+    ExprEvalStageConstraint constraint{core::EvaluationStage::kConstant, "@input_attachment_index"};
+    TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
+
+    auto* materialized = Materialize(ValueExpression(attr->expr));
+    if (!materialized) {
+        return Failure{};
+    }
+    if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
+        AddError(attr->source) << style::Attribute("@input_attachment_index") << " must be an "
+                               << style::Type("i32") << " or " << style::Type("u32") << " value";
+        return Failure{};
+    }
+
+    auto const_value = materialized->ConstantValue();
+    auto value = const_value->ValueAs<AInt>();
+    if (value < 0) {
+        AddError(attr->source) << style::Attribute("@input_attachment_index")
+                               << " value must be non-negative";
+        return Failure{};
+    }
+    return static_cast<uint32_t>(value);
+}
+
 tint::Result<sem::WorkgroupSize> Resolver::WorkgroupAttribute(const ast::WorkgroupAttribute* attr) {
     // Set work-group size defaults.
     sem::WorkgroupSize ws;
diff --git a/src/tint/lang/wgsl/resolver/resolver.h b/src/tint/lang/wgsl/resolver/resolver.h
index df18f91..1781493 100644
--- a/src/tint/lang/wgsl/resolver/resolver.h
+++ b/src/tint/lang/wgsl/resolver/resolver.h
@@ -442,6 +442,11 @@
     /// @returns the group value on success.
     tint::Result<uint32_t> GroupAttribute(const ast::GroupAttribute* attr);
 
+    /// Resolves the `@input_attachment_index` attribute @p attr
+    /// @returns the index value on success.
+    tint::Result<uint32_t> InputAttachmentIndexAttribute(
+        const ast::InputAttachmentIndexAttribute* attr);
+
     /// Resolves the `@workgroup_size` attribute @p attr
     /// @returns the workgroup size on success.
     tint::Result<sem::WorkgroupSize> WorkgroupAttribute(const ast::WorkgroupAttribute* attr);
diff --git a/src/tint/lang/wgsl/resolver/uniformity.cc b/src/tint/lang/wgsl/resolver/uniformity.cc
index d6a5307..2faeb21 100644
--- a/src/tint/lang/wgsl/resolver/uniformity.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity.cc
@@ -1188,14 +1188,21 @@
                                                    const ast::IdentifierExpression* ident,
                                                    bool load_rule = false) {
         // Helper to check if the entry point attribute of `obj` indicates non-uniformity.
-        auto has_nonuniform_entry_point_attribute = [&](auto* obj) {
-            // Only the num_workgroups and workgroup_id builtins are uniform.
+        auto has_nonuniform_entry_point_attribute = [&](auto* obj, auto* entry_point) {
+            // Only the num_workgroups and workgroup_id builtins, and subgroup_size builtin used in
+            // compute stage are uniform.
             if (auto* builtin_attr = ast::GetAttribute<ast::BuiltinAttribute>(obj->attributes)) {
                 auto builtin = b.Sem().Get(builtin_attr)->Value();
                 if (builtin == core::BuiltinValue::kNumWorkgroups ||
                     builtin == core::BuiltinValue::kWorkgroupId) {
                     return false;
                 }
+                if (builtin == core::BuiltinValue::kSubgroupSize) {
+                    // Currently Tint only allow using subgroup_size builtin as a compute shader
+                    // input.
+                    TINT_ASSERT(entry_point->PipelineStage() == ast::PipelineStage::kCompute);
+                    return false;
+                }
             }
             return true;
         };
@@ -1216,14 +1223,16 @@
                         // is non-uniform.
                         bool uniform = true;
                         for (auto* member : str->Members()) {
-                            if (has_nonuniform_entry_point_attribute(member->Declaration())) {
+                            if (has_nonuniform_entry_point_attribute(member->Declaration(),
+                                                                     user_func->Declaration())) {
                                 uniform = false;
                             }
                         }
                         node->AddEdge(uniform ? cf : current_function_->may_be_non_uniform);
                         return std::make_pair(cf, node);
                     } else {
-                        if (has_nonuniform_entry_point_attribute(param->Declaration())) {
+                        if (has_nonuniform_entry_point_attribute(param->Declaration(),
+                                                                 user_func->Declaration())) {
                             node->AddEdge(current_function_->may_be_non_uniform);
                         } else {
                             node->AddEdge(cf);
diff --git a/src/tint/lang/wgsl/resolver/uniformity_test.cc b/src/tint/lang/wgsl/resolver/uniformity_test.cc
index 9c15fa6..092cfd0 100644
--- a/src/tint/lang/wgsl/resolver/uniformity_test.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity_test.cc
@@ -514,7 +514,11 @@
 class ComputeBuiltin : public UniformityAnalysisTestBase,
                        public ::testing::TestWithParam<BuiltinEntry> {};
 TEST_P(ComputeBuiltin, AsParam) {
-    std::string src = R"(
+    std::string src = std::string((GetParam().name == "subgroup_size")
+                                      ? R"(enable chromium_experimental_subgroups;
+)"
+                                      : "") +
+                      R"(
 @compute @workgroup_size(64)
 fn main(@builtin()" + GetParam().name +
                       R"() b : )" + GetParam().type + R"() {
@@ -545,7 +549,11 @@
 }
 
 TEST_P(ComputeBuiltin, InStruct) {
-    std::string src = R"(
+    std::string src = std::string((GetParam().name == "subgroup_size")
+                                      ? R"(enable chromium_experimental_subgroups;
+)"
+                                      : "") +
+                      R"(
 struct S {
   @builtin()" + GetParam().name +
                       R"() b : )" + GetParam().type + R"(
@@ -585,7 +593,8 @@
                                            BuiltinEntry{"local_invocation_index", "u32", false},
                                            BuiltinEntry{"global_invocation_id", "vec3<u32>", false},
                                            BuiltinEntry{"workgroup_id", "vec3<u32>", true},
-                                           BuiltinEntry{"num_workgroups", "vec3<u32>", true}),
+                                           BuiltinEntry{"num_workgroups", "vec3<u32>", true},
+                                           BuiltinEntry{"subgroup_size", "u32", true}),
                          [](const ::testing::TestParamInfo<ComputeBuiltin::ParamType>& p) {
                              return p.param.name;
                          });
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index f482497..559aa50 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -442,6 +442,12 @@
 }
 
 bool Validator::InputAttachment(const core::type::InputAttachment* t, const Source& source) const {
+    if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalInputAttachments)) {
+        AddError(source) << "use of " << style::Type("input_attachment")
+                         << " requires enabling extension "
+                         << style::Code("chromium_internal_input_attachments");
+        return false;
+    }
     if (!t->type()->UnwrapRef()->IsAnyOf<core::type::F32, core::type::I32, core::type::U32>()) {
         AddError(source) << "input_attachment<type>: type must be f32, i32 or u32";
         return false;
@@ -450,6 +456,29 @@
     return true;
 }
 
+bool Validator::InputAttachmentIndexAttribute(const ast::InputAttachmentIndexAttribute* attr,
+                                              const core::type::Type* type,
+                                              const Source& source) const {
+    if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalInputAttachments)) {
+        AddError(source) << "use of " << style::Attribute("@input_attachment_index")
+                         << " requires enabling extension "
+                         << style::Code("chromium_internal_input_attachments");
+        return false;
+    }
+
+    if (!type->Is<core::type::InputAttachment>()) {
+        std::string invalid_type = sem_.TypeNameOf(type);
+        AddError(source) << "cannot apply " << style::Attribute("@input_attachment_index")
+                         << " to declaration of type " << style::Type(invalid_type);
+        AddNote(attr->source) << style::Attribute("@input_attachment_index")
+                              << " must only be applied to declarations of "
+                              << style::Type("input_attachment") << " type";
+        return false;
+    }
+
+    return true;
+}
+
 bool Validator::Materialize(const core::type::Type* to,
                             const core::type::Type* from,
                             const Source& source) const {
@@ -708,6 +737,13 @@
         return false;
     }
 
+    auto* input_attachment_index_attr =
+        ast::GetAttribute<ast::InputAttachmentIndexAttribute>(decl->attributes);
+    if (input_attachment_index_attr &&
+        !InputAttachmentIndexAttribute(input_attachment_index_attr, global->Type()->UnwrapRef(),
+                                       decl->source)) {
+        return false;
+    }
     switch (global->AddressSpace()) {
         case core::AddressSpace::kUniform:
         case core::AddressSpace::kStorage:
@@ -720,6 +756,13 @@
                     << style::Attribute("@binding") << " attributes";
                 return false;
             }
+            if (global->Type()->UnwrapRef()->Is<core::type::InputAttachment>() &&
+                !input_attachment_index_attr) {
+                AddError(decl->source)
+                    << style::Type("input_attachment") << " variables require "
+                    << style::Attribute("@input_attachment_index") << " attribute";
+                return false;
+            }
             break;
         }
         default: {
diff --git a/src/tint/lang/wgsl/resolver/validator.h b/src/tint/lang/wgsl/resolver/validator.h
index 6cdbfaf..fb54048 100644
--- a/src/tint/lang/wgsl/resolver/validator.h
+++ b/src/tint/lang/wgsl/resolver/validator.h
@@ -434,6 +434,15 @@
     /// @returns true on success, false otherwise
     bool InputAttachment(const core::type::InputAttachment* t, const Source& source) const;
 
+    /// Validates a input attachment index attribute
+    /// @param attr the input attachment index attribute to validate
+    /// @param type the variable type
+    /// @param source the source of declaration using the attribute
+    /// @returns true on success, false otherwise.
+    bool InputAttachmentIndexAttribute(const ast::InputAttachmentIndexAttribute* attr,
+                                       const core::type::Type* type,
+                                       const Source& source) const;
+
     /// Validates a structure
     /// @param str the structure to validate
     /// @param stage the current pipeline stage
diff --git a/src/tint/lang/wgsl/sem/variable.h b/src/tint/lang/wgsl/sem/variable.h
index bf97929..51cac2a 100644
--- a/src/tint/lang/wgsl/sem/variable.h
+++ b/src/tint/lang/wgsl/sem/variable.h
@@ -169,6 +169,8 @@
     /// @note a GlobalVariable generally doesn't have a `color` in WGSL, as it isn't allowed by
     /// the spec. The location maybe attached by transforms such as CanonicalizeEntryPointIO.
     std::optional<uint32_t> color;
+    /// The `input_attachment_index` attribute value for the variable, if set
+    std::optional<uint32_t> input_attachment_index;
 };
 
 /// GlobalVariable is a module-scope variable
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index b404397..7243f25 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -1069,9 +1069,7 @@
     /// name.
     void Bind(const core::ir::Value* value, Symbol name) {
         TINT_ASSERT(value);
-
-        bool added = bindings_.Add(value, VariableValue{name});
-        if (TINT_UNLIKELY(!added)) {
+        if (TINT_UNLIKELY(!bindings_.Add(value, VariableValue{name}))) {
             TINT_ICE() << "Bind(" << value->TypeInfo().name << ") called twice for same value";
         }
     }
diff --git a/src/tint/utils/command/command_posix.cc b/src/tint/utils/command/command_posix.cc
index f3c30de..fae21be 100644
--- a/src/tint/utils/command/command_posix.cc
+++ b/src/tint/utils/command/command_posix.cc
@@ -75,7 +75,7 @@
     operator int() { return handle_; }
 
     /// @returns true if the file is not closed
-    operator bool() { return handle_ != kClosed; }
+    explicit operator bool() { return handle_ != kClosed; }
 
   private:
     File(const File&) = delete;
@@ -103,7 +103,7 @@
     }
 
     /// @returns true if the pipe has an open read or write file
-    operator bool() { return read || write; }
+    explicit operator bool() { return read || write; }
 
     /// The reader end of the pipe
     File read;
diff --git a/src/tint/utils/command/command_windows.cc b/src/tint/utils/command/command_windows.cc
index a6a577b..e591416 100644
--- a/src/tint/utils/command/command_windows.cc
+++ b/src/tint/utils/command/command_windows.cc
@@ -73,7 +73,7 @@
     operator HANDLE() { return handle_; }
 
     /// @returns true if the handle is not invalid
-    operator bool() { return handle_ != nullptr; }
+    explicit operator bool() { return handle_ != nullptr; }
 
   private:
     Handle(const Handle&) = delete;
@@ -106,7 +106,7 @@
     }
 
     /// @returns true if the pipe has an open read or write file
-    operator bool() { return read || write; }
+    explicit operator bool() { return read || write; }
 
     /// The reader end of the pipe
     Handle read;
diff --git a/src/tint/utils/containers/hashmap.h b/src/tint/utils/containers/hashmap.h
index 69790f3..df1ed08 100644
--- a/src/tint/utils/containers/hashmap.h
+++ b/src/tint/utils/containers/hashmap.h
@@ -90,7 +90,7 @@
     T* value = nullptr;
 
     /// @returns `true` if #value is not null.
-    operator bool() const { return value; }
+    explicit operator bool() const { return value; }
 
     /// @returns the dereferenced value, which must not be null.
     T& operator*() const { return *value; }
@@ -129,7 +129,7 @@
     bool added = false;
 
     /// @returns #added
-    operator bool() const { return added; }
+    explicit operator bool() const { return added; }
 };
 
 /// An unordered hashmap, with a fixed-size capacity that avoids heap allocations.
diff --git a/src/tint/utils/containers/hashmap_test.cc b/src/tint/utils/containers/hashmap_test.cc
index af926a2..f3b91bd 100644
--- a/src/tint/utils/containers/hashmap_test.cc
+++ b/src/tint/utils/containers/hashmap_test.cc
@@ -421,7 +421,7 @@
         switch (rnd() % 7) {
             case 0: {  // Add
                 auto expected = reference.emplace(key, value).second;
-                EXPECT_EQ(map.Add(key, value), expected) << "i:" << i;
+                EXPECT_EQ(map.Add(key, value).added, expected) << "i:" << i;
                 EXPECT_EQ(map.Get(key), value) << "i:" << i;
                 EXPECT_TRUE(map.Contains(key)) << "i:" << i;
                 break;
diff --git a/src/tint/utils/file/tmpfile.h b/src/tint/utils/file/tmpfile.h
index 0ae8e95..6bcd663 100644
--- a/src/tint/utils/file/tmpfile.h
+++ b/src/tint/utils/file/tmpfile.h
@@ -50,7 +50,7 @@
     ~TmpFile();
 
     /// @return true if the temporary file was successfully created.
-    operator bool() { return !path_.empty(); }
+    explicit operator bool() { return !path_.empty(); }
 
     /// @return the path to the temporary file
     std::string Path() const { return path_; }
diff --git a/src/tint/utils/id/generation_id.h b/src/tint/utils/id/generation_id.h
index 018b4a2..b3078df 100644
--- a/src/tint/utils/id/generation_id.h
+++ b/src/tint/utils/id/generation_id.h
@@ -67,7 +67,7 @@
     uint32_t Value() const { return val; }
 
     /// @returns true if this GenerationID is valid
-    operator bool() const { return val != 0; }
+    explicit operator bool() const { return val != 0; }
 
   private:
     explicit GenerationID(uint32_t);
diff --git a/src/tint/utils/symbol/symbol.h b/src/tint/utils/symbol/symbol.h
index 9f26bb0..10b7bd7 100644
--- a/src/tint/utils/symbol/symbol.h
+++ b/src/tint/utils/symbol/symbol.h
@@ -83,7 +83,7 @@
     bool IsValid() const { return val_ != static_cast<uint32_t>(-1); }
 
     /// @returns true if the symbol is valid
-    operator bool() const { return IsValid(); }
+    explicit operator bool() const { return IsValid(); }
 
     /// @returns the value for the symbol
     uint32_t value() const { return val_; }