diff --git a/src/dawn_native/CommandBuffer.cpp b/src/dawn_native/CommandBuffer.cpp
index 18fef0d..a964f40 100644
--- a/src/dawn_native/CommandBuffer.cpp
+++ b/src/dawn_native/CommandBuffer.cpp
@@ -73,6 +73,10 @@
         return mResourceUsages;
     }
 
+    CommandIterator* CommandBufferBase::GetCommandIteratorForTesting() {
+        return &mCommands;
+    }
+
     bool IsCompleteSubresourceCopiedTo(const TextureBase* texture,
                                        const Extent3D copySize,
                                        const uint32_t mipLevel) {
diff --git a/src/dawn_native/CommandBuffer.h b/src/dawn_native/CommandBuffer.h
index 2800929..c6d47ae 100644
--- a/src/dawn_native/CommandBuffer.h
+++ b/src/dawn_native/CommandBuffer.h
@@ -43,6 +43,8 @@
 
         const CommandBufferResourceUsage& GetResourceUsages() const;
 
+        CommandIterator* GetCommandIteratorForTesting();
+
       protected:
         ~CommandBufferBase() override;
 
diff --git a/src/dawn_native/CommandBufferStateTracker.cpp b/src/dawn_native/CommandBufferStateTracker.cpp
index 5210936..45892a6 100644
--- a/src/dawn_native/CommandBufferStateTracker.cpp
+++ b/src/dawn_native/CommandBufferStateTracker.cpp
@@ -17,8 +17,10 @@
 #include "common/Assert.h"
 #include "common/BitSetIterator.h"
 #include "dawn_native/BindGroup.h"
+#include "dawn_native/ComputePassEncoder.h"
 #include "dawn_native/ComputePipeline.h"
 #include "dawn_native/Forward.h"
+#include "dawn_native/ObjectType_autogen.h"
 #include "dawn_native/PipelineLayout.h"
 #include "dawn_native/RenderPipeline.h"
 
@@ -83,13 +85,15 @@
     MaybeError CommandBufferStateTracker::ValidateBufferInRangeForVertexBuffer(
         uint32_t vertexCount,
         uint32_t firstVertex) {
+        RenderPipelineBase* lastRenderPipeline = GetRenderPipeline();
+
         const ityp::bitset<VertexBufferSlot, kMaxVertexBuffers>&
             vertexBufferSlotsUsedAsVertexBuffer =
-                mLastRenderPipeline->GetVertexBufferSlotsUsedAsVertexBuffer();
+                lastRenderPipeline->GetVertexBufferSlotsUsedAsVertexBuffer();
 
         for (auto usedSlotVertex : IterateBitSet(vertexBufferSlotsUsedAsVertexBuffer)) {
             const VertexBufferInfo& vertexBuffer =
-                mLastRenderPipeline->GetVertexBuffer(usedSlotVertex);
+                lastRenderPipeline->GetVertexBuffer(usedSlotVertex);
             uint64_t arrayStride = vertexBuffer.arrayStride;
             uint64_t bufferSize = mVertexBufferSizes[usedSlotVertex];
 
@@ -120,13 +124,15 @@
     MaybeError CommandBufferStateTracker::ValidateBufferInRangeForInstanceBuffer(
         uint32_t instanceCount,
         uint32_t firstInstance) {
+        RenderPipelineBase* lastRenderPipeline = GetRenderPipeline();
+
         const ityp::bitset<VertexBufferSlot, kMaxVertexBuffers>&
             vertexBufferSlotsUsedAsInstanceBuffer =
-                mLastRenderPipeline->GetVertexBufferSlotsUsedAsInstanceBuffer();
+                lastRenderPipeline->GetVertexBufferSlotsUsedAsInstanceBuffer();
 
         for (auto usedSlotInstance : IterateBitSet(vertexBufferSlotsUsedAsInstanceBuffer)) {
             const VertexBufferInfo& vertexBuffer =
-                mLastRenderPipeline->GetVertexBuffer(usedSlotInstance);
+                lastRenderPipeline->GetVertexBuffer(usedSlotInstance);
             uint64_t arrayStride = vertexBuffer.arrayStride;
             uint64_t bufferSize = mVertexBufferSizes[usedSlotInstance];
             if (arrayStride == 0) {
@@ -209,18 +215,19 @@
         }
 
         if (aspects[VALIDATION_ASPECT_VERTEX_BUFFERS]) {
-            ASSERT(mLastRenderPipeline != nullptr);
+            RenderPipelineBase* lastRenderPipeline = GetRenderPipeline();
 
             const ityp::bitset<VertexBufferSlot, kMaxVertexBuffers>& requiredVertexBuffers =
-                mLastRenderPipeline->GetVertexBufferSlotsUsed();
+                lastRenderPipeline->GetVertexBufferSlotsUsed();
             if (IsSubset(requiredVertexBuffers, mVertexBufferSlotsUsed)) {
                 mAspects.set(VALIDATION_ASPECT_VERTEX_BUFFERS);
             }
         }
 
         if (aspects[VALIDATION_ASPECT_INDEX_BUFFER] && mIndexBufferSet) {
-            if (!IsStripPrimitiveTopology(mLastRenderPipeline->GetPrimitiveTopology()) ||
-                mIndexFormat == mLastRenderPipeline->GetStripIndexFormat()) {
+            RenderPipelineBase* lastRenderPipeline = GetRenderPipeline();
+            if (!IsStripPrimitiveTopology(lastRenderPipeline->GetPrimitiveTopology()) ||
+                mIndexFormat == lastRenderPipeline->GetStripIndexFormat()) {
                 mAspects.set(VALIDATION_ASPECT_INDEX_BUFFER);
             }
         }
@@ -234,12 +241,13 @@
         if (DAWN_UNLIKELY(aspects[VALIDATION_ASPECT_INDEX_BUFFER])) {
             DAWN_INVALID_IF(!mIndexBufferSet, "Index buffer was not set.");
 
-            wgpu::IndexFormat pipelineIndexFormat = mLastRenderPipeline->GetStripIndexFormat();
+            RenderPipelineBase* lastRenderPipeline = GetRenderPipeline();
+            wgpu::IndexFormat pipelineIndexFormat = lastRenderPipeline->GetStripIndexFormat();
             DAWN_INVALID_IF(
-                IsStripPrimitiveTopology(mLastRenderPipeline->GetPrimitiveTopology()) &&
+                IsStripPrimitiveTopology(lastRenderPipeline->GetPrimitiveTopology()) &&
                     mIndexFormat != pipelineIndexFormat,
                 "Strip index format (%s) of %s does not match index buffer format (%s).",
-                pipelineIndexFormat, mLastRenderPipeline, mIndexFormat);
+                pipelineIndexFormat, lastRenderPipeline, mIndexFormat);
 
             // The chunk of code above should be similar to the one in |RecomputeLazyAspects|.
             // It returns the first invalid state found. We shouldn't be able to reach this line
@@ -251,7 +259,7 @@
 
         // TODO(dawn:563): Indicate which slots were not set.
         DAWN_INVALID_IF(aspects[VALIDATION_ASPECT_VERTEX_BUFFERS],
-                        "Vertex buffer slots required by %s were not set.", mLastRenderPipeline);
+                        "Vertex buffer slots required by %s were not set.", GetRenderPipeline());
 
         if (DAWN_UNLIKELY(aspects[VALIDATION_ASPECT_BIND_GROUPS])) {
             for (BindGroupIndex i : IterateBitSet(mLastPipelineLayout->GetBindGroupLayoutsMask())) {
@@ -290,12 +298,15 @@
     }
 
     void CommandBufferStateTracker::SetRenderPipeline(RenderPipelineBase* pipeline) {
-        mLastRenderPipeline = pipeline;
         SetPipelineCommon(pipeline);
     }
 
-    void CommandBufferStateTracker::SetBindGroup(BindGroupIndex index, BindGroupBase* bindgroup) {
+    void CommandBufferStateTracker::SetBindGroup(BindGroupIndex index,
+                                                 BindGroupBase* bindgroup,
+                                                 uint32_t dynamicOffsetCount,
+                                                 const uint32_t* dynamicOffsets) {
         mBindgroups[index] = bindgroup;
+        mDynamicOffsets[index].assign(dynamicOffsets, dynamicOffsets + dynamicOffsetCount);
         mAspects.reset(VALIDATION_ASPECT_BIND_GROUPS);
     }
 
@@ -311,8 +322,9 @@
     }
 
     void CommandBufferStateTracker::SetPipelineCommon(PipelineBase* pipeline) {
-        mLastPipelineLayout = pipeline->GetLayout();
-        mMinBufferSizes = &pipeline->GetMinBufferSizes();
+        mLastPipeline = pipeline;
+        mLastPipelineLayout = pipeline != nullptr ? pipeline->GetLayout() : nullptr;
+        mMinBufferSizes = pipeline != nullptr ? &pipeline->GetMinBufferSizes() : nullptr;
 
         mAspects.set(VALIDATION_ASPECT_PIPELINE);
 
@@ -324,6 +336,25 @@
         return mBindgroups[index];
     }
 
+    const std::vector<uint32_t>& CommandBufferStateTracker::GetDynamicOffsets(
+        BindGroupIndex index) const {
+        return mDynamicOffsets[index];
+    }
+
+    bool CommandBufferStateTracker::HasPipeline() const {
+        return mLastPipeline != nullptr;
+    }
+
+    RenderPipelineBase* CommandBufferStateTracker::GetRenderPipeline() const {
+        ASSERT(HasPipeline() && mLastPipeline->GetType() == ObjectType::RenderPipeline);
+        return static_cast<RenderPipelineBase*>(mLastPipeline);
+    }
+
+    ComputePipelineBase* CommandBufferStateTracker::GetComputePipeline() const {
+        ASSERT(HasPipeline() && mLastPipeline->GetType() == ObjectType::ComputePipeline);
+        return static_cast<ComputePipelineBase*>(mLastPipeline);
+    }
+
     PipelineLayoutBase* CommandBufferStateTracker::GetPipelineLayout() const {
         return mLastPipelineLayout;
     }
diff --git a/src/dawn_native/CommandBufferStateTracker.h b/src/dawn_native/CommandBufferStateTracker.h
index 0a6c587..5686956 100644
--- a/src/dawn_native/CommandBufferStateTracker.h
+++ b/src/dawn_native/CommandBufferStateTracker.h
@@ -38,7 +38,10 @@
         // State-modifying methods
         void SetComputePipeline(ComputePipelineBase* pipeline);
         void SetRenderPipeline(RenderPipelineBase* pipeline);
-        void SetBindGroup(BindGroupIndex index, BindGroupBase* bindgroup);
+        void SetBindGroup(BindGroupIndex index,
+                          BindGroupBase* bindgroup,
+                          uint32_t dynamicOffsetCount,
+                          const uint32_t* dynamicOffsets);
         void SetIndexBuffer(wgpu::IndexFormat format, uint64_t size);
         void SetVertexBuffer(VertexBufferSlot slot, uint64_t size);
 
@@ -46,6 +49,10 @@
         using ValidationAspects = std::bitset<kNumAspects>;
 
         BindGroupBase* GetBindGroup(BindGroupIndex index) const;
+        const std::vector<uint32_t>& GetDynamicOffsets(BindGroupIndex index) const;
+        bool HasPipeline() const;
+        RenderPipelineBase* GetRenderPipeline() const;
+        ComputePipelineBase* GetComputePipeline() const;
         PipelineLayoutBase* GetPipelineLayout() const;
         wgpu::IndexFormat GetIndexFormat() const;
         uint64_t GetIndexBufferSize() const;
@@ -60,6 +67,7 @@
         ValidationAspects mAspects;
 
         ityp::array<BindGroupIndex, BindGroupBase*, kMaxBindGroups> mBindgroups = {};
+        ityp::array<BindGroupIndex, std::vector<uint32_t>, kMaxBindGroups> mDynamicOffsets = {};
         ityp::bitset<VertexBufferSlot, kMaxVertexBuffers> mVertexBufferSlotsUsed;
         bool mIndexBufferSet = false;
         wgpu::IndexFormat mIndexFormat;
@@ -68,7 +76,7 @@
         ityp::array<VertexBufferSlot, uint64_t, kMaxVertexBuffers> mVertexBufferSizes = {};
 
         PipelineLayoutBase* mLastPipelineLayout = nullptr;
-        RenderPipelineBase* mLastRenderPipeline = nullptr;
+        PipelineBase* mLastPipeline = nullptr;
 
         const RequiredBufferSizes* mMinBufferSizes = nullptr;
     };
diff --git a/src/dawn_native/ComputePassEncoder.cpp b/src/dawn_native/ComputePassEncoder.cpp
index 1aa4845..05c68fb 100644
--- a/src/dawn_native/ComputePassEncoder.cpp
+++ b/src/dawn_native/ComputePassEncoder.cpp
@@ -14,18 +14,107 @@
 
 #include "dawn_native/ComputePassEncoder.h"
 
+#include "dawn_native/BindGroup.h"
+#include "dawn_native/BindGroupLayout.h"
 #include "dawn_native/Buffer.h"
 #include "dawn_native/CommandEncoder.h"
 #include "dawn_native/CommandValidation.h"
 #include "dawn_native/Commands.h"
 #include "dawn_native/ComputePipeline.h"
 #include "dawn_native/Device.h"
+#include "dawn_native/InternalPipelineStore.h"
 #include "dawn_native/ObjectType_autogen.h"
 #include "dawn_native/PassResourceUsageTracker.h"
 #include "dawn_native/QuerySet.h"
 
 namespace dawn_native {
 
+    namespace {
+
+        ResultOrError<ComputePipelineBase*> GetOrCreateIndirectDispatchValidationPipeline(
+            DeviceBase* device) {
+            InternalPipelineStore* store = device->GetInternalPipelineStore();
+
+            if (store->dispatchIndirectValidationPipeline != nullptr) {
+                return store->dispatchIndirectValidationPipeline.Get();
+            }
+
+            ShaderModuleDescriptor descriptor;
+            ShaderModuleWGSLDescriptor wgslDesc;
+            descriptor.nextInChain = reinterpret_cast<ChainedStruct*>(&wgslDesc);
+
+            // TODO(https://crbug.com/dawn/1108): Propagate validation feedback from this
+            // shader in various failure modes.
+            wgslDesc.source = R"(
+                [[block]] struct UniformParams {
+                    maxComputeWorkgroupsPerDimension: u32;
+                    clientOffsetInU32: u32;
+                };
+
+                [[block]] struct IndirectParams {
+                    data: array<u32>;
+                };
+
+                [[block]] struct ValidatedParams {
+                    data: array<u32, 3>;
+                };
+
+                [[group(0), binding(0)]] var<uniform> uniformParams: UniformParams;
+                [[group(0), binding(1)]] var<storage, read_write> clientParams: IndirectParams;
+                [[group(0), binding(2)]] var<storage, write> validatedParams: ValidatedParams;
+
+                [[stage(compute), workgroup_size(1, 1, 1)]]
+                fn main() {
+                    for (var i = 0u; i < 3u; i = i + 1u) {
+                        var numWorkgroups = clientParams.data[uniformParams.clientOffsetInU32 + i];
+                        if (numWorkgroups > uniformParams.maxComputeWorkgroupsPerDimension) {
+                            numWorkgroups = 0u;
+                        }
+                        validatedParams.data[i] = numWorkgroups;
+                    }
+                }
+            )";
+
+            Ref<ShaderModuleBase> shaderModule;
+            DAWN_TRY_ASSIGN(shaderModule, device->CreateShaderModule(&descriptor));
+
+            std::array<BindGroupLayoutEntry, 3> entries;
+            entries[0].binding = 0;
+            entries[0].visibility = wgpu::ShaderStage::Compute;
+            entries[0].buffer.type = wgpu::BufferBindingType::Uniform;
+            entries[1].binding = 1;
+            entries[1].visibility = wgpu::ShaderStage::Compute;
+            entries[1].buffer.type = kInternalStorageBufferBinding;
+            entries[2].binding = 2;
+            entries[2].visibility = wgpu::ShaderStage::Compute;
+            entries[2].buffer.type = wgpu::BufferBindingType::Storage;
+
+            BindGroupLayoutDescriptor bindGroupLayoutDescriptor;
+            bindGroupLayoutDescriptor.entryCount = entries.size();
+            bindGroupLayoutDescriptor.entries = entries.data();
+            Ref<BindGroupLayoutBase> bindGroupLayout;
+            DAWN_TRY_ASSIGN(bindGroupLayout,
+                            device->CreateBindGroupLayout(&bindGroupLayoutDescriptor, true));
+
+            PipelineLayoutDescriptor pipelineDescriptor;
+            pipelineDescriptor.bindGroupLayoutCount = 1;
+            pipelineDescriptor.bindGroupLayouts = &bindGroupLayout.Get();
+            Ref<PipelineLayoutBase> pipelineLayout;
+            DAWN_TRY_ASSIGN(pipelineLayout, device->CreatePipelineLayout(&pipelineDescriptor));
+
+            ComputePipelineDescriptor computePipelineDescriptor = {};
+            computePipelineDescriptor.layout = pipelineLayout.Get();
+            computePipelineDescriptor.compute.module = shaderModule.Get();
+            computePipelineDescriptor.compute.entryPoint = "main";
+
+            DAWN_TRY_ASSIGN(store->dispatchIndirectValidationPipeline,
+                            device->CreateComputePipeline(&computePipelineDescriptor));
+
+            return store->dispatchIndirectValidationPipeline.Get();
+        }
+
+    }  // namespace
+
     ComputePassEncoder::ComputePassEncoder(DeviceBase* device,
                                            CommandEncoder* commandEncoder,
                                            EncodingContext* encodingContext)
@@ -107,6 +196,95 @@
             "encoding Dispatch (x: %u, y: %u, z: %u)", x, y, z);
     }
 
+    ResultOrError<std::pair<Ref<BufferBase>, uint64_t>>
+    ComputePassEncoder::ValidateIndirectDispatch(BufferBase* indirectBuffer,
+                                                 uint64_t indirectOffset) {
+        DeviceBase* device = GetDevice();
+        auto* const store = device->GetInternalPipelineStore();
+
+        Ref<ComputePipelineBase> validationPipeline;
+        DAWN_TRY_ASSIGN(validationPipeline, GetOrCreateIndirectDispatchValidationPipeline(device));
+
+        Ref<BindGroupLayoutBase> layout;
+        DAWN_TRY_ASSIGN(layout, validationPipeline->GetBindGroupLayout(0));
+
+        uint32_t storageBufferOffsetAlignment =
+            device->GetLimits().v1.minStorageBufferOffsetAlignment;
+
+        std::array<BindGroupEntry, 3> bindings;
+
+        // Storage binding holding the client's indirect buffer.
+        BindGroupEntry& clientIndirectBinding = bindings[0];
+        clientIndirectBinding.binding = 1;
+        clientIndirectBinding.buffer = indirectBuffer;
+
+        // Let the offset be the indirectOffset, aligned down to |storageBufferOffsetAlignment|.
+        const uint32_t clientOffsetFromAlignedBoundary =
+            indirectOffset % storageBufferOffsetAlignment;
+        const uint64_t clientOffsetAlignedDown = indirectOffset - clientOffsetFromAlignedBoundary;
+        clientIndirectBinding.offset = clientOffsetAlignedDown;
+
+        // Let the size of the binding be the additional offset, plus the size.
+        clientIndirectBinding.size = kDispatchIndirectSize + clientOffsetFromAlignedBoundary;
+
+        struct UniformParams {
+            uint32_t maxComputeWorkgroupsPerDimension;
+            uint32_t clientOffsetInU32;
+        };
+
+        // Create a uniform buffer to hold parameters for the shader.
+        Ref<BufferBase> uniformBuffer;
+        {
+            BufferDescriptor uniformDesc = {};
+            uniformDesc.size = sizeof(UniformParams);
+            uniformDesc.usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst;
+            uniformDesc.mappedAtCreation = true;
+            DAWN_TRY_ASSIGN(uniformBuffer, device->CreateBuffer(&uniformDesc));
+
+            UniformParams* params = static_cast<UniformParams*>(
+                uniformBuffer->GetMappedRange(0, sizeof(UniformParams)));
+            params->maxComputeWorkgroupsPerDimension =
+                device->GetLimits().v1.maxComputeWorkgroupsPerDimension;
+            params->clientOffsetInU32 = clientOffsetFromAlignedBoundary / sizeof(uint32_t);
+            uniformBuffer->Unmap();
+        }
+
+        // Uniform buffer binding pointing to the uniform parameters.
+        BindGroupEntry& uniformBinding = bindings[1];
+        uniformBinding.binding = 0;
+        uniformBinding.buffer = uniformBuffer.Get();
+        uniformBinding.offset = 0;
+        uniformBinding.size = sizeof(UniformParams);
+
+        // Reserve space in the scratch buffer to hold the validated indirect params.
+        ScratchBuffer& scratchBuffer = store->scratchIndirectStorage;
+        DAWN_TRY(scratchBuffer.EnsureCapacity(kDispatchIndirectSize));
+        Ref<BufferBase> validatedIndirectBuffer = scratchBuffer.GetBuffer();
+
+        // Binding for the validated indirect params.
+        BindGroupEntry& validatedParamsBinding = bindings[2];
+        validatedParamsBinding.binding = 2;
+        validatedParamsBinding.buffer = validatedIndirectBuffer.Get();
+        validatedParamsBinding.offset = 0;
+        validatedParamsBinding.size = kDispatchIndirectSize;
+
+        BindGroupDescriptor bindGroupDescriptor = {};
+        bindGroupDescriptor.layout = layout.Get();
+        bindGroupDescriptor.entryCount = bindings.size();
+        bindGroupDescriptor.entries = bindings.data();
+
+        Ref<BindGroupBase> validationBindGroup;
+        DAWN_TRY_ASSIGN(validationBindGroup, device->CreateBindGroup(&bindGroupDescriptor));
+
+        // Issue commands to validate the indirect buffer.
+        APISetPipeline(validationPipeline.Get());
+        APISetBindGroup(0, validationBindGroup.Get());
+        APIDispatch(1);
+
+        // Return the new indirect buffer and indirect buffer offset.
+        return std::make_pair(std::move(validatedIndirectBuffer), uint64_t(0));
+    }
+
     void ComputePassEncoder::APIDispatchIndirect(BufferBase* indirectBuffer,
                                                  uint64_t indirectOffset) {
         mEncodingContext->TryEncode(
@@ -136,18 +314,46 @@
                         indirectOffset, kDispatchIndirectSize, indirectBuffer->GetSize());
                 }
 
-                // Record the synchronization scope for Dispatch, both the bindgroups and the
-                // indirect buffer.
                 SyncScopeUsageTracker scope;
                 scope.BufferUsedAs(indirectBuffer, wgpu::BufferUsage::Indirect);
                 mUsageTracker.AddReferencedBuffer(indirectBuffer);
+                // TODO(crbug.com/dawn/1166): If validation is enabled, adding |indirectBuffer|
+                // is needed for correct usage validation even though it will only be bound for
+                // storage. This will unecessarily transition the |indirectBuffer| in
+                // the backend.
+
+                Ref<BufferBase> indirectBufferRef = indirectBuffer;
+                if (IsValidationEnabled()) {
+                    // Save the previous command buffer state so it can be restored after the
+                    // validation inserts additional commands.
+                    CommandBufferStateTracker previousState = mCommandBufferState;
+
+                    // Validate each indirect dispatch with a single dispatch to copy the indirect
+                    // buffer params into a scratch buffer if they're valid, and otherwise zero them
+                    // out. We could consider moving the validation earlier in the pass after the
+                    // last point the indirect buffer was used with writable usage, as well as batch
+                    // validation for multiple dispatches into one, but inserting commands at
+                    // arbitrary points in the past is not possible right now.
+                    DAWN_TRY_ASSIGN(
+                        std::tie(indirectBufferRef, indirectOffset),
+                        ValidateIndirectDispatch(indirectBufferRef.Get(), indirectOffset));
+
+                    // Restore the state.
+                    RestoreCommandBufferState(std::move(previousState));
+
+                    // |indirectBufferRef| was replaced with a scratch buffer. Add it to the
+                    // synchronization scope.
+                    ASSERT(indirectBufferRef.Get() != indirectBuffer);
+                    scope.BufferUsedAs(indirectBufferRef.Get(), wgpu::BufferUsage::Indirect);
+                    mUsageTracker.AddReferencedBuffer(indirectBufferRef.Get());
+                }
+
                 AddDispatchSyncScope(std::move(scope));
 
                 DispatchIndirectCmd* dispatch =
                     allocator->Allocate<DispatchIndirectCmd>(Command::DispatchIndirect);
-                dispatch->indirectBuffer = indirectBuffer;
+                dispatch->indirectBuffer = std::move(indirectBufferRef);
                 dispatch->indirectOffset = indirectOffset;
-
                 return {};
             },
             "encoding DispatchIndirect with %s", indirectBuffer);
@@ -187,10 +393,10 @@
                 }
 
                 mUsageTracker.AddResourcesReferencedByBindGroup(group);
