diff --git a/src/dawn/native/CommandBufferStateTracker.cpp b/src/dawn/native/CommandBufferStateTracker.cpp
index 2828752..0418f07 100644
--- a/src/dawn/native/CommandBufferStateTracker.cpp
+++ b/src/dawn/native/CommandBufferStateTracker.cpp
@@ -754,10 +754,13 @@
     mAspects.reset(VALIDATION_ASPECT_BIND_GROUPS);
 }
 
-void CommandBufferStateTracker::SetIndexBuffer(wgpu::IndexFormat format, uint64_t size) {
+void CommandBufferStateTracker::SetIndexBuffer(wgpu::IndexFormat format,
+                                               uint64_t offset,
+                                               uint64_t size) {
     mIndexBufferSet = true;
     mIndexFormat = format;
     mIndexBufferSize = size;
+    mIndexBufferOffset = offset;
 }
 
 void CommandBufferStateTracker::UnsetVertexBuffer(VertexBufferSlot slot) {
@@ -817,4 +820,8 @@
     return mIndexBufferSize;
 }
 
+uint64_t CommandBufferStateTracker::GetIndexBufferOffset() const {
+    return mIndexBufferOffset;
+}
+
 }  // namespace dawn::native
diff --git a/src/dawn/native/CommandBufferStateTracker.h b/src/dawn/native/CommandBufferStateTracker.h
index 0e25e41..5df4fae 100644
--- a/src/dawn/native/CommandBufferStateTracker.h
+++ b/src/dawn/native/CommandBufferStateTracker.h
@@ -68,7 +68,7 @@
                       BindGroupBase* bindgroup,
                       uint32_t dynamicOffsetCount,
                       const uint32_t* dynamicOffsets);
-    void SetIndexBuffer(wgpu::IndexFormat format, uint64_t size);
+    void SetIndexBuffer(wgpu::IndexFormat format, uint64_t offset, uint64_t size);
     void UnsetVertexBuffer(VertexBufferSlot slot);
     void SetVertexBuffer(VertexBufferSlot slot, uint64_t size);
 
@@ -83,6 +83,7 @@
     PipelineLayoutBase* GetPipelineLayout() const;
     wgpu::IndexFormat GetIndexFormat() const;
     uint64_t GetIndexBufferSize() const;
+    uint64_t GetIndexBufferOffset() const;
 
   private:
     MaybeError ValidateOperation(ValidationAspects requiredAspects);
@@ -102,6 +103,7 @@
     bool mIndexBufferSet = false;
     wgpu::IndexFormat mIndexFormat;
     uint64_t mIndexBufferSize = 0;
+    uint64_t mIndexBufferOffset = 0;
 
     // TODO(https://crbug.com/dawn/2349): Investigate DanglingUntriaged in dawn/native.
     raw_ptr<PipelineLayoutBase, DanglingUntriaged> mLastPipelineLayout = nullptr;
diff --git a/src/dawn/native/Device.cpp b/src/dawn/native/Device.cpp
index 135b928..1d4f624 100644
--- a/src/dawn/native/Device.cpp
+++ b/src/dawn/native/Device.cpp
@@ -2291,6 +2291,10 @@
     return false;
 }
 
+bool DeviceBase::ShouldApplyIndexBufferOffsetToFirstIndex() const {
+    return false;
+}
+
 bool DeviceBase::IsResolveTextureBlitWithDrawSupported() const {
     return false;
 }
diff --git a/src/dawn/native/Device.h b/src/dawn/native/Device.h
index fa1ac4f..986dabd 100644
--- a/src/dawn/native/Device.h
+++ b/src/dawn/native/Device.h
@@ -416,6 +416,12 @@
     virtual bool ShouldDuplicateParametersForDrawIndirect(
         const RenderPipelineBase* renderPipelineBase) const;
 
