Compat: Handle index buffer offset with indexed indirect draw.
The validation is already copying and updating all of the
indirect draw buffers so this just adds the adjustment to
update the firstIndex parameter when copying.
Validation would always be on for WebGPU in the browser.
Emits an error if validation is off and the index buffer
offset is non-zero.
Bug: dawn:161
Change-Id: I385f5a8907edd85a9c44652d495b11dd4710b5a8
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/174400
Reviewed-by: Austin Eng <enga@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Stephen White <senorblanco@chromium.org>
Commit-Queue: Gregg Tavares <gman@chromium.org>
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());