-
                 RecordSetBindGroup(allocator, groupIndex, group, dynamicOffsetCount,
                                    dynamicOffsets);
-                mCommandBufferState.SetBindGroup(groupIndex, group);
+                mCommandBufferState.SetBindGroup(groupIndex, group, dynamicOffsetCount,
+                                                 dynamicOffsets);
 
                 return {};
             },
@@ -226,4 +432,29 @@
         mUsageTracker.AddDispatch(scope.AcquireSyncScopeUsage());
     }
 
+    void ComputePassEncoder::RestoreCommandBufferState(CommandBufferStateTracker state) {
+        // Encode commands for the backend to restore the pipeline and bind groups.
+        if (state.HasPipeline()) {
+            APISetPipeline(state.GetComputePipeline());
+        }
+        for (BindGroupIndex i(0); i < kMaxBindGroupsTyped; ++i) {
+            BindGroupBase* bg = state.GetBindGroup(i);
+            if (bg != nullptr) {
+                const std::vector<uint32_t>& offsets = state.GetDynamicOffsets(i);
+                if (offsets.empty()) {
+                    APISetBindGroup(static_cast<uint32_t>(i), bg);
+                } else {
+                    APISetBindGroup(static_cast<uint32_t>(i), bg, offsets.size(), offsets.data());
+                }
+            }
+        }
+
+        // Restore the frontend state tracking information.
+        mCommandBufferState = std::move(state);
+    }
+
+    CommandBufferStateTracker* ComputePassEncoder::GetCommandBufferStateTrackerForTesting() {
+        return &mCommandBufferState;
+    }
+
 }  // namespace dawn_native
