Add D3D12 Backend for MultiDrawIndirect Feature
This CL adds support for the D3D12 backend, which requires emulation of
base vertex and base instance parameters in the indirect arguments.
The indirect draw validation compute pass duplicates the parameters if it is used in the shaders.
MultiDraw capability is supported on all devices with D3D12.
New tests added for MultiDrawIndexedIndirect to test baseVertex and firstInstance.
Change-Id: I75bc48243e4801f49e6e50091cf1560593a3d14c
Bug: 356461286
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/203254
Reviewed-by: Loko Kung <lokokung@google.com>
Commit-Queue: Srijan Dhungana <srijan.dhungana6@gmail.com>
Reviewed-by: Austin Eng <enga@chromium.org>
diff --git a/src/dawn/native/IndirectDrawMetadata.cpp b/src/dawn/native/IndirectDrawMetadata.cpp
index 04f652d..3a65571 100644
--- a/src/dawn/native/IndirectDrawMetadata.cpp
+++ b/src/dawn/native/IndirectDrawMetadata.cpp
@@ -242,10 +242,12 @@
mIndexedIndirectBufferValidationInfo.clear();
}
-void IndirectDrawMetadata::AddMultiDrawIndirect(MultiDrawIndirectCmd* cmd) {
+void IndirectDrawMetadata::AddMultiDrawIndirect(bool duplicateBaseVertexInstance,
+ MultiDrawIndirectCmd* cmd) {
IndirectMultiDraw multiDraw;
multiDraw.type = DrawType::NonIndexed;
multiDraw.cmd = cmd;
+ multiDraw.duplicateBaseVertexInstance = duplicateBaseVertexInstance;
mMultiDraws.push_back(multiDraw);
}
@@ -253,12 +255,14 @@
wgpu::IndexFormat indexFormat,
uint64_t indexBufferSize,
uint64_t indexBufferOffset,
+ bool duplicateBaseVertexInstance,
MultiDrawIndexedIndirectCmd* cmd) {
IndirectMultiDraw multiDraw;
multiDraw.type = DrawType::Indexed;
multiDraw.cmd = cmd;
multiDraw.indexBufferSize = indexBufferSize;
multiDraw.indexFormat = indexFormat;
+ multiDraw.duplicateBaseVertexInstance = duplicateBaseVertexInstance;
mMultiDraws.push_back(multiDraw);
}
diff --git a/src/dawn/native/IndirectDrawMetadata.h b/src/dawn/native/IndirectDrawMetadata.h
index 657dd08..ad42513 100644
--- a/src/dawn/native/IndirectDrawMetadata.h
+++ b/src/dawn/native/IndirectDrawMetadata.h
@@ -82,8 +82,9 @@
struct IndirectMultiDraw {
DrawType type;
- uint64_t indexBufferSize;
- wgpu::IndexFormat indexFormat;
+ uint64_t indexBufferSize = 0;
+ wgpu::IndexFormat indexFormat = wgpu::IndexFormat::Undefined;
+ bool duplicateBaseVertexInstance;
// When validation is enabled, the original indirect buffer is validated and copied to a new
// indirect buffer containing only valid commands. The pointer to the command allocated in
@@ -166,12 +167,13 @@
bool duplicateBaseVertexInstance,
DrawIndirectCmd* cmd);
- void AddMultiDrawIndirect(MultiDrawIndirectCmd* cmd);
+ void AddMultiDrawIndirect(bool duplicateBaseVertexInstance, MultiDrawIndirectCmd* cmd);
void AddMultiDrawIndexedIndirect(BufferBase* indexBuffer,
wgpu::IndexFormat indexFormat,
uint64_t indexBufferSize,
uint64_t indexBufferOffset,
+ bool duplicateBaseVertexInstance,
MultiDrawIndexedIndirectCmd* cmd);
void ClearIndexedIndirectBufferValidationInfo();
diff --git a/src/dawn/native/IndirectDrawValidationEncoder.cpp b/src/dawn/native/IndirectDrawValidationEncoder.cpp
index 182fbc6..b95cea8 100644
--- a/src/dawn/native/IndirectDrawValidationEncoder.cpp
+++ b/src/dawn/native/IndirectDrawValidationEncoder.cpp
@@ -204,6 +204,15 @@
let inIndex = drawIndex * numInputParams;
let inputOffset = drawConstants.indirectOffsetInElements;
+ if (bool(drawConstants.flags & kDuplicateBaseVertexInstance)) {
+ // first/baseVertex and firstInstance are always last two parameters
+ let dupIndex = inputOffset + inIndex + numInputParams - 2u;
+ outputParams.data[outIndex] = inputParams.data[dupIndex];
+ outputParams.data[outIndex + 1u] = inputParams.data[dupIndex + 1u];
+
+ outIndex = outIndex + 2u;
+ }
+
for(var i = 0u; i < numInputParams; i = i + 1u) {
outputParams.data[outIndex + i] = inputParams.data[inputOffset + inIndex + i];
}
@@ -266,9 +275,7 @@
@compute @workgroup_size(kWorkgroupSize, 1, 1)
fn validate_multi_draw(@builtin(global_invocation_id) id : vec3u) {
-
var drawCount = drawConstants.maxDrawCount;
-
var drawCountOffset = drawConstants.drawCountOffsetInElements;
if(bool(drawConstants.flags & kIndirectDrawCountBuffer)) {
@@ -280,6 +287,16 @@
return;
}
+ if(!bool(drawConstants.flags & kValidationEnabled)) {
+ set_pass_multi(id.x);
+ return;
+ }
+
+ if (!bool(drawConstants.flags & kIndexedDraw)) {
+ set_pass_multi(id.x);
+ return;
+ }
+
let numIndexBufferElementsHigh = drawConstants.numIndexBufferElementsHigh;
if (numIndexBufferElementsHigh >= 2u) {
@@ -313,6 +330,17 @@
)";
+static constexpr uint32_t GetOutputIndirectDrawSize(IndirectDrawMetadata::DrawType drawType,
+ bool duplicateBaseVertexInstance) {
+ uint32_t drawSize = drawType == IndirectDrawMetadata::DrawType::Indexed
+ ? kDrawIndexedIndirectSize
+ : kDrawIndirectSize;
+ if (duplicateBaseVertexInstance) {
+ drawSize += 2 * sizeof(uint32_t);
+ }
+ return drawSize;
+}
+
ResultOrError<dawn::Ref<ComputePipelineBase>> CreateRenderValidationPipelines(
DeviceBase* device,
const char* entryPoint,
@@ -473,10 +501,8 @@
config.drawType == IndirectDrawMetadata::DrawType::Indexed ? kDrawIndexedIndirectSize
: kDrawIndirectSize;
- uint64_t outputIndirectSize = indirectDrawCommandSize;
- if (config.duplicateBaseVertexInstance) {
- outputIndirectSize += 2 * sizeof(uint32_t);
- }
+ uint64_t outputIndirectSize =
+ GetOutputIndirectDrawSize(config.drawType, config.duplicateBaseVertexInstance);
for (const IndirectDrawMetadata::IndirectValidationBatch& batch :
validationInfo.GetBatches()) {
@@ -552,11 +578,22 @@
uint64_t outputParamsSizeForMultiDraw = 0;
// Calculate size of output params for multi draws
for (auto& draw : multiDraws) {
- // Don't need to validate non-indexed draws.
- if (draw.type == IndirectDrawMetadata::DrawType::NonIndexed) {
+ // Multi draw metadatas are added even if validation is disabled, because the Metal backend
+ // needs to convert all multi draws into an ICB. If validation is disabled, and the draw
+ // doesn't need duplication of base vertex and instance, we can skip the compute pass.
+ // In general, non-indexed multi draws don't need validation.
+ if ((draw.type == IndirectDrawMetadata::DrawType::NonIndexed ||
+ !device->IsValidationEnabled()) &&
+ !draw.duplicateBaseVertexInstance) {
continue;
}
- outputParamsSizeForMultiDraw += draw.cmd->maxDrawCount * kDrawIndexedIndirectSize;
+
+ outputParamsSizeForMultiDraw +=
+ draw.cmd->maxDrawCount *
+ GetOutputIndirectDrawSize(draw.type, draw.duplicateBaseVertexInstance);
+
+ outputParamsSizeForMultiDraw =
+ Align(outputParamsSizeForMultiDraw, minStorageBufferOffsetAlignment);
if (outputParamsSizeForMultiDraw > maxStorageBufferBindingSize) {
return DAWN_INTERNAL_ERROR("Too many multiDrawIndexedIndirect calls to validate");
@@ -580,7 +617,8 @@
for (const Pass& pass : passes) {
requiredBatchDataBufferSize = std::max(requiredBatchDataBufferSize, pass.batchDataSize);
}
- // Needs to at least be able to store a MultiDrawConstants struct for the multi draw validation.
+ // Needs to at least be able to store a MultiDrawConstants struct for the multi draw
+ // validation.
requiredBatchDataBufferSize =
std::max(requiredBatchDataBufferSize, static_cast<uint64_t>(sizeof(MultiDrawConstants)));
@@ -656,10 +694,10 @@
bindGroupDescriptor.entryCount = 3;
bindGroupDescriptor.entries = bindings;
- // Finally, we can now encode our validation and duplication passes. Each pass first does
- // two WriteBuffer to get batch and pass data over to the GPU, followed by a single compute
- // pass. The compute pass encodes a separate SetBindGroup and Dispatch command for each
- // batch.
+ // Finally, we can now encode our validation and duplication passes. Each pass first
+ // does a WriteBuffer to get batch and pass data over to the GPU, followed by a single
+ // compute pass. The compute pass encodes a separate SetBindGroup and Dispatch command
+ // for each batch.
for (const Pass& pass : passes) {
commandEncoder->APIWriteBuffer(batchDataBuffer.GetBuffer(), 0,
static_cast<const uint8_t*>(pass.batchData.get()),
@@ -724,12 +762,19 @@
uint64_t outputOffset = multiDrawOutputParamsOffset;
for (auto& draw : multiDraws) {
- if (draw.type == IndirectDrawMetadata::DrawType::NonIndexed) {
+ // If the draw meets these conditions, there is no need to run the compute pass,
+ // and there is no space allocated for the output params
+ if ((draw.type == IndirectDrawMetadata::DrawType::NonIndexed ||
+ !device->IsValidationEnabled()) &&
+ !draw.duplicateBaseVertexInstance) {
continue;
}
- const size_t formatSize = IndexFormatSize(draw.indexFormat);
- uint64_t numIndexBufferElements = draw.indexBufferSize / formatSize;
+ uint64_t numIndexBufferElements = 0;
+ if (draw.type == IndirectDrawMetadata::DrawType::Indexed) {
+ const size_t formatSize = IndexFormatSize(draw.indexFormat);
+ numIndexBufferElements = draw.indexBufferSize / formatSize;
+ }
// Same struct for both indexed and non-indexed draws.
MultiDrawIndirectCmd* cmd = draw.cmd;
@@ -748,10 +793,20 @@
static_cast<uint32_t>(numIndexBufferElements & 0xFFFFFFFF);
drawConstants.numIndexBufferElementsHigh =
static_cast<uint32_t>((numIndexBufferElements >> 32) & 0xFFFFFFFF);
- drawConstants.flags = kIndexedDraw;
+
+ drawConstants.flags = 0;
+ if (device->IsValidationEnabled()) {
+ drawConstants.flags |= kValidationEnabled;
+ }
+ if (draw.type == IndirectDrawMetadata::DrawType::Indexed) {
+ drawConstants.flags |= kIndexedDraw;
+ }
if (cmd->drawCountBuffer != nullptr) {
drawConstants.flags |= kIndirectDrawCountBuffer;
}
+ if (draw.duplicateBaseVertexInstance) {
+ drawConstants.flags |= kDuplicateBaseVertexInstance;
+ }
inputIndirectBinding.buffer = cmd->indirectBuffer.Get();
// We can't use the offset directly because the indirect offset is guaranteed to
@@ -763,19 +818,23 @@
outputParamsBinding.buffer = outputParamsBuffer.GetBuffer();
outputParamsBinding.offset = outputOffset;
+ outputParamsBinding.size =
+ draw.cmd->maxDrawCount *
+ GetOutputIndirectDrawSize(draw.type, draw.duplicateBaseVertexInstance);
if (cmd->drawCountBuffer != nullptr) {
// If the drawCountBuffer is set, we need to bind it to the bind group.
// The drawCountBuffer is used to read the drawCount for the multi draw call.
- // If the drawCount exceeds the maxDrawCount, it will be clamped to maxDrawCount.
+ // If the drawCount exceeds the maxDrawCount, it will be clamped to
+ // maxDrawCount.
drawCountBinding.buffer = cmd->drawCountBuffer.Get();
drawCountBinding.offset =
AlignDown(cmd->drawCountOffset, minStorageBufferOffsetAlignment);
} else {
// This is an unused binding.
- // Bind group entry for the drawCountBuffer is not needed however we need to bind
- // something else than nullptr to the bind group entry to avoid validation errors.
- // This buffer is never used in the shader, since there is a flag
+ // Bind group entry for the drawCountBuffer is not needed however we need to
+ // bind something else than nullptr to the bind group entry to avoid validation
+ // errors. This buffer is never used in the shader, since there is a flag
// (kIndirectDrawCountBuffer) to check if the drawCountBuffer is set.
drawCountBinding.buffer = cmd->indirectBuffer.Get();
drawCountBinding.offset = 0;
@@ -792,9 +851,6 @@
passEncoder->APISetPipeline(pipeline);
passEncoder->APISetBindGroup(0, bindGroup.Get());
- // TODO(crbug.com/356461286): After maxDrawCount has a limit we can
- // dispatch exact number of workgroups without worrying about overflow:
- // uint32_t workgroupCount = (cmd->maxDrawCount + kWorkgroupSize - 1u) / kWorkgroupSize;
uint32_t workgroupCount = cmd->maxDrawCount / kWorkgroupSize;
// Integer division rounds down so adding 1 if there is a remainder.
workgroupCount += cmd->maxDrawCount % kWorkgroupSize == 0 ? 0 : 1;
@@ -802,12 +858,15 @@
passEncoder->APIEnd();
// Update the draw command to use the validated indirect buffer.
- // The drawCountBuffer doesn't need to be updated because if it exceeds the maxDrawCount
- // it will be clamped to maxDrawCount.
+ // The drawCountBuffer doesn't need to be updated because if it exceeds the
+ // maxDrawCount it will be clamped to maxDrawCount.
cmd->indirectBuffer = outputParamsBuffer.GetBuffer();
cmd->indirectOffset = outputOffset;
- outputOffset += cmd->maxDrawCount * kDrawIndexedIndirectSize;
+ // Proceed to the next output offset.
+ outputOffset += cmd->maxDrawCount *
+ GetOutputIndirectDrawSize(draw.type, draw.duplicateBaseVertexInstance);
+ outputOffset = Align(outputOffset, minStorageBufferOffsetAlignment);
}
}
diff --git a/src/dawn/native/RenderEncoderBase.cpp b/src/dawn/native/RenderEncoderBase.cpp
index 1995db6..acc8e93 100644
--- a/src/dawn/native/RenderEncoderBase.cpp
+++ b/src/dawn/native/RenderEncoderBase.cpp
@@ -381,7 +381,11 @@
cmd->drawCountBuffer = drawCountBuffer;
cmd->drawCountOffset = drawCountBufferOffset;
- mIndirectDrawMetadata.AddMultiDrawIndirect(cmd);
+ bool duplicateBaseVertexInstance =
+ GetDevice()->ShouldDuplicateParametersForDrawIndirect(
+ mCommandBufferState.GetRenderPipeline());
+
+ mIndirectDrawMetadata.AddMultiDrawIndirect(duplicateBaseVertexInstance, cmd);
// TODO(crbug.com/dawn/1166): Adding the indirectBuffer is needed for correct usage
// validation, but it will unecessarily transition to indirectBuffer usage in the
@@ -470,10 +474,14 @@
cmd->drawCountBuffer = drawCountBuffer;
cmd->drawCountOffset = drawCountBufferOffset;
+ bool duplicateBaseVertexInstance =
+ GetDevice()->ShouldDuplicateParametersForDrawIndirect(
+ mCommandBufferState.GetRenderPipeline());
+
mIndirectDrawMetadata.AddMultiDrawIndexedIndirect(
mCommandBufferState.GetIndexBuffer(), mCommandBufferState.GetIndexFormat(),
mCommandBufferState.GetIndexBufferSize(),
- mCommandBufferState.GetIndexBufferOffset(), cmd);
+ mCommandBufferState.GetIndexBufferOffset(), duplicateBaseVertexInstance, cmd);
// TODO(crbug.com/dawn/1166): Adding the indirectBuffer is needed for correct usage
// validation, but it will unecessarily transition to indirectBuffer usage in the
diff --git a/src/dawn/native/d3d12/CommandBufferD3D12.cpp b/src/dawn/native/d3d12/CommandBufferD3D12.cpp
index 78dc3a2..4098468 100644
--- a/src/dawn/native/d3d12/CommandBufferD3D12.cpp
+++ b/src/dawn/native/d3d12/CommandBufferD3D12.cpp
@@ -1594,6 +1594,56 @@
break;
}
+ case Command::MultiDrawIndirect: {
+ MultiDrawIndirectCmd* draw = iter->NextCommand<MultiDrawIndirectCmd>();
+
+ DAWN_TRY(bindingTracker->Apply(commandContext));
+ vertexBufferTracker.Apply(commandList, lastPipeline);
+
+ Buffer* indirectBuffer = ToBackend(draw->indirectBuffer.Get());
+ DAWN_ASSERT(indirectBuffer != nullptr);
+
+ Buffer* countBuffer = ToBackend(draw->drawCountBuffer.Get());
+
+ // There is no distinction between DrawIndirect and MultiDrawIndirect in D3D12.
+ // This is why we can use the same command signature for both.
+ ComPtr<ID3D12CommandSignature> signature =
+ lastPipeline->GetDrawIndirectCommandSignature();
+
+ commandList->ExecuteIndirect(
+ signature.Get(), draw->maxDrawCount, indirectBuffer->GetD3D12Resource(),
+ draw->indirectOffset,
+ countBuffer != nullptr ? countBuffer->GetD3D12Resource() : nullptr,
+ countBuffer != nullptr ? draw->drawCountOffset : 0);
+
+ break;
+ }
+
+ case Command::MultiDrawIndexedIndirect: {
+ MultiDrawIndexedIndirectCmd* draw =
+ iter->NextCommand<MultiDrawIndexedIndirectCmd>();
+
+ DAWN_TRY(bindingTracker->Apply(commandContext));
+ vertexBufferTracker.Apply(commandList, lastPipeline);
+
+ Buffer* indirectBuffer = ToBackend(draw->indirectBuffer.Get());
+ DAWN_ASSERT(indirectBuffer != nullptr);
+
+ Buffer* countBuffer = ToBackend(draw->drawCountBuffer.Get());
+
+ // There is no distinction between DrawIndexedIndirect and MultiDrawIndexedIndirect
+ // in D3D12. This is why we can use the same command signature for both.
+ ComPtr<ID3D12CommandSignature> signature =
+ lastPipeline->GetDrawIndexedIndirectCommandSignature();
+
+ commandList->ExecuteIndirect(
+ signature.Get(), draw->maxDrawCount, indirectBuffer->GetD3D12Resource(),
+ draw->indirectOffset,
+ countBuffer != nullptr ? countBuffer->GetD3D12Resource() : nullptr,
+ countBuffer != nullptr ? draw->drawCountOffset : 0);
+ break;
+ }
+
case Command::InsertDebugMarker: {
InsertDebugMarkerCmd* cmd = iter->NextCommand<InsertDebugMarkerCmd>();
const char* label = iter->NextData<char>(cmd->length + 1);
diff --git a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
index eeb5b66..64231d6 100644
--- a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
+++ b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
@@ -146,6 +146,7 @@
EnableFeature(Feature::SharedBufferMemoryD3D12Resource);
EnableFeature(Feature::ShaderModuleCompilationOptions);
EnableFeature(Feature::StaticSamplers);
+ EnableFeature(Feature::MultiDrawIndirect);
if (AreTimestampQueriesSupported()) {
EnableFeature(Feature::TimestampQuery);
diff --git a/src/dawn/tests/end2end/MultiDrawIndexedIndirectTests.cpp b/src/dawn/tests/end2end/MultiDrawIndexedIndirectTests.cpp
index 50ebc68..c625566 100644
--- a/src/dawn/tests/end2end/MultiDrawIndexedIndirectTests.cpp
+++ b/src/dawn/tests/end2end/MultiDrawIndexedIndirectTests.cpp
@@ -430,7 +430,188 @@
EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 3, 1);
}
-DAWN_INSTANTIATE_TEST(MultiDrawIndexedIndirectTest, VulkanBackend());
+DAWN_INSTANTIATE_TEST(MultiDrawIndexedIndirectTest, VulkanBackend(), D3D12Backend());
+
+class MultiDrawIndexedIndirectUsingFirstVertexTest : public DawnTest {
+ protected:
+ std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
+ if (!SupportsFeatures({wgpu::FeatureName::MultiDrawIndirect})) {
+ return {};
+ }
+ return {wgpu::FeatureName::MultiDrawIndirect};
+ }
+ virtual void SetupShaderModule() {
+ vsModule = utils::CreateShaderModule(device, R"(
+ struct VertexInput {
+ @builtin(vertex_index) id : u32,
+ @location(0) pos: vec4f,
+ };
+ @group(0) @binding(0) var<uniform> offset: array<vec4f, 2>;
+ @vertex
+ fn main(input: VertexInput) -> @builtin(position) vec4f {
+ return input.pos + offset[input.id / 3u];
+ })");
+ fsModule = utils::CreateShaderModule(device, R"(
+ @fragment fn main() -> @location(0) vec4f {
+ return vec4f(0.0, 1.0, 0.0, 1.0);
+ })");
+ }
+ void GeneralSetup() {
+ renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize);
+ SetupShaderModule();
+ utils::ComboRenderPipelineDescriptor descriptor;
+ descriptor.vertex.module = vsModule;
+ descriptor.cFragment.module = fsModule;
+ descriptor.primitive.topology = wgpu::PrimitiveTopology::TriangleStrip;
+ descriptor.primitive.stripIndexFormat = wgpu::IndexFormat::Uint32;
+ descriptor.vertex.bufferCount = 1;
+ descriptor.cBuffers[0].arrayStride = 4 * sizeof(float);
+ descriptor.cBuffers[0].attributeCount = 1;
+ descriptor.cAttributes[0].format = wgpu::VertexFormat::Float32x4;
+ descriptor.cTargets[0].format = renderPass.colorFormat;
+
+ pipeline = device.CreateRenderPipeline(&descriptor);
+
+ // Offset to the vertices, that needs correcting by the calibration offset from uniform
+ // buffer referenced by instance index to get filled triangle on screen.
+ constexpr float calibration = 99.0f;
+ vertexBuffer = dawn::utils::CreateBufferFromData<float>(
+ device, wgpu::BufferUsage::Vertex,
+ {// The bottom left triangle
+ -1.0f - calibration, 1.0f, 0.0f, 1.0f, 1.0f - calibration, -1.0f, 0.0f, 1.0f,
+ -1.0f - calibration, -1.0f, 0.0f, 1.0f,
+ // The top right triangle
+ -1.0f - calibration, 1.0f, 0.0f, 1.0f, 1.0f - calibration, -1.0f, 0.0f, 1.0f,
+ 1.0f - calibration, 1.0f, 0.0f, 1.0f});
+
+ indexBuffer = dawn::utils::CreateBufferFromData<uint32_t>(device, wgpu::BufferUsage::Index,
+ {0, 1, 2});
+
+ // Providing calibration vec4f offset values
+ wgpu::Buffer uniformBuffer =
+ utils::CreateBufferFromData<float>(device, wgpu::BufferUsage::Uniform,
+ {
+ // Bad calibration at [0]
+ 0.0,
+ 0.0,
+ 0.0,
+ 0.0,
+ // Good calibration at [1]
+ calibration,
+ 0.0,
+ 0.0,
+ 0.0,
+ });
+
+ bindGroup =
+ utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, uniformBuffer}});
+ }
+ void SetUp() override {
+ DawnTest::SetUp();
+ DAWN_TEST_UNSUPPORTED_IF(!device.HasFeature(wgpu::FeatureName::MultiDrawIndirect));
+ GeneralSetup();
+ }
+ utils::BasicRenderPass renderPass;
+ wgpu::RenderPipeline pipeline;
+ wgpu::Buffer vertexBuffer;
+ wgpu::Buffer indexBuffer;
+ wgpu::BindGroup bindGroup;
+ wgpu::ShaderModule vsModule;
+ wgpu::ShaderModule fsModule;
+ // Test two DrawIndirect calls with different indirect offsets within one pass.
+ void Test(std::initializer_list<uint32_t> bufferList,
+ uint32_t maxDrawCount,
+ utils::RGBA8 bottomLeftExpected,
+ utils::RGBA8 topRightExpected) {
+ wgpu::Buffer indirectBuffer =
+ utils::CreateBufferFromData<uint32_t>(device, wgpu::BufferUsage::Indirect, bufferList);
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ {
+ wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo);
+ pass.SetPipeline(pipeline);
+ pass.SetVertexBuffer(0, vertexBuffer);
+ pass.SetBindGroup(0, bindGroup);
+ pass.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0);
+ pass.MultiDrawIndexedIndirect(indirectBuffer, 0, maxDrawCount, nullptr, 0);
+ pass.End();
+ }
+ wgpu::CommandBuffer commands = encoder.Finish();
+ queue.Submit(1, &commands);
+ EXPECT_PIXEL_RGBA8_EQ(bottomLeftExpected, renderPass.color, 1, 3);
+ EXPECT_PIXEL_RGBA8_EQ(topRightExpected, renderPass.color, 3, 1);
+ }
+};
+
+TEST_P(MultiDrawIndexedIndirectUsingFirstVertexTest, IndirectOffset) {
+ utils::RGBA8 filled(0, 255, 0, 255);
+ utils::RGBA8 notFilled(0, 0, 0, 0);
+
+ // Test an offset draw call, with indirect buffer containing 2 calls:
+ // 1) only the first 3 indices (bottom left triangle)
+ // 2) only the last 3 indices (top right triangle)
+ // #2 draw has the correct offset applied by vertex index
+ Test({3, 1, 0, 0, 0, 3, 1, 0, 3, 0}, 2, notFilled, filled);
+}
+
+DAWN_INSTANTIATE_TEST(MultiDrawIndexedIndirectUsingFirstVertexTest,
+ VulkanBackend(),
+ D3D12Backend());
+
+class MultiDrawIndexedIndirectUsingInstanceIndexTest
+ : public MultiDrawIndexedIndirectUsingFirstVertexTest {
+ protected:
+ std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
+ if (!SupportsFeatures({wgpu::FeatureName::MultiDrawIndirect})) {
+ return {};
+ }
+ return {wgpu::FeatureName::MultiDrawIndirect};
+ }
+
+ void SetupShaderModule() override {
+ vsModule = utils::CreateShaderModule(device, R"(
+ struct VertexInput {
+ @builtin(instance_index) id : u32,
+ @location(0) pos: vec4f,
+ };
+
+ @group(0) @binding(0) var<uniform> offset: array<vec4f, 2>;
+
+ @vertex
+ fn main(input: VertexInput) -> @builtin(position) vec4f {
+ return input.pos + offset[input.id];
+ })");
+
+ fsModule = utils::CreateShaderModule(device, R"(
+ @fragment fn main() -> @location(0) vec4f {
+ return vec4f(0.0, 1.0, 0.0, 1.0);
+ })");
+ }
+
+ void SetUp() override {
+ DawnTest::SetUp();
+ DAWN_TEST_UNSUPPORTED_IF(!device.HasFeature(wgpu::FeatureName::MultiDrawIndirect));
+ GeneralSetup();
+ }
+};
+
+TEST_P(MultiDrawIndexedIndirectUsingInstanceIndexTest, IndirectOffset) {
+ utils::RGBA8 filled(0, 255, 0, 255);
+ utils::RGBA8 notFilled(0, 0, 0, 0);
+
+ // Test an offset draw call, with indirect buffer containing 2 calls:
+ // 1) only the first 3 indices (bottom left triangle)
+ // 2) only the last 3 indices (top right triangle)
+
+ // Test 1: #1 draw has the correct calibration referenced by instance index
+ Test({3, 1, 0, 0, 1, 3, 1, 0, 3, 0}, 2, filled, notFilled);
+
+ // Test 2: #2 draw has the correct offset applied by instance index
+ Test({3, 1, 0, 0, 0, 3, 1, 0, 3, 1}, 2, notFilled, filled);
+}
+
+DAWN_INSTANTIATE_TEST(MultiDrawIndexedIndirectUsingInstanceIndexTest,
+ VulkanBackend(),
+ D3D12Backend());
} // anonymous namespace
} // namespace dawn
diff --git a/src/dawn/tests/end2end/MultiDrawIndirectTests.cpp b/src/dawn/tests/end2end/MultiDrawIndirectTests.cpp
index b127fc7..8c7feb1 100644
--- a/src/dawn/tests/end2end/MultiDrawIndirectTests.cpp
+++ b/src/dawn/tests/end2end/MultiDrawIndirectTests.cpp
@@ -167,7 +167,7 @@
Test({3, 1, 0, 0, 3, 1, 3, 0}, kDrawIndirectSize, 1, notFilled, filled);
}
-DAWN_INSTANTIATE_TEST(MultiDrawIndirectTest, VulkanBackend());
+DAWN_INSTANTIATE_TEST(MultiDrawIndirectTest, VulkanBackend(), D3D12Backend());
class MultiDrawIndirectUsingFirstVertexTest : public DawnTest {
protected:
@@ -280,7 +280,7 @@
Test({3, 1, 0, 0, 3, 1, 3, 0}, 2, notFilled, filled);
}
-DAWN_INSTANTIATE_TEST(MultiDrawIndirectUsingFirstVertexTest, VulkanBackend());
+DAWN_INSTANTIATE_TEST(MultiDrawIndirectUsingFirstVertexTest, VulkanBackend(), D3D12Backend());
class MultiDrawIndirectUsingInstanceIndexTest : public MultiDrawIndirectUsingFirstVertexTest {
protected:
@@ -330,7 +330,7 @@
Test({3, 1, 0, 0, 3, 1, 3, 1}, 2, notFilled, filled);
}
-DAWN_INSTANTIATE_TEST(MultiDrawIndirectUsingInstanceIndexTest, VulkanBackend());
+DAWN_INSTANTIATE_TEST(MultiDrawIndirectUsingInstanceIndexTest, VulkanBackend(), D3D12Backend());
} // namespace
} // namespace dawn