+    // For OpenGL/OpenGL ES, we must apply the index buffer offset from SetIndexBuffer to the
+    // firstIndex parameter in indirect buffers. This happens in the validation since it
+    // copies the indirect buffers and updates them while validating.
+    // See https://crbug.com/dawn/161
+    virtual bool ShouldApplyIndexBufferOffsetToFirstIndex() const;
+
     // Whether the backend supports blitting the resolve texture with draw calls in the same render
     // pass that it will be resolved into.
     virtual bool IsResolveTextureBlitWithDrawSupported() const;
diff --git a/src/dawn/native/IndirectDrawMetadata.cpp b/src/dawn/native/IndirectDrawMetadata.cpp
index 5963655..1572466 100644
--- a/src/dawn/native/IndirectDrawMetadata.cpp
+++ b/src/dawn/native/IndirectDrawMetadata.cpp
@@ -170,17 +170,21 @@
 
 void IndirectDrawMetadata::AddIndexedIndirectDraw(wgpu::IndexFormat indexFormat,
                                                   uint64_t indexBufferSize,
+                                                  uint64_t indexBufferOffset,
                                                   BufferBase* indirectBuffer,
                                                   uint64_t indirectOffset,
                                                   bool duplicateBaseVertexInstance,
                                                   DrawIndexedIndirectCmd* cmd) {
     uint64_t numIndexBufferElements;
+    uint64_t indexBufferOffsetInElements;
     switch (indexFormat) {
         case wgpu::IndexFormat::Uint16:
             numIndexBufferElements = indexBufferSize / 2;
+            indexBufferOffsetInElements = indexBufferOffset / 2;
             break;
         case wgpu::IndexFormat::Uint32:
             numIndexBufferElements = indexBufferSize / 4;
+            indexBufferOffsetInElements = indexBufferOffset / 4;
             break;
         case wgpu::IndexFormat::Undefined:
             DAWN_UNREACHABLE();
@@ -198,6 +202,7 @@
     IndirectDraw draw{};
     draw.inputBufferOffset = indirectOffset;
     draw.numIndexBufferElements = numIndexBufferElements;
+    draw.indexBufferOffsetInElements = indexBufferOffsetInElements;
     draw.cmd = cmd;
     it->second.AddIndirectDraw(mMaxDrawCallsPerBatch, mMaxBatchOffsetRange, draw);
 }
diff --git a/src/dawn/native/IndirectDrawMetadata.h b/src/dawn/native/IndirectDrawMetadata.h
index 5c36f1d..c2966b1 100644
--- a/src/dawn/native/IndirectDrawMetadata.h
+++ b/src/dawn/native/IndirectDrawMetadata.h
@@ -59,6 +59,7 @@
     struct IndirectDraw {
         uint64_t inputBufferOffset;
         uint64_t numIndexBufferElements;
+        uint64_t indexBufferOffsetInElements;
         // This is a pointer to the command that should be populated with the validated
         // indirect scratch buffer. It is only valid up until the encoded command buffer
         // is submitted.
@@ -136,6 +137,7 @@
     void AddBundle(RenderBundleBase* bundle);
     void AddIndexedIndirectDraw(wgpu::IndexFormat indexFormat,
                                 uint64_t indexBufferSize,
+                                uint64_t indexBufferOffset,
                                 BufferBase* indirectBuffer,
                                 uint64_t indirectOffset,
                                 bool duplicateBaseVertexInstance,
diff --git a/src/dawn/native/IndirectDrawValidationEncoder.cpp b/src/dawn/native/IndirectDrawValidationEncoder.cpp
index d86df19..cd68b14 100644
--- a/src/dawn/native/IndirectDrawValidationEncoder.cpp
+++ b/src/dawn/native/IndirectDrawValidationEncoder.cpp
@@ -58,14 +58,16 @@
 constexpr uint32_t kIndexedDraw = 2;
 constexpr uint32_t kValidationEnabled = 4;
 constexpr uint32_t kIndirectFirstInstanceEnabled = 8;
+constexpr uint32_t kUseFirstIndexToEmulateIndexBufferOffset = 16;
 
 // Equivalent to the IndirectDraw struct defined in the shader below.
 struct IndirectDraw {
     uint32_t indirectOffset;
     uint32_t numIndexBufferElementsLow;
     uint32_t numIndexBufferElementsHigh;
+    uint32_t indexOffsetAsNumElements;
 };
-static_assert(sizeof(IndirectDraw) == sizeof(uint32_t) * 3);
+static_assert(sizeof(IndirectDraw) == sizeof(uint32_t) * 4);
 static_assert(alignof(IndirectDraw) == alignof(uint32_t));
 
 // Equivalent to the BatchInfo struct defined in the shader below.
@@ -75,7 +77,7 @@
 };
 
 // The size, in bytes, of the IndirectDraw struct defined in the shader below.
-constexpr uint32_t kIndirectDrawByteSize = sizeof(uint32_t) * 3;
+constexpr uint32_t kIndirectDrawByteSize = sizeof(uint32_t) * 4;
 
 // TODO(https://crbug.com/dawn/1108): Propagate validation feedback from this shader in
 // various failure modes.
@@ -91,11 +93,13 @@
             const kIndexedDraw = 2u;
             const kValidationEnabled = 4u;
             const kIndirectFirstInstanceEnabled = 8u;
+            const kUseFirstIndexToEmulateIndexBufferOffset = 16u;
 
             struct IndirectDraw {
                 indirectOffset: u32,
                 numIndexBufferElementsLow: u32,
                 numIndexBufferElementsHigh: u32,
+                indexOffsetAsNumElements: u32,
             }
 
             struct BatchInfo {
@@ -157,6 +161,10 @@
                 for(var i = 0u; i < numInputParams; i = i + 1u) {
                     outputParams.data[outIndex + i] = inputParams.data[inIndex + i];
                 }
+
+                if (bool(batch.flags & kUseFirstIndexToEmulateIndexBufferOffset)) {
+                    outputParams.data[outIndex + kFirstIndexEntry] += batch.draws[drawIndex].indexOffsetAsNumElements;
+                }
             }
 
             @compute @workgroup_size(64, 1, 1)
@@ -219,6 +227,23 @@
     InternalPipelineStore* store = device->GetInternalPipelineStore();
 
     if (store->renderValidationPipeline == nullptr) {
+        // If we need to apply the index buffer offset to the first index then
+        // we can't handle buffers larger than 4gig otherwise we'll overflow first_index
+        // which is a 32bit value.
+        //
+        // When a buffer is less than 4gig the largest index buffer offset you can pass to
+        // SetIndexBuffer is 0xffff_fffe. Otherwise you'll get a validation error. This
+        // is converted to count of indices and so at most 0x7fff_ffff.
+        //
+        // The largest valid first_index would be 0x7fff_ffff. Anything larger will fail
+        // the validation used in this compute shader and the validated indirect buffer
+        // will have 0,0,0,0,0.
+        //
+        // Adding 0x7fff_ffff + 0x7fff_ffff does not overflow so as long as we keep
+        // maxBufferSize < 4gig we're safe.
+        DAWN_ASSERT(!device->ShouldApplyIndexBufferOffsetToFirstIndex() ||
+                    device->GetLimits().v1.maxBufferSize < 0x100000000u);
+
         // Create compute shader module if not cached before.
         if (store->renderValidationShader == nullptr) {
             DAWN_TRY_ASSIGN(store->renderValidationShader,
@@ -317,6 +342,9 @@
     const uint32_t minStorageBufferOffsetAlignment =
         device->GetLimits().v1.minStorageBufferOffsetAlignment;
 
+    const bool applyIndexBufferOffsetToFirstIndex =
+        device->ShouldApplyIndexBufferOffsetToFirstIndex();
+
     for (auto& [config, validationInfo] : bufferInfoMap) {
         const uint64_t indirectDrawCommandSize =
             config.drawType == IndirectDrawMetadata::DrawType::Indexed ? kDrawIndexedIndirectSize
@@ -376,6 +404,10 @@
             }
             if (config.drawType == IndirectDrawMetadata::DrawType::Indexed) {
                 newPass.flags |= kIndexedDraw;
+
+                if (applyIndexBufferOffsetToFirstIndex) {
+                    newPass.flags |= kUseFirstIndexToEmulateIndexBufferOffset;
+                }
             }
             if (device->IsValidationEnabled()) {
                 newPass.flags |= kValidationEnabled;
@@ -424,6 +456,9 @@
                     static_cast<uint32_t>(draw.numIndexBufferElements & 0xFFFFFFFF);
                 indirectDraw->numIndexBufferElementsHigh =
                     static_cast<uint32_t>((draw.numIndexBufferElements >> 32) & 0xFFFFFFFF);
+
+                // This is only used in the GL backend.
+                indirectDraw->indexOffsetAsNumElements = draw.indexBufferOffsetInElements;
                 indirectDraw++;
 
                 draw.cmd->indirectBuffer = outputParamsBuffer.GetBuffer();
diff --git a/src/dawn/native/RenderEncoderBase.cpp b/src/dawn/native/RenderEncoderBase.cpp
index 75a9f1f..15c749e 100644
--- a/src/dawn/native/RenderEncoderBase.cpp
+++ b/src/dawn/native/RenderEncoderBase.cpp
@@ -279,7 +279,10 @@
             bool duplicateBaseVertexInstance =
                 GetDevice()->ShouldDuplicateParametersForDrawIndirect(
                     mCommandBufferState.GetRenderPipeline());
-            if (IsValidationEnabled() || duplicateBaseVertexInstance) {
+            bool applyIndexBufferOffsetToFirstIndex =
+                GetDevice()->ShouldApplyIndexBufferOffsetToFirstIndex();
+            if (IsValidationEnabled() || duplicateBaseVertexInstance ||
+                applyIndexBufferOffsetToFirstIndex) {
                 // Later, EncodeIndirectDrawValidationCommands will allocate a scratch storage
                 // buffer which will store the validated or duplicated indirect data. The buffer
                 // and offset will be updated to point to it.
@@ -289,7 +292,8 @@
 
                 mIndirectDrawMetadata.AddIndexedIndirectDraw(
                     mCommandBufferState.GetIndexFormat(), mCommandBufferState.GetIndexBufferSize(),
-                    indirectBuffer, indirectOffset, duplicateBaseVertexInstance, cmd);
+                    mCommandBufferState.GetIndexBufferOffset(), indirectBuffer, indirectOffset,
+                    duplicateBaseVertexInstance, cmd);
             } else {
                 cmd->indirectBuffer = indirectBuffer;
                 cmd->indirectOffset = indirectOffset;
@@ -384,7 +388,7 @@
                 }
             }
 
-            mCommandBufferState.SetIndexBuffer(format, size);
+            mCommandBufferState.SetIndexBuffer(format, offset, size);
 
             SetIndexBufferCmd* cmd =
                 allocator->Allocate<SetIndexBufferCmd>(Command::SetIndexBuffer);
diff --git a/src/dawn/native/opengl/DeviceGL.cpp b/src/dawn/native/opengl/DeviceGL.cpp
index 4249f8a..133552f 100644
--- a/src/dawn/native/opengl/DeviceGL.cpp
+++ b/src/dawn/native/opengl/DeviceGL.cpp
@@ -421,6 +421,14 @@
     return 1.0f;
 }
 
+bool Device::MayRequireDuplicationOfIndirectParameters() const {
+    return true;
+}
+
+bool Device::ShouldApplyIndexBufferOffsetToFirstIndex() const {
+    return true;
+}
+
 const OpenGLFunctions& Device::GetGL() const {
     mContext->MakeCurrent();
     ToBackend(GetQueue())->OnGLUsed();
diff --git a/src/dawn/native/opengl/DeviceGL.h b/src/dawn/native/opengl/DeviceGL.h
index 22af159..4f99ec2 100644
--- a/src/dawn/native/opengl/DeviceGL.h
+++ b/src/dawn/native/opengl/DeviceGL.h
@@ -94,6 +94,9 @@
 
     float GetTimestampPeriodInNS() const override;
 
+    bool MayRequireDuplicationOfIndirectParameters() const override;
+    bool ShouldApplyIndexBufferOffsetToFirstIndex() const override;
+
     class Context {
       public:
         virtual ~Context() {}
diff --git a/src/dawn/native/opengl/PhysicalDeviceGL.cpp b/src/dawn/native/opengl/PhysicalDeviceGL.cpp
index 699ad03..c92b41b 100644
--- a/src/dawn/native/opengl/PhysicalDeviceGL.cpp
+++ b/src/dawn/native/opengl/PhysicalDeviceGL.cpp
@@ -292,6 +292,11 @@
     limits->v1.minStorageBufferOffsetAlignment = Get(gl, GL_SHADER_STORAGE_BUFFER_OFFSET_ALIGNMENT);
     limits->v1.maxVertexBuffers = Get(gl, GL_MAX_VERTEX_ATTRIB_BINDINGS);
     limits->v1.maxBufferSize = kAssumedMaxBufferSize;
+    // The code that handles adding the index buffer offset to first_index
+    // used in drawIndexedIndirect can not handle a max buffer size larger than 4gig.
+    // See IndirectDrawValidationEncoder.cpp
+    static_assert(kAssumedMaxBufferSize < 0x100000000u);
+
     limits->v1.maxVertexAttributes = Get(gl, GL_MAX_VERTEX_ATTRIBS);
     limits->v1.maxVertexBufferArrayStride = Get(gl, GL_MAX_VERTEX_ATTRIB_STRIDE);
     limits->v1.maxInterStageShaderComponents = Get(gl, GL_MAX_VARYING_COMPONENTS);
diff --git a/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp b/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp
index 333f2d7..d461d69 100644
--- a/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp
+++ b/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp
@@ -169,11 +169,6 @@
 
 // Test the parameter 'baseVertex' of DrawIndexed() works.
 TEST_P(DrawIndexedIndirectTest, BaseVertex) {
-    // TODO(crbug.com/dawn/161): add workaround for OpenGL index buffer offset (could be compute
-    // shader that adds it to the draw calls)
-    DAWN_TEST_UNSUPPORTED_IF(IsOpenGL());
-    DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
-
     // TODO(crbug.com/dawn/1292): Some Intel OpenGL drivers don't seem to like
     // the offsets that Tint/GLSL produces.
     DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL() && IsLinux());
@@ -197,6 +192,10 @@
 
     // Test a draw with only the last 3 indices of the first quad (top right triangle)
     Test({3, 1, 3, unsignedNegFour, 0}, 6 * sizeof(uint32_t), 0, notFilled, filled);
+
+    // Test a draw with only the last 3 indices of the first quad (top right triangle) and offset
+    Test({0, 3, 1, 3, unsignedNegFour, 0}, 6 * sizeof(uint32_t), 1 * sizeof(uint32_t), notFilled,
+         filled);
 }
 
 TEST_P(DrawIndexedIndirectTest, IndirectOffset) {
@@ -248,10 +247,6 @@
 }
 
 TEST_P(DrawIndexedIndirectTest, ValidateWithOffsets) {
-    // TODO(crbug.com/dawn/161): The GL/GLES backend doesn't support indirect index buffer offsets
-    // yet.
-    DAWN_SUPPRESS_TEST_IF(IsOpenGL() || IsOpenGLES());
-
     // TODO(crbug.com/dawn/1292): Some Intel OpenGL drivers don't seem to like
     // the offsets that Tint/GLSL produces.
     DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL() && IsLinux());