diff --git a/src/dawn_native/ComputePassEncoder.h b/src/dawn_native/ComputePassEncoder.h
index b0962f4..03997ce 100644
--- a/src/dawn_native/ComputePassEncoder.h
+++ b/src/dawn_native/ComputePassEncoder.h
@@ -50,6 +50,11 @@
 
         void APIWriteTimestamp(QuerySetBase* querySet, uint32_t queryIndex);
 
+        CommandBufferStateTracker* GetCommandBufferStateTrackerForTesting();
+        void RestoreCommandBufferStateForTesting(CommandBufferStateTracker state) {
+            RestoreCommandBufferState(std::move(state));
+        }
+
       protected:
         ComputePassEncoder(DeviceBase* device,
                            CommandEncoder* commandEncoder,
@@ -57,6 +62,12 @@
                            ErrorTag errorTag);
 
       private:
+        ResultOrError<std::pair<Ref<BufferBase>, uint64_t>> ValidateIndirectDispatch(
+            BufferBase* indirectBuffer,
+            uint64_t indirectOffset);
+
+        void RestoreCommandBufferState(CommandBufferStateTracker state);
+
         CommandBufferStateTracker mCommandBufferState;
 
         // Adds the bindgroups used for the current dispatch to the SyncScopeResourceUsage and
