diff --git a/src/common/HashUtils.h b/src/common/HashUtils.h
index a316dbc..9d10ca7 100644
--- a/src/common/HashUtils.h
+++ b/src/common/HashUtils.h
@@ -17,6 +17,7 @@
 
 #include "common/Platform.h"
 #include "common/TypedInteger.h"
+#include "common/ityp_bitset.h"
 
 #include <bitset>
 #include <functional>
@@ -87,4 +88,14 @@
 }
 #endif
 
+namespace std {
+    template <typename Index, size_t N>
+    class hash<ityp::bitset<Index, N>> {
+      public:
+        size_t operator()(const ityp::bitset<Index, N>& value) const {
+            return Hash(static_cast<const std::bitset<N>&>(value));
+        }
+    };
+}  // namespace std
+
 #endif  // COMMON_HASHUTILS_H_
diff --git a/src/common/TypedInteger.h b/src/common/TypedInteger.h
index dd06331..5474d9a 100644
--- a/src/common/TypedInteger.h
+++ b/src/common/TypedInteger.h
@@ -62,12 +62,14 @@
 
 namespace detail {
     template <typename Tag, typename T>
-    class TypedIntegerImpl {
+    class alignas(T) TypedIntegerImpl {
         static_assert(std::is_integral<T>::value, "TypedInteger must be integral");
         T mValue;
 
       public:
         constexpr TypedIntegerImpl() : mValue(0) {
+            static_assert(alignof(TypedIntegerImpl) == alignof(T), "");
+            static_assert(sizeof(TypedIntegerImpl) == sizeof(T), "");
         }
 
         // Construction from non-narrowing integral types.
diff --git a/src/common/ityp_bitset.h b/src/common/ityp_bitset.h
index 46badc9..ef351d4 100644
--- a/src/common/ityp_bitset.h
+++ b/src/common/ityp_bitset.h
@@ -58,6 +58,14 @@
         using Base::none;
         using Base::size;
 
+        bool operator==(const bitset& other) const noexcept {
+            return Base::operator==(static_cast<const Base&>(other));
+        }
+
+        bool operator!=(const bitset& other) const noexcept {
+            return Base::operator!=(static_cast<const Base&>(other));
+        }
+
         bitset& operator&=(const bitset& other) noexcept {
             return static_cast<bitset&>(Base::operator&=(static_cast<const Base&>(other)));
         }
@@ -117,6 +125,8 @@
         friend BitSetIterator<N, Index> IterateBitSet(const bitset& bitset) {
             return BitSetIterator<N, Index>(static_cast<const Base&>(bitset));
         }
+
+        friend class std::hash<bitset>;
     };
 
 }  // namespace ityp
diff --git a/src/dawn_native/BindGroupAndStorageBarrierTracker.h b/src/dawn_native/BindGroupAndStorageBarrierTracker.h
index f04dd64..e17d263 100644
--- a/src/dawn_native/BindGroupAndStorageBarrierTracker.h
+++ b/src/dawn_native/BindGroupAndStorageBarrierTracker.h
@@ -32,10 +32,12 @@
       public:
         BindGroupAndStorageBarrierTrackerBase() = default;
 
-        void OnSetBindGroup(uint32_t index,
+        void OnSetBindGroup(BindGroupIndex index,
                             BindGroupBase* bindGroup,
                             uint32_t dynamicOffsetCount,
                             uint32_t* dynamicOffsets) {
+            ASSERT(index < kMaxBindGroupsTyped);
+
             if (this->mBindGroups[index] != bindGroup) {
                 mBindings[index] = {};
                 mBindingsNeedingBarrier[index] = {};
@@ -89,12 +91,16 @@
         }
 
       protected:
-        std::array<ityp::bitset<BindingIndex, kMaxBindingsPerGroup>, kMaxBindGroups>
-            mBindingsNeedingBarrier = {};
-        std::array<ityp::array<BindingIndex, wgpu::BindingType, kMaxBindingsPerGroup>,
-                   kMaxBindGroups>
+        ityp::
+            array<BindGroupIndex, ityp::bitset<BindingIndex, kMaxBindingsPerGroup>, kMaxBindGroups>
+                mBindingsNeedingBarrier = {};
+        ityp::array<BindGroupIndex,
+                    ityp::array<BindingIndex, wgpu::BindingType, kMaxBindingsPerGroup>,
+                    kMaxBindGroups>
             mBindingTypes = {};
-        std::array<ityp::array<BindingIndex, ObjectBase*, kMaxBindingsPerGroup>, kMaxBindGroups>
+        ityp::array<BindGroupIndex,
+                    ityp::array<BindingIndex, ObjectBase*, kMaxBindingsPerGroup>,
+                    kMaxBindGroups>
             mBindings = {};
     };
 
diff --git a/src/dawn_native/BindGroupTracker.h b/src/dawn_native/BindGroupTracker.h
index 121dd0f..8d03ebf 100644
--- a/src/dawn_native/BindGroupTracker.h
+++ b/src/dawn_native/BindGroupTracker.h
@@ -32,11 +32,11 @@
     template <bool CanInheritBindGroups, typename DynamicOffset>
     class BindGroupTrackerBase {
       public:
-        void OnSetBindGroup(uint32_t index,
+        void OnSetBindGroup(BindGroupIndex index,
                             BindGroupBase* bindGroup,
                             uint32_t dynamicOffsetCount,
                             uint32_t* dynamicOffsets) {
-            ASSERT(index < kMaxBindGroups);
+            ASSERT(index < kMaxBindGroupsTyped);
 
             if (mBindGroupLayoutsMask[index]) {
                 // It is okay to only dirty bind groups that are used by the current pipeline
@@ -73,7 +73,7 @@
             // the first |k| matching bind groups may be inherited.
             if (CanInheritBindGroups && mLastAppliedPipelineLayout != nullptr) {
                 // Dirty bind groups that cannot be inherited.
-                std::bitset<kMaxBindGroups> dirtiedGroups =
+                BindGroupLayoutMask dirtiedGroups =
                     ~mPipelineLayout->InheritedGroupsMask(mLastAppliedPipelineLayout);
 
                 mDirtyBindGroups |= dirtiedGroups;
@@ -98,12 +98,12 @@
             mLastAppliedPipelineLayout = mPipelineLayout;
         }
 
-        std::bitset<kMaxBindGroups> mDirtyBindGroups = 0;
-        std::bitset<kMaxBindGroups> mDirtyBindGroupsObjectChangedOrIsDynamic = 0;
-        std::bitset<kMaxBindGroups> mBindGroupLayoutsMask = 0;
-        std::array<BindGroupBase*, kMaxBindGroups> mBindGroups = {};
-        std::array<uint32_t, kMaxBindGroups> mDynamicOffsetCounts = {};
-        std::array<std::array<DynamicOffset, kMaxBindingsPerGroup>, kMaxBindGroups>
+        BindGroupLayoutMask mDirtyBindGroups = 0;
+        BindGroupLayoutMask mDirtyBindGroupsObjectChangedOrIsDynamic = 0;
+        BindGroupLayoutMask mBindGroupLayoutsMask = 0;
+        ityp::array<BindGroupIndex, BindGroupBase*, kMaxBindGroups> mBindGroups = {};
+        ityp::array<BindGroupIndex, uint32_t, kMaxBindGroups> mDynamicOffsetCounts = {};
+        ityp::array<BindGroupIndex, std::array<DynamicOffset, kMaxBindingsPerGroup>, kMaxBindGroups>
             mDynamicOffsets = {};
 
         // |mPipelineLayout| is the current pipeline layout set on the command buffer.
diff --git a/src/dawn_native/BindingInfo.h b/src/dawn_native/BindingInfo.h
index 5cd6aad..cac4f4c 100644
--- a/src/dawn_native/BindingInfo.h
+++ b/src/dawn_native/BindingInfo.h
@@ -17,6 +17,7 @@
 
 #include "common/Constants.h"
 #include "common/TypedInteger.h"
+#include "common/ityp_array.h"
 #include "dawn_native/Format.h"
 #include "dawn_native/dawn_platform.h"
 
@@ -30,7 +31,10 @@
     // Binding numbers get mapped to a packed range of indices
     using BindingIndex = TypedInteger<struct BindingIndexT, uint32_t>;
 
+    using BindGroupIndex = TypedInteger<struct BindGroupIndexT, uint32_t>;
+
     static constexpr BindingIndex kMaxBindingsPerGroupTyped = BindingIndex(kMaxBindingsPerGroup);
+    static constexpr BindGroupIndex kMaxBindGroupsTyped = BindGroupIndex(kMaxBindGroups);
 
     struct BindingInfo {
         BindingNumber binding;
@@ -45,7 +49,7 @@
     };
 
     // For buffer size validation
-    using RequiredBufferSizes = std::array<std::vector<uint64_t>, kMaxBindGroups>;
+    using RequiredBufferSizes = ityp::array<BindGroupIndex, std::vector<uint64_t>, kMaxBindGroups>;
 
 }  // namespace dawn_native
 
diff --git a/src/dawn_native/CommandBufferStateTracker.cpp b/src/dawn_native/CommandBufferStateTracker.cpp
index 20643ed..7c4a327 100644
--- a/src/dawn_native/CommandBufferStateTracker.cpp
+++ b/src/dawn_native/CommandBufferStateTracker.cpp
@@ -100,7 +100,7 @@
         if (aspects[VALIDATION_ASPECT_BIND_GROUPS]) {
             bool matches = true;
 
-            for (uint32_t i : IterateBitSet(mLastPipelineLayout->GetBindGroupLayoutsMask())) {
+            for (BindGroupIndex i : IterateBitSet(mLastPipelineLayout->GetBindGroupLayoutsMask())) {
                 if (mBindgroups[i] == nullptr ||
                     mLastPipelineLayout->GetBindGroupLayout(i) != mBindgroups[i]->GetLayout() ||
                     !BufferSizesAtLeastAsBig(mBindgroups[i]->GetUnverifiedBufferSizes(),
@@ -140,18 +140,19 @@
         }
 
         if (aspects[VALIDATION_ASPECT_BIND_GROUPS]) {
-            for (uint32_t i : IterateBitSet(mLastPipelineLayout->GetBindGroupLayoutsMask())) {
+            for (BindGroupIndex i : IterateBitSet(mLastPipelineLayout->GetBindGroupLayoutsMask())) {
                 if (mBindgroups[i] == nullptr) {
-                    return DAWN_VALIDATION_ERROR("Missing bind group " + std::to_string(i));
+                    return DAWN_VALIDATION_ERROR("Missing bind group " +
+                                                 std::to_string(static_cast<uint32_t>(i)));
                 } else if (mLastPipelineLayout->GetBindGroupLayout(i) !=
                            mBindgroups[i]->GetLayout()) {
                     return DAWN_VALIDATION_ERROR(
                         "Pipeline and bind group layout doesn't match for bind group " +
-                        std::to_string(i));
+                        std::to_string(static_cast<uint32_t>(i)));
                 } else if (!BufferSizesAtLeastAsBig(mBindgroups[i]->GetUnverifiedBufferSizes(),
                                                     (*mMinimumBufferSizes)[i])) {
                     return DAWN_VALIDATION_ERROR("Binding sizes too small for bind group " +
-                                                 std::to_string(i));
+                                                 std::to_string(static_cast<uint32_t>(i)));
                 }
             }
 
@@ -179,7 +180,7 @@
         SetPipelineCommon(pipeline);
     }
 
-    void CommandBufferStateTracker::SetBindGroup(uint32_t index, BindGroupBase* bindgroup) {
+    void CommandBufferStateTracker::SetBindGroup(BindGroupIndex index, BindGroupBase* bindgroup) {
         mBindgroups[index] = bindgroup;
         mAspects.reset(VALIDATION_ASPECT_BIND_GROUPS);
     }
diff --git a/src/dawn_native/CommandBufferStateTracker.h b/src/dawn_native/CommandBufferStateTracker.h
index 478429c..39d32fd 100644
--- a/src/dawn_native/CommandBufferStateTracker.h
+++ b/src/dawn_native/CommandBufferStateTracker.h
@@ -16,11 +16,11 @@
 #define DAWNNATIVE_COMMANDBUFFERSTATETRACKER_H
 
 #include "common/Constants.h"
+#include "common/ityp_array.h"
 #include "dawn_native/BindingInfo.h"
 #include "dawn_native/Error.h"
 #include "dawn_native/Forward.h"
 
-#include <array>
 #include <bitset>
 #include <map>
 #include <set>
@@ -37,7 +37,7 @@
         // State-modifying methods
         void SetComputePipeline(ComputePipelineBase* pipeline);
         void SetRenderPipeline(RenderPipelineBase* pipeline);
-        void SetBindGroup(uint32_t index, BindGroupBase* bindgroup);
+        void SetBindGroup(BindGroupIndex index, BindGroupBase* bindgroup);
         void SetIndexBuffer();
         void SetVertexBuffer(uint32_t slot);
 
@@ -53,7 +53,7 @@
 
         ValidationAspects mAspects;
 
-        std::array<BindGroupBase*, kMaxBindGroups> mBindgroups = {};
+        ityp::array<BindGroupIndex, BindGroupBase*, kMaxBindGroups> mBindgroups = {};
         std::bitset<kMaxVertexBuffers> mVertexBufferSlotsUsed;
 
         PipelineLayoutBase* mLastPipelineLayout = nullptr;
diff --git a/src/dawn_native/Commands.h b/src/dawn_native/Commands.h
index 6ccd2f3..155a21d 100644
--- a/src/dawn_native/Commands.h
+++ b/src/dawn_native/Commands.h
@@ -18,6 +18,7 @@
 #include "common/Constants.h"
 
 #include "dawn_native/AttachmentState.h"
+#include "dawn_native/BindingInfo.h"
 #include "dawn_native/Texture.h"
 
 #include "dawn_native/dawn_platform.h"
@@ -210,7 +211,7 @@
     };
 
     struct SetBindGroupCmd {
-        uint32_t index;
+        BindGroupIndex index;
         Ref<BindGroupBase> group;
         uint32_t dynamicOffsetCount;
     };
diff --git a/src/dawn_native/Pipeline.cpp b/src/dawn_native/Pipeline.cpp
index 09771e7..df34441 100644
--- a/src/dawn_native/Pipeline.cpp
+++ b/src/dawn_native/Pipeline.cpp
@@ -85,11 +85,12 @@
         return {};
     }
 
-    BindGroupLayoutBase* PipelineBase::GetBindGroupLayout(uint32_t groupIndex) {
-        if (GetDevice()->ConsumedError(ValidateGetBindGroupLayout(groupIndex))) {
+    BindGroupLayoutBase* PipelineBase::GetBindGroupLayout(uint32_t groupIndexIn) {
+        if (GetDevice()->ConsumedError(ValidateGetBindGroupLayout(groupIndexIn))) {
             return BindGroupLayoutBase::MakeError(GetDevice());
         }
 
+        BindGroupIndex groupIndex(groupIndexIn);
         if (!mLayout->GetBindGroupLayoutsMask()[groupIndex]) {
             // Get or create an empty bind group layout.
             // TODO(enga): Consider caching this object on the Device and reusing it.
diff --git a/src/dawn_native/PipelineLayout.cpp b/src/dawn_native/PipelineLayout.cpp
index eedd6a9..def4875 100644
--- a/src/dawn_native/PipelineLayout.cpp
+++ b/src/dawn_native/PipelineLayout.cpp
@@ -97,8 +97,9 @@
                                            const PipelineLayoutDescriptor* descriptor)
         : CachedObject(device) {
         ASSERT(descriptor->bindGroupLayoutCount <= kMaxBindGroups);
-        for (uint32_t group = 0; group < descriptor->bindGroupLayoutCount; ++group) {
-            mBindGroupLayouts[group] = descriptor->bindGroupLayouts[group];
+        for (BindGroupIndex group(0); group < BindGroupIndex(descriptor->bindGroupLayoutCount);
+             ++group) {
+            mBindGroupLayouts[group] = descriptor->bindGroupLayouts[static_cast<uint32_t>(group)];
             mMask.set(group);
         }
     }
@@ -127,22 +128,24 @@
         ASSERT(count > 0);
 
         // Data which BindGroupLayoutDescriptor will point to for creation
-        std::array<ityp::array<BindingIndex, BindGroupLayoutEntry, kMaxBindingsPerGroup>,
-                   kMaxBindGroups>
+        ityp::array<BindGroupIndex,
+                    ityp::array<BindingIndex, BindGroupLayoutEntry, kMaxBindingsPerGroup>,
+                    kMaxBindGroups>
             entryData = {};
 
         // A map of bindings to the index in |entryData|
-        std::array<std::map<BindingNumber, BindingIndex>, kMaxBindGroups> usedBindingsMap = {};
+        ityp::array<BindGroupIndex, std::map<BindingNumber, BindingIndex>, kMaxBindGroups>
+            usedBindingsMap = {};
 
         // A counter of how many bindings we've populated in |entryData|
-        std::array<BindingIndex, kMaxBindGroups> entryCounts = {};
+        ityp::array<BindGroupIndex, BindingIndex, kMaxBindGroups> entryCounts = {};
 
-        uint32_t bindGroupLayoutCount = 0;
+        BindGroupIndex bindGroupLayoutCount(0);
         for (uint32_t moduleIndex = 0; moduleIndex < count; ++moduleIndex) {
             const ShaderModuleBase* module = modules[moduleIndex];
             const ShaderModuleBase::ModuleBindingInfo& info = module->GetBindingInfo();
 
-            for (uint32_t group = 0; group < info.size(); ++group) {
+            for (BindGroupIndex group(0); group < info.size(); ++group) {
                 for (const auto& it : info[group]) {
                     BindingNumber bindingNumber = it.first;
                     const ShaderModuleBase::ShaderBindingInfo& bindingInfo = it.second;
@@ -205,13 +208,14 @@
 
                     entryCounts[group]++;
 
-                    bindGroupLayoutCount = std::max(bindGroupLayoutCount, group + 1);
+                    bindGroupLayoutCount =
+                        std::max(bindGroupLayoutCount, group + BindGroupIndex(1));
                 }
             }
         }
 
-        std::array<BindGroupLayoutBase*, kMaxBindGroups> bindGroupLayouts = {};
-        for (uint32_t group = 0; group < bindGroupLayoutCount; ++group) {
+        ityp::array<BindGroupIndex, BindGroupLayoutBase*, kMaxBindGroups> bindGroupLayouts = {};
+        for (BindGroupIndex group(0); group < bindGroupLayoutCount; ++group) {
             BindGroupLayoutDescriptor desc = {};
             desc.entries = entryData[group].data();
             desc.entryCount = static_cast<uint32_t>(entryCounts[group]);
@@ -223,13 +227,13 @@
 
         PipelineLayoutDescriptor desc = {};
         desc.bindGroupLayouts = bindGroupLayouts.data();
-        desc.bindGroupLayoutCount = bindGroupLayoutCount;
+        desc.bindGroupLayoutCount = static_cast<uint32_t>(bindGroupLayoutCount);
         PipelineLayoutBase* pipelineLayout = device->CreatePipelineLayout(&desc);
         ASSERT(!pipelineLayout->IsError());
 
         // These bind group layouts are created internally and referenced by the pipeline layout.
         // Release the external refcount.
-        for (uint32_t group = 0; group < bindGroupLayoutCount; ++group) {
+        for (BindGroupIndex group(0); group < bindGroupLayoutCount; ++group) {
             if (bindGroupLayouts[group] != nullptr) {
                 bindGroupLayouts[group]->Release();
             }
@@ -244,50 +248,50 @@
         return pipelineLayout;
     }
 
-    const BindGroupLayoutBase* PipelineLayoutBase::GetBindGroupLayout(uint32_t group) const {
+    const BindGroupLayoutBase* PipelineLayoutBase::GetBindGroupLayout(BindGroupIndex group) const {
         ASSERT(!IsError());
-        ASSERT(group < kMaxBindGroups);
+        ASSERT(group < kMaxBindGroupsTyped);
         ASSERT(mMask[group]);
         const BindGroupLayoutBase* bgl = mBindGroupLayouts[group].Get();
         ASSERT(bgl != nullptr);
         return bgl;
     }
 
-    BindGroupLayoutBase* PipelineLayoutBase::GetBindGroupLayout(uint32_t group) {
+    BindGroupLayoutBase* PipelineLayoutBase::GetBindGroupLayout(BindGroupIndex group) {
         ASSERT(!IsError());
-        ASSERT(group < kMaxBindGroups);
+        ASSERT(group < kMaxBindGroupsTyped);
         ASSERT(mMask[group]);
         BindGroupLayoutBase* bgl = mBindGroupLayouts[group].Get();
         ASSERT(bgl != nullptr);
         return bgl;
     }
 
-    const std::bitset<kMaxBindGroups> PipelineLayoutBase::GetBindGroupLayoutsMask() const {
+    const BindGroupLayoutMask& PipelineLayoutBase::GetBindGroupLayoutsMask() const {
         ASSERT(!IsError());
         return mMask;
     }
 
-    std::bitset<kMaxBindGroups> PipelineLayoutBase::InheritedGroupsMask(
+    BindGroupLayoutMask PipelineLayoutBase::InheritedGroupsMask(
         const PipelineLayoutBase* other) const {
         ASSERT(!IsError());
-        return {(1 << GroupsInheritUpTo(other)) - 1u};
+        return {(1 << static_cast<uint32_t>(GroupsInheritUpTo(other))) - 1u};
     }
 
-    uint32_t PipelineLayoutBase::GroupsInheritUpTo(const PipelineLayoutBase* other) const {
+    BindGroupIndex PipelineLayoutBase::GroupsInheritUpTo(const PipelineLayoutBase* other) const {
         ASSERT(!IsError());
 
-        for (uint32_t i = 0; i < kMaxBindGroups; ++i) {
+        for (BindGroupIndex i(0); i < kMaxBindGroupsTyped; ++i) {
             if (!mMask[i] || mBindGroupLayouts[i].Get() != other->mBindGroupLayouts[i].Get()) {
                 return i;
             }
         }
-        return kMaxBindGroups;
+        return kMaxBindGroupsTyped;
     }
 
     size_t PipelineLayoutBase::HashFunc::operator()(const PipelineLayoutBase* pl) const {
         size_t hash = Hash(pl->mMask);
 
-        for (uint32_t group : IterateBitSet(pl->mMask)) {
+        for (BindGroupIndex group : IterateBitSet(pl->mMask)) {
             HashCombine(&hash, pl->GetBindGroupLayout(group));
         }
 
@@ -300,7 +304,7 @@
             return false;
         }
 
-        for (uint32_t group : IterateBitSet(a->mMask)) {
+        for (BindGroupIndex group : IterateBitSet(a->mMask)) {
             if (a->GetBindGroupLayout(group) != b->GetBindGroupLayout(group)) {
                 return false;
             }
diff --git a/src/dawn_native/PipelineLayout.h b/src/dawn_native/PipelineLayout.h
index f919eff..862caaf 100644
--- a/src/dawn_native/PipelineLayout.h
+++ b/src/dawn_native/PipelineLayout.h
@@ -16,6 +16,9 @@
 #define DAWNNATIVE_PIPELINELAYOUT_H_
 
 #include "common/Constants.h"
+#include "common/ityp_array.h"
+#include "common/ityp_bitset.h"
+#include "dawn_native/BindingInfo.h"
 #include "dawn_native/CachedObject.h"
 #include "dawn_native/Error.h"
 #include "dawn_native/Forward.h"
@@ -30,7 +33,9 @@
     MaybeError ValidatePipelineLayoutDescriptor(DeviceBase*,
                                                 const PipelineLayoutDescriptor* descriptor);
 
-    using BindGroupLayoutArray = std::array<Ref<BindGroupLayoutBase>, kMaxBindGroups>;
+    using BindGroupLayoutArray =
+        ityp::array<BindGroupIndex, Ref<BindGroupLayoutBase>, kMaxBindGroups>;
+    using BindGroupLayoutMask = ityp::bitset<BindGroupIndex, kMaxBindGroups>;
 
     class PipelineLayoutBase : public CachedObject {
       public:
@@ -41,17 +46,17 @@
         static ResultOrError<PipelineLayoutBase*>
         CreateDefault(DeviceBase* device, const ShaderModuleBase* const* modules, uint32_t count);
 
-        const BindGroupLayoutBase* GetBindGroupLayout(uint32_t group) const;
-        BindGroupLayoutBase* GetBindGroupLayout(uint32_t group);
-        const std::bitset<kMaxBindGroups> GetBindGroupLayoutsMask() const;
+        const BindGroupLayoutBase* GetBindGroupLayout(BindGroupIndex group) const;
+        BindGroupLayoutBase* GetBindGroupLayout(BindGroupIndex group);
+        const BindGroupLayoutMask& GetBindGroupLayoutsMask() const;
 
         // Utility functions to compute inherited bind groups.
         // Returns the inherited bind groups as a mask.
-        std::bitset<kMaxBindGroups> InheritedGroupsMask(const PipelineLayoutBase* other) const;
+        BindGroupLayoutMask InheritedGroupsMask(const PipelineLayoutBase* other) const;
 
         // Returns the index of the first incompatible bind group in the range
-        // [1, kMaxBindGroups + 1]
-        uint32_t GroupsInheritUpTo(const PipelineLayoutBase* other) const;
+        // [0, kMaxBindGroups]
+        BindGroupIndex GroupsInheritUpTo(const PipelineLayoutBase* other) const;
 
         // Functors necessary for the unordered_set<PipelineLayoutBase*>-based cache.
         struct HashFunc {
@@ -65,7 +70,7 @@
         PipelineLayoutBase(DeviceBase* device, ObjectBase::ErrorTag tag);
 
         BindGroupLayoutArray mBindGroupLayouts;
-        std::bitset<kMaxBindGroups> mMask;
+        BindGroupLayoutMask mMask;
     };
 
 }  // namespace dawn_native
diff --git a/src/dawn_native/ProgrammablePassEncoder.cpp b/src/dawn_native/ProgrammablePassEncoder.cpp
index 5ab4ccb..83aedad 100644
--- a/src/dawn_native/ProgrammablePassEncoder.cpp
+++ b/src/dawn_native/ProgrammablePassEncoder.cpp
@@ -130,15 +130,17 @@
         });
     }
 
-    void ProgrammablePassEncoder::SetBindGroup(uint32_t groupIndex,
+    void ProgrammablePassEncoder::SetBindGroup(uint32_t groupIndexIn,
                                                BindGroupBase* group,
                                                uint32_t dynamicOffsetCountIn,
                                                const uint32_t* dynamicOffsetsIn) {
         mEncodingContext->TryEncode(this, [&](CommandAllocator* allocator) -> MaybeError {
+            BindGroupIndex groupIndex(groupIndexIn);
+
             if (GetDevice()->IsValidationEnabled()) {
                 DAWN_TRY(GetDevice()->ValidateObject(group));
 
-                if (groupIndex >= kMaxBindGroups) {
+                if (groupIndex >= kMaxBindGroupsTyped) {
                     return DAWN_VALIDATION_ERROR("Setting bind group over the max");
                 }
 
diff --git a/src/dawn_native/RenderPipeline.cpp b/src/dawn_native/RenderPipeline.cpp
index df0a8d7..ceeb1eb 100644
--- a/src/dawn_native/RenderPipeline.cpp
+++ b/src/dawn_native/RenderPipeline.cpp
@@ -194,7 +194,7 @@
                     descriptor->fragmentStage->module->ComputeRequiredBufferSizesForLayout(
                         descriptor->layout);
 
-                for (uint32_t group = 0; group < bufferSizes.size(); ++group) {
+                for (BindGroupIndex group(0); group < bufferSizes.size(); ++group) {
                     ASSERT(bufferSizes[group].size() == fragmentSizes[group].size());
                     for (size_t i = 0; i < bufferSizes[group].size(); ++i) {
                         bufferSizes[group][i] =
diff --git a/src/dawn_native/ShaderModule.cpp b/src/dawn_native/ShaderModule.cpp
index dc79d75..04b6dfc 100644
--- a/src/dawn_native/ShaderModule.cpp
+++ b/src/dawn_native/ShaderModule.cpp
@@ -289,10 +289,10 @@
             }
         }
 
-        std::string GetShaderDeclarationString(size_t group, BindingNumber binding) {
+        std::string GetShaderDeclarationString(BindGroupIndex group, BindingNumber binding) {
             std::ostringstream ostream;
-            ostream << "the shader module declaration at set " << group << " binding "
-                    << static_cast<uint32_t>(binding);
+            ostream << "the shader module declaration at set " << static_cast<uint32_t>(group)
+                    << " binding " << static_cast<uint32_t>(binding);
             return ostream.str();
         }
     }  // anonymous namespace
@@ -515,12 +515,14 @@
         auto ExtractResourcesBinding =
             [this](std::vector<shaderc_spvc_binding_info> bindings) -> MaybeError {
             for (const auto& binding : bindings) {
-                if (binding.set >= kMaxBindGroups) {
+                BindGroupIndex bindGroupIndex(binding.set);
+
+                if (bindGroupIndex >= kMaxBindGroupsTyped) {
                     return DAWN_VALIDATION_ERROR("Bind group index over limits in the SPIRV");
                 }
 
-                const auto& it = mBindingInfo[binding.set].emplace(BindingNumber(binding.binding),
-                                                                   ShaderBindingInfo{});
+                const auto& it = mBindingInfo[bindGroupIndex].emplace(
+                    BindingNumber(binding.binding), ShaderBindingInfo{});
                 if (!it.second) {
                     return DAWN_VALIDATION_ERROR("Shader has duplicate bindings");
                 }
@@ -702,13 +704,15 @@
 
                 BindingNumber bindingNumber(
                     compiler.get_decoration(resource.id, spv::DecorationBinding));
-                uint32_t set = compiler.get_decoration(resource.id, spv::DecorationDescriptorSet);
+                BindGroupIndex bindGroupIndex(
+                    compiler.get_decoration(resource.id, spv::DecorationDescriptorSet));
 
-                if (set >= kMaxBindGroups) {
+                if (bindGroupIndex >= kMaxBindGroupsTyped) {
                     return DAWN_VALIDATION_ERROR("Bind group index over limits in the SPIRV");
                 }
 
-                const auto& it = mBindingInfo[set].emplace(bindingNumber, ShaderBindingInfo{});
+                const auto& it =
+                    mBindingInfo[bindGroupIndex].emplace(bindingNumber, ShaderBindingInfo{});
                 if (!it.second) {
                     return DAWN_VALIDATION_ERROR("Shader has duplicate bindings");
                 }
@@ -884,7 +888,7 @@
     RequiredBufferSizes ShaderModuleBase::ComputeRequiredBufferSizesForLayout(
         const PipelineLayoutBase* layout) const {
         RequiredBufferSizes bufferSizes;
-        for (uint32_t group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
             bufferSizes[group] =
                 GetBindGroupMinBufferSizes(mBindingInfo[group], layout->GetBindGroupLayout(group));
         }
@@ -926,16 +930,16 @@
         const PipelineLayoutBase* layout) const {
         ASSERT(!IsError());
 
-        for (uint32_t group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
             DAWN_TRY(
                 ValidateCompatibilityWithBindGroupLayout(group, layout->GetBindGroupLayout(group)));
         }
 
-        for (uint32_t group : IterateBitSet(~layout->GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(~layout->GetBindGroupLayoutsMask())) {
             if (mBindingInfo[group].size() > 0) {
                 std::ostringstream ostream;
-                ostream << "No bind group layout entry matches the declaration set " << group
-                        << " in the shader module";
+                ostream << "No bind group layout entry matches the declaration set "
+                        << static_cast<uint32_t>(group) << " in the shader module";
                 return DAWN_VALIDATION_ERROR(ostream.str());
             }
         }
@@ -944,7 +948,7 @@
     }
 
     MaybeError ShaderModuleBase::ValidateCompatibilityWithBindGroupLayout(
-        uint32_t group,
+        BindGroupIndex group,
         const BindGroupLayoutBase* layout) const {
         ASSERT(!IsError());
 
diff --git a/src/dawn_native/ShaderModule.h b/src/dawn_native/ShaderModule.h
index 9e68896..4e771aa 100644
--- a/src/dawn_native/ShaderModule.h
+++ b/src/dawn_native/ShaderModule.h
@@ -16,6 +16,7 @@
 #define DAWNNATIVE_SHADERMODULE_H_
 
 #include "common/Constants.h"
+#include "common/ityp_array.h"
 #include "dawn_native/BindingInfo.h"
 #include "dawn_native/CachedObject.h"
 #include "dawn_native/Error.h"
@@ -27,7 +28,6 @@
 
 #include "spvc/spvc.hpp"
 
-#include <array>
 #include <bitset>
 #include <map>
 #include <vector>
@@ -64,7 +64,7 @@
         };
 
         using BindingInfoMap = std::map<BindingNumber, ShaderBindingInfo>;
-        using ModuleBindingInfo = std::array<BindingInfoMap, kMaxBindGroups>;
+        using ModuleBindingInfo = ityp::array<BindGroupIndex, BindingInfoMap, kMaxBindGroups>;
 
         const ModuleBindingInfo& GetBindingInfo() const;
         const std::bitset<kMaxVertexAttributes>& GetUsedVertexAttributes() const;
@@ -102,7 +102,7 @@
         ShaderModuleBase(DeviceBase* device, ObjectBase::ErrorTag tag);
 
         MaybeError ValidateCompatibilityWithBindGroupLayout(
-            uint32_t group,
+            BindGroupIndex group,
             const BindGroupLayoutBase* layout) const;
 
         std::vector<uint64_t> GetBindGroupMinBufferSizes(const BindingInfoMap& shaderMap,
diff --git a/src/dawn_native/d3d12/CommandBufferD3D12.cpp b/src/dawn_native/d3d12/CommandBufferD3D12.cpp
index 8af7a9e..b4d171a 100644
--- a/src/dawn_native/d3d12/CommandBufferD3D12.cpp
+++ b/src/dawn_native/d3d12/CommandBufferD3D12.cpp
@@ -106,7 +106,7 @@
             // TODO(bryan.bernhart@intel.com): Consider further optimization.
             bool didCreateBindGroupViews = true;
             bool didCreateBindGroupSamplers = true;
-            for (uint32_t index : IterateBitSet(mDirtyBindGroups)) {
+            for (BindGroupIndex index : IterateBitSet(mDirtyBindGroups)) {
                 BindGroup* group = ToBackend(mBindGroups[index]);
                 didCreateBindGroupViews = group->PopulateViews(mViewAllocator);
                 didCreateBindGroupSamplers = group->PopulateSamplers(mDevice, mSamplerAllocator);
@@ -132,7 +132,7 @@
                 // Must be called before applying the bindgroups.
                 SetID3D12DescriptorHeaps(commandList);
 
-                for (uint32_t index : IterateBitSet(mBindGroupLayoutsMask)) {
+                for (BindGroupIndex index : IterateBitSet(mBindGroupLayoutsMask)) {
                     BindGroup* group = ToBackend(mBindGroups[index]);
                     didCreateBindGroupViews = group->PopulateViews(mViewAllocator);
                     didCreateBindGroupSamplers =
@@ -142,14 +142,14 @@
                 }
             }
 
-            for (uint32_t index : IterateBitSet(mDirtyBindGroupsObjectChangedOrIsDynamic)) {
+            for (BindGroupIndex index : IterateBitSet(mDirtyBindGroupsObjectChangedOrIsDynamic)) {
                 BindGroup* group = ToBackend(mBindGroups[index]);
                 ApplyBindGroup(commandList, ToBackend(mPipelineLayout), index, group,
                                mDynamicOffsetCounts[index], mDynamicOffsets[index].data());
             }
 
             if (mInCompute) {
-                for (uint32_t index : IterateBitSet(mBindGroupLayoutsMask)) {
+                for (BindGroupIndex index : IterateBitSet(mBindGroupLayoutsMask)) {
                     for (BindingIndex binding : IterateBitSet(mBindingsNeedingBarrier[index])) {
                         wgpu::BindingType bindingType = mBindingTypes[index][binding];
                         switch (bindingType) {
@@ -211,7 +211,7 @@
       private:
         void ApplyBindGroup(ID3D12GraphicsCommandList* commandList,
                             const PipelineLayout* pipelineLayout,
-                            uint32_t index,
+                            BindGroupIndex index,
                             BindGroup* group,
                             uint32_t dynamicOffsetCountIn,
                             const uint64_t* dynamicOffsetsIn) {
diff --git a/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp b/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp
index 456466b..5e3460b 100644
--- a/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp
+++ b/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp
@@ -79,7 +79,7 @@
 
         uint32_t rangeIndex = 0;
 
-        for (uint32_t group : IterateBitSet(GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(GetBindGroupLayoutsMask())) {
             const BindGroupLayout* bindGroupLayout = ToBackend(GetBindGroupLayout(group));
 
             // Set the root descriptor table parameter and copy ranges. Ranges are offset by the
@@ -99,7 +99,7 @@
 
                 for (uint32_t i = 0; i < rangeCount; ++i) {
                     ranges[rangeIndex] = descriptorRanges[i];
-                    ranges[rangeIndex].RegisterSpace = group;
+                    ranges[rangeIndex].RegisterSpace = static_cast<uint32_t>(group);
                     rangeIndex++;
                 }
 
@@ -140,7 +140,7 @@
                 // Setup root descriptor.
                 D3D12_ROOT_DESCRIPTOR rootDescriptor;
                 rootDescriptor.ShaderRegister = shaderRegisters[dynamicBindingIndex];
-                rootDescriptor.RegisterSpace = group;
+                rootDescriptor.RegisterSpace = static_cast<uint32_t>(group);
 
                 // Set root descriptors in root signatures.
                 rootParameter.Descriptor = rootDescriptor;
@@ -177,13 +177,13 @@
         return {};
     }
 
-    uint32_t PipelineLayout::GetCbvUavSrvRootParameterIndex(uint32_t group) const {
-        ASSERT(group < kMaxBindGroups);
+    uint32_t PipelineLayout::GetCbvUavSrvRootParameterIndex(BindGroupIndex group) const {
+        ASSERT(group < kMaxBindGroupsTyped);
         return mCbvUavSrvRootParameterInfo[group];
     }
 
-    uint32_t PipelineLayout::GetSamplerRootParameterIndex(uint32_t group) const {
-        ASSERT(group < kMaxBindGroups);
+    uint32_t PipelineLayout::GetSamplerRootParameterIndex(BindGroupIndex group) const {
+        ASSERT(group < kMaxBindGroupsTyped);
         return mSamplerRootParameterInfo[group];
     }
 
@@ -191,9 +191,9 @@
         return mRootSignature.Get();
     }
 
-    uint32_t PipelineLayout::GetDynamicRootParameterIndex(uint32_t group,
+    uint32_t PipelineLayout::GetDynamicRootParameterIndex(BindGroupIndex group,
                                                           BindingIndex bindingIndex) const {
-        ASSERT(group < kMaxBindGroups);
+        ASSERT(group < kMaxBindGroupsTyped);
         ASSERT(bindingIndex < kMaxBindingsPerGroupTyped);
         ASSERT(GetBindGroupLayout(group)->GetBindingInfo(bindingIndex).hasDynamicOffset);
         ASSERT(GetBindGroupLayout(group)->GetBindingInfo(bindingIndex).visibility !=
diff --git a/src/dawn_native/d3d12/PipelineLayoutD3D12.h b/src/dawn_native/d3d12/PipelineLayoutD3D12.h
index a539174..8de41b0 100644
--- a/src/dawn_native/d3d12/PipelineLayoutD3D12.h
+++ b/src/dawn_native/d3d12/PipelineLayoutD3D12.h
@@ -29,11 +29,12 @@
         static ResultOrError<PipelineLayout*> Create(Device* device,
                                                      const PipelineLayoutDescriptor* descriptor);
 
-        uint32_t GetCbvUavSrvRootParameterIndex(uint32_t group) const;
-        uint32_t GetSamplerRootParameterIndex(uint32_t group) const;
+        uint32_t GetCbvUavSrvRootParameterIndex(BindGroupIndex group) const;
+        uint32_t GetSamplerRootParameterIndex(BindGroupIndex group) const;
 
         // Returns the index of the root parameter reserved for a dynamic buffer binding
-        uint32_t GetDynamicRootParameterIndex(uint32_t group, BindingIndex bindingIndex) const;
+        uint32_t GetDynamicRootParameterIndex(BindGroupIndex group,
+                                              BindingIndex bindingIndex) const;
 
         ID3D12RootSignature* GetRootSignature() const;
 
@@ -41,9 +42,11 @@
         ~PipelineLayout() override = default;
         using PipelineLayoutBase::PipelineLayoutBase;
         MaybeError Initialize();
-        std::array<uint32_t, kMaxBindGroups> mCbvUavSrvRootParameterInfo;
-        std::array<uint32_t, kMaxBindGroups> mSamplerRootParameterInfo;
-        std::array<ityp::array<BindingIndex, uint32_t, kMaxBindingsPerGroup>, kMaxBindGroups>
+        ityp::array<BindGroupIndex, uint32_t, kMaxBindGroups> mCbvUavSrvRootParameterInfo;
+        ityp::array<BindGroupIndex, uint32_t, kMaxBindGroups> mSamplerRootParameterInfo;
+        ityp::array<BindGroupIndex,
+                    ityp::array<BindingIndex, uint32_t, kMaxBindingsPerGroup>,
+                    kMaxBindGroups>
             mDynamicRootParameterIndices;
         ComPtr<ID3D12RootSignature> mRootSignature;
     };
diff --git a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
index 8901d20..770a006 100644
--- a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
@@ -155,7 +155,7 @@
         }
 
         const ModuleBindingInfo& moduleBindingInfo = GetBindingInfo();
-        for (uint32_t group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
             const BindGroupLayout* bgl = ToBackend(layout->GetBindGroupLayout(group));
             const auto& bindingOffsets = bgl->GetBindingOffsets();
             const auto& groupBindingInfo = moduleBindingInfo[group];
@@ -181,14 +181,14 @@
                     if (forceStorageBufferAsUAV) {
                         DAWN_TRY(CheckSpvcSuccess(
                             mSpvcContext.SetHLSLForceStorageBufferAsUAV(
-                                group, static_cast<uint32_t>(bindingNumber)),
+                                static_cast<uint32_t>(group), static_cast<uint32_t>(bindingNumber)),
                             "Unable to force read-only storage buffer as UAV w/ spvc"));
                     }
                 } else {
                     compiler->set_decoration(bindingInfo.id, spv::DecorationBinding, bindingOffset);
                     if (forceStorageBufferAsUAV) {
                         compiler->set_hlsl_force_storage_buffer_as_uav(
-                            group, static_cast<uint32_t>(bindingNumber));
+                            static_cast<uint32_t>(group), static_cast<uint32_t>(bindingNumber));
                     }
                 }
             }
diff --git a/src/dawn_native/metal/CommandBufferMTL.mm b/src/dawn_native/metal/CommandBufferMTL.mm
index c52b63d..2bbcdd4 100644
--- a/src/dawn_native/metal/CommandBufferMTL.mm
+++ b/src/dawn_native/metal/CommandBufferMTL.mm
@@ -471,7 +471,8 @@
 
             template <typename Encoder>
             void Apply(Encoder encoder) {
-                for (uint32_t index : IterateBitSet(mDirtyBindGroupsObjectChangedOrIsDynamic)) {
+                for (BindGroupIndex index :
+                     IterateBitSet(mDirtyBindGroupsObjectChangedOrIsDynamic)) {
                     ApplyBindGroup(encoder, index, ToBackend(mBindGroups[index]),
                                    mDynamicOffsetCounts[index], mDynamicOffsets[index].data(),
                                    ToBackend(mPipelineLayout));
@@ -486,7 +487,7 @@
             // two encoder types.
             void ApplyBindGroupImpl(id<MTLRenderCommandEncoder> render,
                                     id<MTLComputeCommandEncoder> compute,
-                                    uint32_t index,
+                                    BindGroupIndex index,
                                     BindGroup* group,
                                     uint32_t dynamicOffsetCount,
                                     uint64_t* dynamicOffsets,
diff --git a/src/dawn_native/metal/PipelineLayoutMTL.h b/src/dawn_native/metal/PipelineLayoutMTL.h
index e5acafc..7a3ad80 100644
--- a/src/dawn_native/metal/PipelineLayoutMTL.h
+++ b/src/dawn_native/metal/PipelineLayoutMTL.h
@@ -43,7 +43,9 @@
         PipelineLayout(Device* device, const PipelineLayoutDescriptor* descriptor);
 
         using BindingIndexInfo =
-            std::array<ityp::array<BindingIndex, uint32_t, kMaxBindingsPerGroup>, kMaxBindGroups>;
+            ityp::array<BindGroupIndex,
+                        ityp::array<BindingIndex, uint32_t, kMaxBindingsPerGroup>,
+                        kMaxBindGroups>;
         const BindingIndexInfo& GetBindingIndexInfo(SingleShaderStage stage) const;
 
         // The number of Metal vertex stage buffers used for the whole pipeline layout.
diff --git a/src/dawn_native/metal/PipelineLayoutMTL.mm b/src/dawn_native/metal/PipelineLayoutMTL.mm
index 996bd1f..3ee7d92 100644
--- a/src/dawn_native/metal/PipelineLayoutMTL.mm
+++ b/src/dawn_native/metal/PipelineLayoutMTL.mm
@@ -28,7 +28,7 @@
             uint32_t samplerIndex = 0;
             uint32_t textureIndex = 0;
 
-            for (uint32_t group : IterateBitSet(GetBindGroupLayoutsMask())) {
+            for (BindGroupIndex group : IterateBitSet(GetBindGroupLayoutsMask())) {
                 for (BindingIndex bindingIndex{0};
                      bindingIndex < GetBindGroupLayout(group)->GetBindingCount(); ++bindingIndex) {
                     const BindingInfo& bindingInfo =
diff --git a/src/dawn_native/metal/ShaderModuleMTL.mm b/src/dawn_native/metal/ShaderModuleMTL.mm
index be45558..5983e9e 100644
--- a/src/dawn_native/metal/ShaderModuleMTL.mm
+++ b/src/dawn_native/metal/ShaderModuleMTL.mm
@@ -132,7 +132,7 @@
         // a table of MSLResourceBinding to give to SPIRV-Cross.
 
         // Create one resource binding entry per stage per binding.
-        for (uint32_t group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
             const BindGroupLayoutBase::BindingMap& bindingMap =
                 layout->GetBindGroupLayout(group)->GetBindingMap();
 
@@ -148,7 +148,7 @@
                     if (GetDevice()->IsToggleEnabled(Toggle::UseSpvc)) {
                         shaderc_spvc_msl_resource_binding mslBinding;
                         mslBinding.stage = ToSpvcExecutionModel(stage);
-                        mslBinding.desc_set = group;
+                        mslBinding.desc_set = static_cast<uint32_t>(group);
                         mslBinding.binding = static_cast<uint32_t>(bindingNumber);
                         mslBinding.msl_buffer = mslBinding.msl_texture = mslBinding.msl_sampler =
                             shaderIndex;
@@ -157,7 +157,7 @@
                     } else {
                         spirv_cross::MSLResourceBinding mslBinding;
                         mslBinding.stage = SpirvExecutionModelForStage(stage);
-                        mslBinding.desc_set = group;
+                        mslBinding.desc_set = static_cast<uint32_t>(group);
                         mslBinding.binding = static_cast<uint32_t>(bindingNumber);
                         mslBinding.msl_buffer = mslBinding.msl_texture = mslBinding.msl_sampler =
                             shaderIndex;
diff --git a/src/dawn_native/opengl/CommandBufferGL.cpp b/src/dawn_native/opengl/CommandBufferGL.cpp
index 88148e9..fec3705 100644
--- a/src/dawn_native/opengl/CommandBufferGL.cpp
+++ b/src/dawn_native/opengl/CommandBufferGL.cpp
@@ -225,7 +225,8 @@
             }
 
             void Apply(const OpenGLFunctions& gl) {
-                for (uint32_t index : IterateBitSet(mDirtyBindGroupsObjectChangedOrIsDynamic)) {
+                for (BindGroupIndex index :
+                     IterateBitSet(mDirtyBindGroupsObjectChangedOrIsDynamic)) {
                     ApplyBindGroup(gl, index, mBindGroups[index], mDynamicOffsetCounts[index],
                                    mDynamicOffsets[index].data());
                 }
@@ -234,7 +235,7 @@
 
           private:
             void ApplyBindGroup(const OpenGLFunctions& gl,
-                                uint32_t index,
+                                BindGroupIndex index,
                                 BindGroupBase* group,
                                 uint32_t dynamicOffsetCount,
                                 uint64_t* dynamicOffsets) {
diff --git a/src/dawn_native/opengl/PipelineGL.cpp b/src/dawn_native/opengl/PipelineGL.cpp
index 41728be..76c24b6 100644
--- a/src/dawn_native/opengl/PipelineGL.cpp
+++ b/src/dawn_native/opengl/PipelineGL.cpp
@@ -106,7 +106,7 @@
         // etc.
         const auto& indices = layout->GetBindingIndexInfo();
 
-        for (uint32_t group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
             const BindGroupLayoutBase* bgl = layout->GetBindGroupLayout(group);
 
             for (const auto& it : bgl->GetBindingMap()) {
diff --git a/src/dawn_native/opengl/PipelineGL.h b/src/dawn_native/opengl/PipelineGL.h
index 6a081c1..7f681f9 100644
--- a/src/dawn_native/opengl/PipelineGL.h
+++ b/src/dawn_native/opengl/PipelineGL.h
@@ -36,8 +36,9 @@
                         const PipelineLayout* layout,
                         const PerStage<const ShaderModule*>& modules);
 
-        using BindingLocations =
-            std::array<std::array<GLint, kMaxBindingsPerGroup>, kMaxBindGroups>;
+        using BindingLocations = ityp::array<BindGroupIndex,
+                                             ityp::array<BindingIndex, GLint, kMaxBindingsPerGroup>,
+                                             kMaxBindGroups>;
 
         // For each unit a sampler is bound to we need to know if we should use filtering or not
         // because int and uint texture are only complete without filtering.
diff --git a/src/dawn_native/opengl/PipelineLayoutGL.cpp b/src/dawn_native/opengl/PipelineLayoutGL.cpp
index ed4cceb..0e98c61 100644
--- a/src/dawn_native/opengl/PipelineLayoutGL.cpp
+++ b/src/dawn_native/opengl/PipelineLayoutGL.cpp
@@ -28,7 +28,7 @@
         GLuint ssboIndex = 0;
         GLuint storageTextureIndex = 0;
 
-        for (uint32_t group : IterateBitSet(GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex group : IterateBitSet(GetBindGroupLayoutsMask())) {
             const BindGroupLayoutBase* bgl = GetBindGroupLayout(group);
 
             for (BindingIndex bindingIndex{0}; bindingIndex < bgl->GetBindingCount();
diff --git a/src/dawn_native/opengl/PipelineLayoutGL.h b/src/dawn_native/opengl/PipelineLayoutGL.h
index 85458d2..3d511d6 100644
--- a/src/dawn_native/opengl/PipelineLayoutGL.h
+++ b/src/dawn_native/opengl/PipelineLayoutGL.h
@@ -30,7 +30,9 @@
         PipelineLayout(Device* device, const PipelineLayoutDescriptor* descriptor);
 
         using BindingIndexInfo =
-            std::array<ityp::array<BindingIndex, GLuint, kMaxBindingsPerGroup>, kMaxBindGroups>;
+            ityp::array<BindGroupIndex,
+                        ityp::array<BindingIndex, GLuint, kMaxBindingsPerGroup>,
+                        kMaxBindGroups>;
         const BindingIndexInfo& GetBindingIndexInfo() const;
 
         GLuint GetTextureUnitsUsed() const;
diff --git a/src/dawn_native/opengl/ShaderModuleGL.cpp b/src/dawn_native/opengl/ShaderModuleGL.cpp
index d7d9f74..53aa101 100644
--- a/src/dawn_native/opengl/ShaderModuleGL.cpp
+++ b/src/dawn_native/opengl/ShaderModuleGL.cpp
@@ -24,9 +24,10 @@
 
 namespace dawn_native { namespace opengl {
 
-    std::string GetBindingName(uint32_t group, BindingNumber bindingNumber) {
+    std::string GetBindingName(BindGroupIndex group, BindingNumber bindingNumber) {
         std::ostringstream o;
-        o << "dawn_binding_" << group << "_" << static_cast<uint32_t>(bindingNumber);
+        o << "dawn_binding_" << static_cast<uint32_t>(group) << "_"
+          << static_cast<uint32_t>(bindingNumber);
         return o.str();
     }
 
@@ -42,8 +43,9 @@
     std::string CombinedSampler::GetName() const {
         std::ostringstream o;
         o << "dawn_combined";
-        o << "_" << samplerLocation.group << "_" << static_cast<uint32_t>(samplerLocation.binding);
-        o << "_with_" << textureLocation.group << "_"
+        o << "_" << static_cast<uint32_t>(samplerLocation.group) << "_"
+          << static_cast<uint32_t>(samplerLocation.binding);
+        o << "_with_" << static_cast<uint32_t>(textureLocation.group) << "_"
           << static_cast<uint32_t>(textureLocation.binding);
         return o.str();
     }
@@ -141,16 +143,20 @@
                 mCombinedInfo.emplace_back();
                 auto& info = mCombinedInfo.back();
 
+                uint32_t samplerGroup;
                 mSpvcContext.GetDecoration(sampler.sampler_id,
-                                           shaderc_spvc_decoration_descriptorset,
-                                           &info.samplerLocation.group);
+                                           shaderc_spvc_decoration_descriptorset, &samplerGroup);
+                info.samplerLocation.group = BindGroupIndex(samplerGroup);
+
                 uint32_t samplerBinding;
                 mSpvcContext.GetDecoration(sampler.sampler_id, shaderc_spvc_decoration_binding,
                                            &samplerBinding);
                 info.samplerLocation.binding = BindingNumber(samplerBinding);
 
+                uint32_t textureGroup;
                 mSpvcContext.GetDecoration(sampler.image_id, shaderc_spvc_decoration_descriptorset,
-                                           &info.textureLocation.group);
+                                           &textureGroup);
+                info.textureLocation.group = BindGroupIndex(textureGroup);
 
                 uint32_t textureBinding;
                 mSpvcContext.GetDecoration(sampler.image_id, shaderc_spvc_decoration_binding,
@@ -164,12 +170,12 @@
                 mCombinedInfo.emplace_back();
 
                 auto& info = mCombinedInfo.back();
-                info.samplerLocation.group =
-                    compiler->get_decoration(combined.sampler_id, spv::DecorationDescriptorSet);
+                info.samplerLocation.group = BindGroupIndex(
+                    compiler->get_decoration(combined.sampler_id, spv::DecorationDescriptorSet));
                 info.samplerLocation.binding = BindingNumber(
                     compiler->get_decoration(combined.sampler_id, spv::DecorationBinding));
-                info.textureLocation.group =
-                    compiler->get_decoration(combined.image_id, spv::DecorationDescriptorSet);
+                info.textureLocation.group = BindGroupIndex(
+                    compiler->get_decoration(combined.image_id, spv::DecorationDescriptorSet));
                 info.textureLocation.binding = BindingNumber(
                     compiler->get_decoration(combined.image_id, spv::DecorationBinding));
                 compiler->set_name(combined.combined_id, info.GetName());
@@ -179,7 +185,7 @@
         // Change binding names to be "dawn_binding_<group>_<binding>".
         // Also unsets the SPIRV "Binding" decoration as it outputs "layout(binding=)" which
         // isn't supported on OSX's OpenGL.
-        for (uint32_t group = 0; group < kMaxBindGroups; ++group) {
+        for (BindGroupIndex group(0); group < kMaxBindGroupsTyped; ++group) {
             for (const auto& it : bindingInfo[group]) {
                 BindingNumber bindingNumber = it.first;
                 const auto& info = it.second;
diff --git a/src/dawn_native/opengl/ShaderModuleGL.h b/src/dawn_native/opengl/ShaderModuleGL.h
index dc267d2..3ddd19a 100644
--- a/src/dawn_native/opengl/ShaderModuleGL.h
+++ b/src/dawn_native/opengl/ShaderModuleGL.h
@@ -23,10 +23,10 @@
 
     class Device;
 
-    std::string GetBindingName(uint32_t group, BindingNumber bindingNumber);
+    std::string GetBindingName(BindGroupIndex group, BindingNumber bindingNumber);
 
     struct BindingLocation {
-        uint32_t group;
+        BindGroupIndex group;
         BindingNumber binding;
     };
     bool operator<(const BindingLocation& a, const BindingLocation& b);
diff --git a/src/dawn_native/vulkan/CommandBufferVk.cpp b/src/dawn_native/vulkan/CommandBufferVk.cpp
index dca8842..e32725a 100644
--- a/src/dawn_native/vulkan/CommandBufferVk.cpp
+++ b/src/dawn_native/vulkan/CommandBufferVk.cpp
@@ -95,23 +95,25 @@
             return region;
         }
 
-        void ApplyDescriptorSets(Device* device,
-                                 VkCommandBuffer commands,
-                                 VkPipelineBindPoint bindPoint,
-                                 VkPipelineLayout pipelineLayout,
-                                 const std::bitset<kMaxBindGroups>& bindGroupsToApply,
-                                 const std::array<BindGroupBase*, kMaxBindGroups>& bindGroups,
-                                 const std::array<uint32_t, kMaxBindGroups>& dynamicOffsetCounts,
-                                 const std::array<std::array<uint32_t, kMaxBindingsPerGroup>,
-                                                  kMaxBindGroups>& dynamicOffsets) {
-            for (uint32_t dirtyIndex : IterateBitSet(bindGroupsToApply)) {
+        void ApplyDescriptorSets(
+            Device* device,
+            VkCommandBuffer commands,
+            VkPipelineBindPoint bindPoint,
+            VkPipelineLayout pipelineLayout,
+            const BindGroupLayoutMask& bindGroupsToApply,
+            const ityp::array<BindGroupIndex, BindGroupBase*, kMaxBindGroups>& bindGroups,
+            const ityp::array<BindGroupIndex, uint32_t, kMaxBindGroups>& dynamicOffsetCounts,
+            const ityp::array<BindGroupIndex,
+                              std::array<uint32_t, kMaxBindingsPerGroup>,
+                              kMaxBindGroups>& dynamicOffsets) {
+            for (BindGroupIndex dirtyIndex : IterateBitSet(bindGroupsToApply)) {
                 VkDescriptorSet set = ToBackend(bindGroups[dirtyIndex])->GetHandle();
                 const uint32_t* dynamicOffset = dynamicOffsetCounts[dirtyIndex] > 0
                                                     ? dynamicOffsets[dirtyIndex].data()
                                                     : nullptr;
-                device->fn.CmdBindDescriptorSets(commands, bindPoint, pipelineLayout, dirtyIndex, 1,
-                                                 &*set, dynamicOffsetCounts[dirtyIndex],
-                                                 dynamicOffset);
+                device->fn.CmdBindDescriptorSets(commands, bindPoint, pipelineLayout,
+                                                 static_cast<uint32_t>(dirtyIndex), 1, &*set,
+                                                 dynamicOffsetCounts[dirtyIndex], dynamicOffset);
             }
         }
 
@@ -143,7 +145,7 @@
                                     mDirtyBindGroupsObjectChangedOrIsDynamic, mBindGroups,
                                     mDynamicOffsetCounts, mDynamicOffsets);
 
-                for (uint32_t index : IterateBitSet(mBindGroupLayoutsMask)) {
+                for (BindGroupIndex index : IterateBitSet(mBindGroupLayoutsMask)) {
                     for (BindingIndex bindingIndex :
                          IterateBitSet(mBindingsNeedingBarrier[index])) {
                         switch (mBindingTypes[index][bindingIndex]) {
diff --git a/src/dawn_native/vulkan/PipelineLayoutVk.cpp b/src/dawn_native/vulkan/PipelineLayoutVk.cpp
index 87b47ab..80c5aa8 100644
--- a/src/dawn_native/vulkan/PipelineLayoutVk.cpp
+++ b/src/dawn_native/vulkan/PipelineLayoutVk.cpp
@@ -37,7 +37,7 @@
         // this constraints at the Dawn level?
         uint32_t numSetLayouts = 0;
         std::array<VkDescriptorSetLayout, kMaxBindGroups> setLayouts;
-        for (uint32_t setIndex : IterateBitSet(GetBindGroupLayoutsMask())) {
+        for (BindGroupIndex setIndex : IterateBitSet(GetBindGroupLayoutsMask())) {
             setLayouts[numSetLayouts] = ToBackend(GetBindGroupLayout(setIndex))->GetHandle();
             numSetLayouts++;
         }