diff --git a/src/dawn_native/InternalPipelineStore.h b/src/dawn_native/InternalPipelineStore.h
index acf3b13..803e0df 100644
--- a/src/dawn_native/InternalPipelineStore.h
+++ b/src/dawn_native/InternalPipelineStore.h
@@ -52,6 +52,7 @@
 
         Ref<ComputePipelineBase> renderValidationPipeline;
         Ref<ShaderModuleBase> renderValidationShader;
+        Ref<ComputePipelineBase> dispatchIndirectValidationPipeline;
     };
 
 }  // namespace dawn_native
diff --git a/src/dawn_native/RenderEncoderBase.cpp b/src/dawn_native/RenderEncoderBase.cpp
index 0445a97..a8ef2ff 100644
--- a/src/dawn_native/RenderEncoderBase.cpp
+++ b/src/dawn_native/RenderEncoderBase.cpp
@@ -208,6 +208,9 @@
                         BufferLocation::New(indirectBuffer, indirectOffset);
                 }
 
+                // TODO(crbug.com/dawn/1166): Adding the indirectBuffer is needed for correct usage
+                // validation, but it will unecessarily transition to indirectBuffer usage in the
+                // backend.
                 mUsageTracker.BufferUsedAs(indirectBuffer, wgpu::BufferUsage::Indirect);
 
                 return {};
@@ -404,7 +407,8 @@
 
                 RecordSetBindGroup(allocator, groupIndex, group, dynamicOffsetCount,
                                    dynamicOffsets);
-                mCommandBufferState.SetBindGroup(groupIndex, group);
+                mCommandBufferState.SetBindGroup(groupIndex, group, dynamicOffsetCount,
+                                                 dynamicOffsets);
                 mUsageTracker.AddBindGroup(group);
 
                 return {};
diff --git a/src/tests/BUILD.gn b/src/tests/BUILD.gn
index c7f4b7a..8ffd921 100644
--- a/src/tests/BUILD.gn
+++ b/src/tests/BUILD.gn
@@ -221,6 +221,7 @@
     "unittests/SystemUtilsTests.cpp",
     "unittests/ToBackendTests.cpp",
     "unittests/TypedIntegerTests.cpp",
+    "unittests/native/CommandBufferEncodingTests.cpp",
     "unittests/native/DestroyObjectTests.cpp",
     "unittests/validation/BindGroupValidationTests.cpp",
     "unittests/validation/BufferValidationTests.cpp",
diff --git a/src/tests/DawnNativeTest.cpp b/src/tests/DawnNativeTest.cpp
index d39c8e0..28d69bf 100644
--- a/src/tests/DawnNativeTest.cpp
+++ b/src/tests/DawnNativeTest.cpp
@@ -12,9 +12,11 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include <gtest/gtest.h>
+#include "tests/DawnNativeTest.h"
 
 #include "absl/strings/str_cat.h"
+#include "common/Assert.h"
+#include "dawn/dawn_proc.h"
 #include "dawn_native/ErrorData.h"
 
 namespace dawn_native {
@@ -28,3 +30,54 @@
     }
 
 }  // namespace dawn_native
+
+DawnNativeTest::DawnNativeTest() {
+    dawnProcSetProcs(&dawn_native::GetProcs());
+}
+
+DawnNativeTest::~DawnNativeTest() {
+    device = wgpu::Device();
+    dawnProcSetProcs(nullptr);
+}
+
+void DawnNativeTest::SetUp() {
+    instance = std::make_unique<dawn_native::Instance>();
+    instance->DiscoverDefaultAdapters();
+
+    std::vector<dawn_native::Adapter> adapters = instance->GetAdapters();
+
+    // DawnNative unittests run against the null backend, find the corresponding adapter
+    bool foundNullAdapter = false;
+    for (auto& currentAdapter : adapters) {
+        wgpu::AdapterProperties adapterProperties;
+        currentAdapter.GetProperties(&adapterProperties);
+
+        if (adapterProperties.backendType == wgpu::BackendType::Null) {
+            adapter = currentAdapter;
+            foundNullAdapter = true;
+            break;
+        }
+    }
+
+    ASSERT(foundNullAdapter);
+
+    device = wgpu::Device(CreateTestDevice());
+    device.SetUncapturedErrorCallback(DawnNativeTest::OnDeviceError, nullptr);
+}
+
+void DawnNativeTest::TearDown() {
+}
+
+WGPUDevice DawnNativeTest::CreateTestDevice() {
+    // Disabled disallowing unsafe APIs so we can test them.
+    dawn_native::DeviceDescriptor deviceDescriptor;
+    deviceDescriptor.forceDisabledToggles.push_back("disallow_unsafe_apis");
+
+    return adapter.CreateDevice(&deviceDescriptor);
+}
+
+// static
+void DawnNativeTest::OnDeviceError(WGPUErrorType type, const char* message, void* userdata) {
+    ASSERT(type != WGPUErrorType_NoError);
+    FAIL() << "Unexpected error: " << message;
+}
diff --git a/src/tests/DawnNativeTest.h b/src/tests/DawnNativeTest.h
index 94fdafb..91904a3 100644
--- a/src/tests/DawnNativeTest.h
+++ b/src/tests/DawnNativeTest.h
@@ -17,6 +17,8 @@
 
 #include <gtest/gtest.h>
 
+#include "dawn/webgpu_cpp.h"
+#include "dawn_native/DawnNative.h"
 #include "dawn_native/ErrorData.h"
 
 namespace dawn_native {
@@ -29,4 +31,23 @@
 
 }  // namespace dawn_native
 
+class DawnNativeTest : public ::testing::Test {
+  public:
+    DawnNativeTest();
+    ~DawnNativeTest() override;
+
+    void SetUp() override;
+    void TearDown() override;
+
+    virtual WGPUDevice CreateTestDevice();
+
+  protected:
+    std::unique_ptr<dawn_native::Instance> instance;
+    dawn_native::Adapter adapter;
+    wgpu::Device device;
+
+  private:
+    static void OnDeviceError(WGPUErrorType type, const char* message, void* userdata);
+};
+
 #endif  // TESTS_DAWNNATIVETEST_H_
diff --git a/src/tests/end2end/ComputeDispatchTests.cpp b/src/tests/end2end/ComputeDispatchTests.cpp
index 1a8b163..cbc2d86 100644
--- a/src/tests/end2end/ComputeDispatchTests.cpp
+++ b/src/tests/end2end/ComputeDispatchTests.cpp
@@ -158,8 +158,14 @@
         queue.Submit(1, &commands);
 
         std::vector<uint32_t> expected;
+
+        uint32_t maxComputeWorkgroupsPerDimension =
+            GetSupportedLimits().limits.maxComputeWorkgroupsPerDimension;
         if (indirectBufferData[indirectStart] == 0 || indirectBufferData[indirectStart + 1] == 0 ||
-            indirectBufferData[indirectStart + 2] == 0) {
+            indirectBufferData[indirectStart + 2] == 0 ||
+            indirectBufferData[indirectStart] > maxComputeWorkgroupsPerDimension ||
+            indirectBufferData[indirectStart + 1] > maxComputeWorkgroupsPerDimension ||
+            indirectBufferData[indirectStart + 2] > maxComputeWorkgroupsPerDimension) {
             expected = kSentinelData;
         } else {
             expected.assign(indirectBufferData.begin() + indirectStart,
@@ -221,6 +227,52 @@
     IndirectTest({0, 0, 0, 2, 3, 4}, 3 * sizeof(uint32_t));
 }
 
+// Test indirect dispatches at max limit.
+TEST_P(ComputeDispatchTests, MaxWorkgroups) {
+    // TODO(crbug.com/dawn/1165): Fails with WARP
+    DAWN_SUPPRESS_TEST_IF(IsWARP());
+    uint32_t max = GetSupportedLimits().limits.maxComputeWorkgroupsPerDimension;
+
+    // Test that the maximum works in each dimension.
+    // Note: Testing (max, max, max) is very slow.
+    IndirectTest({max, 3, 4}, 0);
+    IndirectTest({2, max, 4}, 0);
+    IndirectTest({2, 3, max}, 0);
+}
+
+// Test indirect dispatches exceeding the max limit are noop-ed.
+TEST_P(ComputeDispatchTests, ExceedsMaxWorkgroupsNoop) {
+    DAWN_TEST_UNSUPPORTED_IF(HasToggleEnabled("skip_validation"));
+
+    uint32_t max = GetSupportedLimits().limits.maxComputeWorkgroupsPerDimension;
+
+    // All dimensions are above the max
+    IndirectTest({max + 1, max + 1, max + 1}, 0);
+
+    // Only x dimension is above the max
+    IndirectTest({max + 1, 3, 4}, 0);
+    IndirectTest({2 * max, 3, 4}, 0);
+
+    // Only y dimension is above the max
+    IndirectTest({2, max + 1, 4}, 0);
+    IndirectTest({2, 2 * max, 4}, 0);
+
+    // Only z dimension is above the max
+    IndirectTest({2, 3, max + 1}, 0);
+    IndirectTest({2, 3, 2 * max}, 0);
+}
+
+// Test indirect dispatches exceeding the max limit with an offset are noop-ed.
+TEST_P(ComputeDispatchTests, ExceedsMaxWorkgroupsWithOffsetNoop) {
+    DAWN_TEST_UNSUPPORTED_IF(HasToggleEnabled("skip_validation"));
+
+    uint32_t max = GetSupportedLimits().limits.maxComputeWorkgroupsPerDimension;
+
+    IndirectTest({1, 2, 3, max + 1, 4, 5}, 1 * sizeof(uint32_t));
+    IndirectTest({1, 2, 3, max + 1, 4, 5}, 2 * sizeof(uint32_t));
+    IndirectTest({1, 2, 3, max + 1, 4, 5}, 3 * sizeof(uint32_t));
+}
+
 DAWN_INSTANTIATE_TEST(ComputeDispatchTests,
                       D3D12Backend(),
                       MetalBackend(),
diff --git a/src/tests/unittests/native/CommandBufferEncodingTests.cpp b/src/tests/unittests/native/CommandBufferEncodingTests.cpp
new file mode 100644
index 0000000..c1ca2d9
--- /dev/null
+++ b/src/tests/unittests/native/CommandBufferEncodingTests.cpp
@@ -0,0 +1,310 @@
+// Copyright 2021 The Dawn Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "tests/DawnNativeTest.h"
+
+#include "dawn_native/CommandBuffer.h"
+#include "dawn_native/Commands.h"
+#include "dawn_native/ComputePassEncoder.h"
+#include "utils/WGPUHelpers.h"
+
+class CommandBufferEncodingTests : public DawnNativeTest {
+  protected:
+    void ExpectCommands(dawn_native::CommandIterator* commands,
+                        std::vector<std::pair<dawn_native::Command,
+                                              std::function<void(dawn_native::CommandIterator*)>>>
+                            expectedCommands) {
+        dawn_native::Command commandId;
+        for (uint32_t commandIndex = 0; commands->NextCommandId(&commandId); ++commandIndex) {
+            ASSERT_LT(commandIndex, expectedCommands.size()) << "Unexpected command";
+            ASSERT_EQ(commandId, expectedCommands[commandIndex].first)
+                << "at command " << commandIndex;
+            expectedCommands[commandIndex].second(commands);
+        }
+    }
+};
+
+// Indirect dispatch validation changes the bind groups in the middle
+// of a pass. Test that bindings are restored after the validation runs.
+TEST_F(CommandBufferEncodingTests, ComputePassEncoderIndirectDispatchStateRestoration) {
+    using namespace dawn_native;
+
+    wgpu::BindGroupLayout staticLayout =
+        utils::MakeBindGroupLayout(device, {{
+                                               0,
+                                               wgpu::ShaderStage::Compute,
+                                               wgpu::BufferBindingType::Uniform,
+                                           }});
+
+    wgpu::BindGroupLayout dynamicLayout =
+        utils::MakeBindGroupLayout(device, {{
+                                               0,
+                                               wgpu::ShaderStage::Compute,
+                                               wgpu::BufferBindingType::Uniform,
+                                               true,
+                                           }});
+
+    // Create a simple pipeline
+    wgpu::ComputePipelineDescriptor csDesc;
+    csDesc.compute.module = utils::CreateShaderModule(device, R"(
+        [[stage(compute), workgroup_size(1, 1, 1)]]
+        fn main() {
+        })");
+    csDesc.compute.entryPoint = "main";
+
+    wgpu::PipelineLayout pl0 = utils::MakePipelineLayout(device, {staticLayout, dynamicLayout});
+    csDesc.layout = pl0;
+    wgpu::ComputePipeline pipeline0 = device.CreateComputePipeline(&csDesc);
+
+    wgpu::PipelineLayout pl1 = utils::MakePipelineLayout(device, {dynamicLayout, staticLayout});
+    csDesc.layout = pl1;
+    wgpu::ComputePipeline pipeline1 = device.CreateComputePipeline(&csDesc);
+
+    // Create buffers to use for both the indirect buffer and the bind groups.
+    wgpu::Buffer indirectBuffer =
+        utils::CreateBufferFromData<uint32_t>(device, wgpu::BufferUsage::Indirect, {1, 2, 3, 4});
+
+    wgpu::BufferDescriptor uniformBufferDesc = {};
+    uniformBufferDesc.size = 512;
+    uniformBufferDesc.usage = wgpu::BufferUsage::Uniform;
+    wgpu::Buffer uniformBuffer = device.CreateBuffer(&uniformBufferDesc);
+
+    wgpu::BindGroup staticBG = utils::MakeBindGroup(device, staticLayout, {{0, uniformBuffer}});
+
+    wgpu::BindGroup dynamicBG =
+        utils::MakeBindGroup(device, dynamicLayout, {{0, uniformBuffer, 0, 256}});
+
+    uint32_t dynamicOffset = 256;
+    std::vector<uint32_t> emptyDynamicOffsets = {};
+    std::vector<uint32_t> singleDynamicOffset = {dynamicOffset};
+
+    // Begin encoding commands.
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+
+    CommandBufferStateTracker* stateTracker =
+        FromAPI(pass.Get())->GetCommandBufferStateTrackerForTesting();
+
+    // Perform a dispatch indirect which will be preceded by a validation dispatch.
+    pass.SetPipeline(pipeline0);
+    pass.SetBindGroup(0, staticBG);
+    pass.SetBindGroup(1, dynamicBG, 1, &dynamicOffset);
+    EXPECT_EQ(ToAPI(stateTracker->GetComputePipeline()), pipeline0.Get());
+
+    pass.DispatchIndirect(indirectBuffer, 0);
+
+    // Expect restored state.
+    EXPECT_EQ(ToAPI(stateTracker->GetComputePipeline()), pipeline0.Get());
+    EXPECT_EQ(ToAPI(stateTracker->GetPipelineLayout()), pl0.Get());
+    EXPECT_EQ(ToAPI(stateTracker->GetBindGroup(BindGroupIndex(0))), staticBG.Get());
+    EXPECT_EQ(stateTracker->GetDynamicOffsets(BindGroupIndex(0)), emptyDynamicOffsets);
+    EXPECT_EQ(ToAPI(stateTracker->GetBindGroup(BindGroupIndex(1))), dynamicBG.Get());
+    EXPECT_EQ(stateTracker->GetDynamicOffsets(BindGroupIndex(1)), singleDynamicOffset);
+
+    // Dispatch again to check that the restored state can be used.
+    // Also pass an indirect offset which should get replaced with the offset
+    // into the scratch indirect buffer (0).
+    pass.DispatchIndirect(indirectBuffer, 4);
+
+    // Expect restored state.
+    EXPECT_EQ(ToAPI(stateTracker->GetComputePipeline()), pipeline0.Get());
+    EXPECT_EQ(ToAPI(stateTracker->GetPipelineLayout()), pl0.Get());
+    EXPECT_EQ(ToAPI(stateTracker->GetBindGroup(BindGroupIndex(0))), staticBG.Get());
+    EXPECT_EQ(stateTracker->GetDynamicOffsets(BindGroupIndex(0)), emptyDynamicOffsets);
+    EXPECT_EQ(ToAPI(stateTracker->GetBindGroup(BindGroupIndex(1))), dynamicBG.Get());
+    EXPECT_EQ(stateTracker->GetDynamicOffsets(BindGroupIndex(1)), singleDynamicOffset);
+
+    // Change the pipeline
+    pass.SetPipeline(pipeline1);
+    pass.SetBindGroup(0, dynamicBG, 1, &dynamicOffset);
+    pass.SetBindGroup(1, staticBG);
+    EXPECT_EQ(ToAPI(stateTracker->GetComputePipeline()), pipeline1.Get());
+    EXPECT_EQ(ToAPI(stateTracker->GetPipelineLayout()), pl1.Get());
+
+    pass.DispatchIndirect(indirectBuffer, 0);
+
+    // Expect restored state.
+    EXPECT_EQ(ToAPI(stateTracker->GetComputePipeline()), pipeline1.Get());
+    EXPECT_EQ(ToAPI(stateTracker->GetPipelineLayout()), pl1.Get());
+    EXPECT_EQ(ToAPI(stateTracker->GetBindGroup(BindGroupIndex(0))), dynamicBG.Get());
+    EXPECT_EQ(stateTracker->GetDynamicOffsets(BindGroupIndex(0)), singleDynamicOffset);
+    EXPECT_EQ(ToAPI(stateTracker->GetBindGroup(BindGroupIndex(1))), staticBG.Get());
+    EXPECT_EQ(stateTracker->GetDynamicOffsets(BindGroupIndex(1)), emptyDynamicOffsets);
+
+    pass.EndPass();
+
+    wgpu::CommandBuffer commandBuffer = encoder.Finish();
+
+    auto ExpectSetPipeline = [](wgpu::ComputePipeline pipeline) {
+        return [pipeline](CommandIterator* commands) {
+            auto* cmd = commands->NextCommand<SetComputePipelineCmd>();
+            EXPECT_EQ(ToAPI(cmd->pipeline.Get()), pipeline.Get());
+        };
+    };
+
+    auto ExpectSetBindGroup = [](uint32_t index, wgpu::BindGroup bg,
+                                 std::vector<uint32_t> offsets = {}) {
+        return [index, bg, offsets](CommandIterator* commands) {
+            auto* cmd = commands->NextCommand<SetBindGroupCmd>();
+            uint32_t* dynamicOffsets = nullptr;
+            if (cmd->dynamicOffsetCount > 0) {
+                dynamicOffsets = commands->NextData<uint32_t>(cmd->dynamicOffsetCount);
+            }
+
+            ASSERT_EQ(cmd->index, BindGroupIndex(index));
+            ASSERT_EQ(ToAPI(cmd->group.Get()), bg.Get());
+            ASSERT_EQ(cmd->dynamicOffsetCount, offsets.size());
+            for (uint32_t i = 0; i < cmd->dynamicOffsetCount; ++i) {
+                ASSERT_EQ(dynamicOffsets[i], offsets[i]);
+            }
+        };
+    };
+
+    // Initialize as null. Once we know the pointer, we'll check
+    // that it's the same buffer every time.
+    WGPUBuffer indirectScratchBuffer = nullptr;
+    auto ExpectDispatchIndirect = [&](CommandIterator* commands) {
+        auto* cmd = commands->NextCommand<DispatchIndirectCmd>();
+        if (indirectScratchBuffer == nullptr) {
+            indirectScratchBuffer = ToAPI(cmd->indirectBuffer.Get());
+        }
+        ASSERT_EQ(ToAPI(cmd->indirectBuffer.Get()), indirectScratchBuffer);
+        ASSERT_EQ(cmd->indirectOffset, uint64_t(0));
+    };
+
+    // Initialize as null. Once we know the pointer, we'll check
+    // that it's the same pipeline every time.
+    WGPUComputePipeline validationPipeline = nullptr;
+    auto ExpectSetValidationPipeline = [&](CommandIterator* commands) {
+        auto* cmd = commands->NextCommand<SetComputePipelineCmd>();
+        WGPUComputePipeline pipeline = ToAPI(cmd->pipeline.Get());
+        if (validationPipeline != nullptr) {
+            EXPECT_EQ(pipeline, validationPipeline);
+        } else {
+            EXPECT_NE(pipeline, nullptr);
+            validationPipeline = pipeline;
+        }
+    };
+
+    auto ExpectSetValidationBindGroup = [&](CommandIterator* commands) {
+        auto* cmd = commands->NextCommand<SetBindGroupCmd>();
+        ASSERT_EQ(cmd->index, BindGroupIndex(0));
+        ASSERT_NE(cmd->group.Get(), nullptr);
+        ASSERT_EQ(cmd->dynamicOffsetCount, 0u);
+    };
+
+    auto ExpectSetValidationDispatch = [&](CommandIterator* commands) {
+        auto* cmd = commands->NextCommand<DispatchCmd>();
+        ASSERT_EQ(cmd->x, 1u);
+        ASSERT_EQ(cmd->y, 1u);
+        ASSERT_EQ(cmd->z, 1u);
+    };
+
+    ExpectCommands(
+        FromAPI(commandBuffer.Get())->GetCommandIteratorForTesting(),
+        {
+            {Command::BeginComputePass,
+             [&](CommandIterator* commands) { SkipCommand(commands, Command::BeginComputePass); }},
+            // Expect the state to be set.
+            {Command::SetComputePipeline, ExpectSetPipeline(pipeline0)},
+            {Command::SetBindGroup, ExpectSetBindGroup(0, staticBG)},
+            {Command::SetBindGroup, ExpectSetBindGroup(1, dynamicBG, {dynamicOffset})},
+
+            // Expect the validation.
+            {Command::SetComputePipeline, ExpectSetValidationPipeline},
+            {Command::SetBindGroup, ExpectSetValidationBindGroup},
+            {Command::Dispatch, ExpectSetValidationDispatch},
+
+            // Expect the state to be restored.
+            {Command::SetComputePipeline, ExpectSetPipeline(pipeline0)},
+            {Command::SetBindGroup, ExpectSetBindGroup(0, staticBG)},
+            {Command::SetBindGroup, ExpectSetBindGroup(1, dynamicBG, {dynamicOffset})},
+
+            // Expect the dispatchIndirect.
+            {Command::DispatchIndirect, ExpectDispatchIndirect},
+
+            // Expect the validation.
+            {Command::SetComputePipeline, ExpectSetValidationPipeline},
+            {Command::SetBindGroup, ExpectSetValidationBindGroup},
+            {Command::Dispatch, ExpectSetValidationDispatch},
+
+            // Expect the state to be restored.
+            {Command::SetComputePipeline, ExpectSetPipeline(pipeline0)},
+            {Command::SetBindGroup, ExpectSetBindGroup(0, staticBG)},
+            {Command::SetBindGroup, ExpectSetBindGroup(1, dynamicBG, {dynamicOffset})},
+
+            // Expect the dispatchIndirect.
+            {Command::DispatchIndirect, ExpectDispatchIndirect},
+
+            // Expect the state to be set (new pipeline).
+            {Command::SetComputePipeline, ExpectSetPipeline(pipeline1)},
+            {Command::SetBindGroup, ExpectSetBindGroup(0, dynamicBG, {dynamicOffset})},
+            {Command::SetBindGroup, ExpectSetBindGroup(1, staticBG)},
+
+            // Expect the validation.
+            {Command::SetComputePipeline, ExpectSetValidationPipeline},
+            {Command::SetBindGroup, ExpectSetValidationBindGroup},
+            {Command::Dispatch, ExpectSetValidationDispatch},
+
+            // Expect the state to be restored.
+            {Command::SetComputePipeline, ExpectSetPipeline(pipeline1)},
+            {Command::SetBindGroup, ExpectSetBindGroup(0, dynamicBG, {dynamicOffset})},
+            {Command::SetBindGroup, ExpectSetBindGroup(1, staticBG)},
+
+            // Expect the dispatchIndirect.
+            {Command::DispatchIndirect, ExpectDispatchIndirect},
+
+            {Command::EndComputePass,
+             [&](CommandIterator* commands) { commands->NextCommand<EndComputePassCmd>(); }},
+        });
+}
+
+// Test that after restoring state, it is fully applied to the state tracker
+// and does not leak state changes that occured between a snapshot and the
+// state restoration.
+TEST_F(CommandBufferEncodingTests, StateNotLeakedAfterRestore) {
+    using namespace dawn_native;
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+
+    CommandBufferStateTracker* stateTracker =
+        FromAPI(pass.Get())->GetCommandBufferStateTrackerForTesting();
+
+    // Snapshot the state.
+    CommandBufferStateTracker snapshot = *stateTracker;
+    // Expect no pipeline in the snapshot
+    EXPECT_FALSE(snapshot.HasPipeline());
+
+    // Create a simple pipeline
+    wgpu::ComputePipelineDescriptor csDesc;
+    csDesc.compute.module = utils::CreateShaderModule(device, R"(
+        [[stage(compute), workgroup_size(1, 1, 1)]]
+        fn main() {
+        })");
+    csDesc.compute.entryPoint = "main";
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
+
+    // Set the pipeline.
+    pass.SetPipeline(pipeline);
+
+    // Expect the pipeline to be set.
+    EXPECT_EQ(ToAPI(stateTracker->GetComputePipeline()), pipeline.Get());
+
+    // Restore the state.
+    FromAPI(pass.Get())->RestoreCommandBufferStateForTesting(std::move(snapshot));
+
+    // Expect no pipeline
+    EXPECT_FALSE(stateTracker->HasPipeline());
+}