encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

D3D12: use the WGSL binding number as shader register 

Currently, the bind group layout tightly packs the shader registers for
each of the sampler and non-sampler descriptors sets. This reduces the
max shader register used and helps targeting shader model 5.0, which has
a relatively low max slot count per resource. It is safe in D3D, since a
shader register collision can be valid if the descriptor types differ.

To support Mesa's SPIR-V to DXIL compiler, we need to avoid possible
collisions between resources' shader registers because it uses SPIR-V as
an intermediate representation (which does not namespace bindings by
type). This change re-works BindGroupLayoutD3D12 to not assume the
resulting shader registers will be tightly packed and group per
descriptor type.

Change-Id: I0bb51106c4683bfe02ce15ecad71716734b7a91f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/60764
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Michael Tang <tangm@microsoft.com>
This commit is contained in:
Michael Tang 2021-08-19 17:17:26 +00:00 committed by Dawn LUCI CQ
parent 05b770ae86
commit 663a16e2fe
6 changed files with 222 additions and 131 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
src/dawn_native/d3d12/BindGroupD3D12.cpp
View File

@ -40,7 +40,7 @@ namespace dawn_native { namespace d3d12 {
mCPUViewAllocation = viewAllocation;
const auto& bindingOffsets = bgl->GetBindingOffsets();
const auto& descriptorHeapOffsets = bgl->GetDescriptorHeapOffsets();
ID3D12Device* d3d12Device = device->GetD3D12Device();
@ -74,8 +74,8 @@ namespace dawn_native { namespace d3d12 {
ToBackend(binding.buffer)->GetVA() + binding.offset;
d3d12Device->CreateConstantBufferView(
&desc, viewAllocation.OffsetFrom(viewSizeIncrement,
bindingOffsets[bindingIndex]));
&desc, viewAllocation.OffsetFrom(
viewSizeIncrement, descriptorHeapOffsets[bindingIndex]));
break;
}
case wgpu::BufferBindingType::Storage:
@ -99,7 +99,7 @@ namespace dawn_native { namespace d3d12 {
d3d12Device->CreateUnorderedAccessView(
resource, nullptr, &desc,
viewAllocation.OffsetFrom(viewSizeIncrement,
bindingOffsets[bindingIndex]));
descriptorHeapOffsets[bindingIndex]));
break;
}
case wgpu::BufferBindingType::ReadOnlyStorage: {
@ -118,7 +118,7 @@ namespace dawn_native { namespace d3d12 {
d3d12Device->CreateShaderResourceView(
resource, &desc,
viewAllocation.OffsetFrom(viewSizeIncrement,
bindingOffsets[bindingIndex]));
descriptorHeapOffsets[bindingIndex]));
break;
}
case wgpu::BufferBindingType::Undefined:
@ -142,7 +142,8 @@ namespace dawn_native { namespace d3d12 {
d3d12Device->CreateShaderResourceView(
resource, &srv,
viewAllocation.OffsetFrom(viewSizeIncrement, bindingOffsets[bindingIndex]));
viewAllocation.OffsetFrom(viewSizeIncrement,
descriptorHeapOffsets[bindingIndex]));
break;
}
@ -165,7 +166,7 @@ namespace dawn_native { namespace d3d12 {
d3d12Device->CreateShaderResourceView(
resource, &srv,
viewAllocation.OffsetFrom(viewSizeIncrement,
bindingOffsets[bindingIndex]));
descriptorHeapOffsets[bindingIndex]));
break;
}
@ -174,7 +175,7 @@ namespace dawn_native { namespace d3d12 {
d3d12Device->CreateUnorderedAccessView(
resource, nullptr, &uav,
viewAllocation.OffsetFrom(viewSizeIncrement,
bindingOffsets[bindingIndex]));
descriptorHeapOffsets[bindingIndex]));
break;
}
@ -201,7 +202,8 @@ namespace dawn_native { namespace d3d12 {
d3d12Device->CreateShaderResourceView(
resource, &srv,
viewAllocation.OffsetFrom(viewSizeIncrement, bindingOffsets[bindingIndex]));
viewAllocation.OffsetFrom(viewSizeIncrement,
descriptorHeapOffsets[bindingIndex]));
break;
}

149
src/dawn_native/d3d12/BindGroupLayoutD3D12.cpp
View File

@ -22,35 +22,35 @@
namespace dawn_native { namespace d3d12 {
namespace {
BindGroupLayout::DescriptorType WGPUBindingInfoToDescriptorType(
D3D12_DESCRIPTOR_RANGE_TYPE WGPUBindingInfoToDescriptorRangeType(
const BindingInfo& bindingInfo) {
switch (bindingInfo.bindingType) {
case BindingInfoType::Buffer:
switch (bindingInfo.buffer.type) {
case wgpu::BufferBindingType::Uniform:
return BindGroupLayout::DescriptorType::CBV;
return D3D12_DESCRIPTOR_RANGE_TYPE_CBV;
case wgpu::BufferBindingType::Storage:
case kInternalStorageBufferBinding:
return BindGroupLayout::DescriptorType::UAV;
return D3D12_DESCRIPTOR_RANGE_TYPE_UAV;
case wgpu::BufferBindingType::ReadOnlyStorage:
return BindGroupLayout::DescriptorType::SRV;
return D3D12_DESCRIPTOR_RANGE_TYPE_SRV;
case wgpu::BufferBindingType::Undefined:
UNREACHABLE();
}
case BindingInfoType::Sampler:
return BindGroupLayout::DescriptorType::Sampler;
return D3D12_DESCRIPTOR_RANGE_TYPE_SAMPLER;
case BindingInfoType::Texture:
case BindingInfoType::ExternalTexture:
return BindGroupLayout::DescriptorType::SRV;
return D3D12_DESCRIPTOR_RANGE_TYPE_SRV;
case BindingInfoType::StorageTexture:
switch (bindingInfo.storageTexture.access) {
case wgpu::StorageTextureAccess::ReadOnly:
return BindGroupLayout::DescriptorType::SRV;
return D3D12_DESCRIPTOR_RANGE_TYPE_SRV;
case wgpu::StorageTextureAccess::WriteOnly:
return BindGroupLayout::DescriptorType::UAV;
return D3D12_DESCRIPTOR_RANGE_TYPE_UAV;
case wgpu::StorageTextureAccess::Undefined:
UNREACHABLE();
}
@ -66,82 +66,62 @@ namespace dawn_native { namespace d3d12 {
BindGroupLayout::BindGroupLayout(Device* device, const BindGroupLayoutDescriptor* descriptor)
: BindGroupLayoutBase(device, descriptor),
mBindingOffsets(GetBindingCount()),
mDescriptorCounts{},
mDescriptorHeapOffsets(GetBindingCount()),
mShaderRegisters(GetBindingCount()),
mCbvUavSrvDescriptorCount(0),
mSamplerDescriptorCount(0),
mBindGroupAllocator(MakeFrontendBindGroupAllocator<BindGroup>(4096)) {
for (BindingIndex bindingIndex = GetDynamicBufferCount(); bindingIndex < GetBindingCount();
++bindingIndex) {
for (BindingIndex bindingIndex{0}; bindingIndex < GetBindingCount(); ++bindingIndex) {
const BindingInfo& bindingInfo = GetBindingInfo(bindingIndex);
// For dynamic resources, Dawn uses root descriptor in D3D12 backend.
// So there is no need to allocate the descriptor from descriptor heap.
// This loop starts after the dynamic buffer indices to skip counting
// dynamic resources in calculating the size of the descriptor heap.
D3D12_DESCRIPTOR_RANGE_TYPE descriptorRangeType =
WGPUBindingInfoToDescriptorRangeType(bindingInfo);
// TODO(dawn:728) In the future, special handling will be needed for external textures
// here because they encompass multiple views.
mShaderRegisters[bindingIndex] = uint32_t(bindingInfo.binding);
if (bindingIndex < GetDynamicBufferCount()) {
continue;
}
// For dynamic resources, Dawn uses root descriptor in D3D12 backend. So there is no
// need to allocate the descriptor from descriptor heap or create descriptor ranges.
ASSERT(!bindingInfo.buffer.hasDynamicOffset);
// TODO(dawn:728) In the future, special handling will be needed for external textures
// here because they encompass multiple views.
DescriptorType descriptorType = WGPUBindingInfoToDescriptorType(bindingInfo);
mBindingOffsets[bindingIndex] = mDescriptorCounts[descriptorType]++;
}
mDescriptorHeapOffsets[bindingIndex] =
descriptorRangeType == D3D12_DESCRIPTOR_RANGE_TYPE_SAMPLER
? mSamplerDescriptorCount++
: mCbvUavSrvDescriptorCount++;
auto SetDescriptorRange = [&](uint32_t index, uint32_t count, uint32_t* baseRegister,
D3D12_DESCRIPTOR_RANGE_TYPE type) -> bool {
if (count == 0) {
return false;
}
auto& range = mRanges[index];
range.RangeType = type;
range.NumDescriptors = count;
range.RegisterSpace = 0;
D3D12_DESCRIPTOR_RANGE range;
range.RangeType = descriptorRangeType;
range.NumDescriptors = 1;
range.BaseShaderRegister = GetShaderRegister(bindingIndex);
range.RegisterSpace = kRegisterSpacePlaceholder;
range.OffsetInDescriptorsFromTableStart = D3D12_DESCRIPTOR_RANGE_OFFSET_APPEND;
range.BaseShaderRegister = *baseRegister;
*baseRegister += count;
// These ranges will be copied and range.BaseShaderRegister will be set in
// d3d12::PipelineLayout to account for bind group register offsets
return true;
};
uint32_t rangeIndex = 0;
uint32_t baseRegister = 0;
std::vector<D3D12_DESCRIPTOR_RANGE>& descriptorRanges =
descriptorRangeType == D3D12_DESCRIPTOR_RANGE_TYPE_SAMPLER
? mSamplerDescriptorRanges
: mCbvUavSrvDescriptorRanges;
std::array<uint32_t, DescriptorType::Count> descriptorOffsets;
// Ranges 0-2 contain the CBV, UAV, and SRV ranges, if they exist, tightly packed
// Range 3 contains the Sampler range, if there is one
if (SetDescriptorRange(rangeIndex, mDescriptorCounts[CBV], &baseRegister,
D3D12_DESCRIPTOR_RANGE_TYPE_CBV)) {
descriptorOffsets[CBV] = mRanges[rangeIndex++].BaseShaderRegister;
}
if (SetDescriptorRange(rangeIndex, mDescriptorCounts[UAV], &baseRegister,
D3D12_DESCRIPTOR_RANGE_TYPE_UAV)) {
descriptorOffsets[UAV] = mRanges[rangeIndex++].BaseShaderRegister;
}
if (SetDescriptorRange(rangeIndex, mDescriptorCounts[SRV], &baseRegister,
D3D12_DESCRIPTOR_RANGE_TYPE_SRV)) {
descriptorOffsets[SRV] = mRanges[rangeIndex++].BaseShaderRegister;
}
uint32_t zero = 0;
SetDescriptorRange(Sampler, mDescriptorCounts[Sampler], &zero,
D3D12_DESCRIPTOR_RANGE_TYPE_SAMPLER);
descriptorOffsets[Sampler] = 0;
for (BindingIndex bindingIndex{0}; bindingIndex < GetBindingCount(); ++bindingIndex) {
const BindingInfo& bindingInfo = GetBindingInfo(bindingIndex);
if (bindingInfo.bindingType == BindingInfoType::Buffer &&
bindingInfo.buffer.hasDynamicOffset) {
// Dawn is using values in mBindingOffsets to decide register number in HLSL.
// Root descriptor needs to set this value to set correct register number in
// generated HLSL shader.
mBindingOffsets[bindingIndex] = baseRegister++;
continue;
// Try to join this range with the previous one, if the current range is a continuation
// of the previous. This is possible because the binding infos in the base type are
// sorted.
if (descriptorRanges.size() >= 2) {
D3D12_DESCRIPTOR_RANGE& previous = descriptorRanges.back();
if (previous.RangeType == range.RangeType &&
previous.BaseShaderRegister + previous.NumDescriptors ==
range.BaseShaderRegister) {
previous.NumDescriptors += range.NumDescriptors;
continue;
}
}
// TODO(dawn:728) In the future, special handling will be needed here for external
// textures because they encompass multiple views.
DescriptorType descriptorType = WGPUBindingInfoToDescriptorType(bindingInfo);
mBindingOffsets[bindingIndex] += descriptorOffsets[descriptorType];
descriptorRanges.push_back(range);
}
mViewAllocator = device->GetViewStagingDescriptorAllocator(GetCbvUavSrvDescriptorCount());
@ -181,34 +161,29 @@ namespace dawn_native { namespace d3d12 {
mBindGroupAllocator.Deallocate(bindGroup);
}
ityp::span<BindingIndex, const uint32_t> BindGroupLayout::GetBindingOffsets() const {
return {mBindingOffsets.data(), mBindingOffsets.size()};
ityp::span<BindingIndex, const uint32_t> BindGroupLayout::GetDescriptorHeapOffsets() const {
return {mDescriptorHeapOffsets.data(), mDescriptorHeapOffsets.size()};
}
uint32_t BindGroupLayout::GetCbvUavSrvDescriptorTableSize() const {
return (static_cast<uint32_t>(mDescriptorCounts[CBV] > 0) +
static_cast<uint32_t>(mDescriptorCounts[UAV] > 0) +
static_cast<uint32_t>(mDescriptorCounts[SRV] > 0));
}
uint32_t BindGroupLayout::GetSamplerDescriptorTableSize() const {
return mDescriptorCounts[Sampler] > 0;
uint32_t BindGroupLayout::GetShaderRegister(BindingIndex bindingIndex) const {
return mShaderRegisters[bindingIndex];
}
uint32_t BindGroupLayout::GetCbvUavSrvDescriptorCount() const {
return mDescriptorCounts[CBV] + mDescriptorCounts[UAV] + mDescriptorCounts[SRV];
return mCbvUavSrvDescriptorCount;
}
uint32_t BindGroupLayout::GetSamplerDescriptorCount() const {
return mDescriptorCounts[Sampler];
return mSamplerDescriptorCount;
}
const D3D12_DESCRIPTOR_RANGE* BindGroupLayout::GetCbvUavSrvDescriptorRanges() const {
return mRanges;
const std::vector<D3D12_DESCRIPTOR_RANGE>& BindGroupLayout::GetCbvUavSrvDescriptorRanges()
const {
return mCbvUavSrvDescriptorRanges;
}
const D3D12_DESCRIPTOR_RANGE* BindGroupLayout::GetSamplerDescriptorRanges() const {
return &mRanges[Sampler];
const std::vector<D3D12_DESCRIPTOR_RANGE>& BindGroupLayout::GetSamplerDescriptorRanges() const {
return mSamplerDescriptorRanges;
}
}} // namespace dawn_native::d3d12

50
src/dawn_native/d3d12/BindGroupLayoutD3D12.h
View File

@ -28,6 +28,13 @@ namespace dawn_native { namespace d3d12 {
class Device;
class StagingDescriptorAllocator;
// A purposefully invalid register space.
//
// We use the bind group index as the register space, but don't know the bind group index until
// pipeline layout creation time. This value should be replaced in PipelineLayoutD3D12.
static constexpr uint32_t kRegisterSpacePlaceholder =
D3D12_DRIVER_RESERVED_REGISTER_SPACE_VALUES_START;
class BindGroupLayout final : public BindGroupLayoutBase {
public:
static Ref<BindGroupLayout> Create(Device* device,
@ -37,28 +44,41 @@ namespace dawn_native { namespace d3d12 {
const BindGroupDescriptor* descriptor);
void DeallocateBindGroup(BindGroup* bindGroup, CPUDescriptorHeapAllocation* viewAllocation);
enum DescriptorType {
CBV,
UAV,
SRV,
Sampler,
Count,
};
// The offset (in descriptor count) into the corresponding descriptor heap. Not valid for
// dynamic binding indexes.
ityp::span<BindingIndex, const uint32_t> GetDescriptorHeapOffsets() const;
ityp::span<BindingIndex, const uint32_t> GetBindingOffsets() const;
uint32_t GetCbvUavSrvDescriptorTableSize() const;
uint32_t GetSamplerDescriptorTableSize() const;
// The D3D shader register that the Dawn binding index is mapped to by this bind group
// layout.
uint32_t GetShaderRegister(BindingIndex bindingIndex) const;
// Counts of descriptors in the descriptor tables.
uint32_t GetCbvUavSrvDescriptorCount() const;
uint32_t GetSamplerDescriptorCount() const;
const D3D12_DESCRIPTOR_RANGE* GetCbvUavSrvDescriptorRanges() const;
const D3D12_DESCRIPTOR_RANGE* GetSamplerDescriptorRanges() const;
const std::vector<D3D12_DESCRIPTOR_RANGE>& GetCbvUavSrvDescriptorRanges() const;
const std::vector<D3D12_DESCRIPTOR_RANGE>& GetSamplerDescriptorRanges() const;
private:
BindGroupLayout(Device* device, const BindGroupLayoutDescriptor* descriptor);
~BindGroupLayout() override = default;
ityp::stack_vec<BindingIndex, uint32_t, kMaxOptimalBindingsPerGroup> mBindingOffsets;
std::array<uint32_t, DescriptorType::Count> mDescriptorCounts;
D3D12_DESCRIPTOR_RANGE mRanges[DescriptorType::Count];
// Contains the offset into the descriptor heap for the given resource view. Samplers and
// non-samplers are stored in separate descriptor heaps, so the offsets should be unique
// within each group and tightly packed.
//
// Dynamic resources are not used here since their descriptors are placed directly in root
// parameters.
ityp::stack_vec<BindingIndex, uint32_t, kMaxOptimalBindingsPerGroup> mDescriptorHeapOffsets;
// Contains the shader register this binding is mapped to.
ityp::stack_vec<BindingIndex, uint32_t, kMaxOptimalBindingsPerGroup> mShaderRegisters;
uint32_t mCbvUavSrvDescriptorCount;
uint32_t mSamplerDescriptorCount;
std::vector<D3D12_DESCRIPTOR_RANGE> mCbvUavSrvDescriptorRanges;
std::vector<D3D12_DESCRIPTOR_RANGE> mSamplerDescriptorRanges;
SlabAllocator<BindGroup> mBindGroupAllocator;

43
src/dawn_native/d3d12/PipelineLayoutD3D12.cpp
View File

@ -70,9 +70,15 @@ namespace dawn_native { namespace d3d12 {
// descriptor.
std::vector<D3D12_ROOT_PARAMETER> rootParameters;
// Ranges are D3D12_DESCRIPTOR_RANGE_TYPE_(SRV|UAV|CBV|SAMPLER)
// They are grouped together so each bind group has at most 4 ranges
D3D12_DESCRIPTOR_RANGE ranges[kMaxBindGroups * 4];
size_t rangesCount = 0;
for (BindGroupIndex group : IterateBitSet(GetBindGroupLayoutsMask())) {
const BindGroupLayout* bindGroupLayout = ToBackend(GetBindGroupLayout(group));
rangesCount += bindGroupLayout->GetCbvUavSrvDescriptorRanges().size() +
bindGroupLayout->GetSamplerDescriptorRanges().size();
}
// We are taking pointers to `ranges`, so we cannot let it resize while we're pushing to it.
std::vector<D3D12_DESCRIPTOR_RANGE> ranges(rangesCount);
uint32_t rangeIndex = 0;
@ -83,7 +89,8 @@ namespace dawn_native { namespace d3d12 {
// bind group index Returns whether or not the parameter was set. A root parameter is
// not set if the number of ranges is 0
auto SetRootDescriptorTable =
[&](uint32_t rangeCount, const D3D12_DESCRIPTOR_RANGE* descriptorRanges) -> bool {
[&](const std::vector<D3D12_DESCRIPTOR_RANGE>& descriptorRanges) -> bool {
auto rangeCount = descriptorRanges.size();
if (rangeCount == 0) {
return false;
}
@ -94,8 +101,9 @@ namespace dawn_native { namespace d3d12 {
rootParameter.DescriptorTable.NumDescriptorRanges = rangeCount;
rootParameter.DescriptorTable.pDescriptorRanges = &ranges[rangeIndex];
for (uint32_t i = 0; i < rangeCount; ++i) {
ranges[rangeIndex] = descriptorRanges[i];
for (auto& range : descriptorRanges) {
ASSERT(range.RegisterSpace == kRegisterSpacePlaceholder);
ranges[rangeIndex] = range;
ranges[rangeIndex].RegisterSpace = static_cast<uint32_t>(group);
rangeIndex++;
}
@ -105,19 +113,13 @@ namespace dawn_native { namespace d3d12 {
return true;
};
if (SetRootDescriptorTable(bindGroupLayout->GetCbvUavSrvDescriptorTableSize(),
bindGroupLayout->GetCbvUavSrvDescriptorRanges())) {
if (SetRootDescriptorTable(bindGroupLayout->GetCbvUavSrvDescriptorRanges())) {
mCbvUavSrvRootParameterInfo[group] = rootParameters.size() - 1;
}
if (SetRootDescriptorTable(bindGroupLayout->GetSamplerDescriptorTableSize(),
bindGroupLayout->GetSamplerDescriptorRanges())) {
if (SetRootDescriptorTable(bindGroupLayout->GetSamplerDescriptorRanges())) {
mSamplerRootParameterInfo[group] = rootParameters.size() - 1;
}
// Get calculated shader register for root descriptors
const auto& shaderRegisters = bindGroupLayout->GetBindingOffsets();
// Init root descriptors in root signatures for dynamic buffer bindings.
// These are packed at the beginning of the layout binding info.
for (BindingIndex dynamicBindingIndex{0};
@ -136,7 +138,8 @@ namespace dawn_native { namespace d3d12 {
// Setup root descriptor.
D3D12_ROOT_DESCRIPTOR rootDescriptor;
rootDescriptor.ShaderRegister = shaderRegisters[dynamicBindingIndex];
rootDescriptor.ShaderRegister =
bindGroupLayout->GetShaderRegister(dynamicBindingIndex);
rootDescriptor.RegisterSpace = static_cast<uint32_t>(group);
// Set root descriptors in root signatures.
@ -153,15 +156,21 @@ namespace dawn_native { namespace d3d12 {
}
}
// Make sure that we added exactly the number of elements we expected. If we added more,
// |ranges| will have resized and the pointers in the |rootParameter|s will be invalid.
ASSERT(rangeIndex == rangesCount);
// Since Tint's HLSL writer doesn't currently map sets to spaces, we use the default space
// (0).
mFirstIndexOffsetRegisterSpace = 0;
BindGroupIndex firstOffsetGroup{mFirstIndexOffsetRegisterSpace};
if (GetBindGroupLayoutsMask()[firstOffsetGroup]) {
// Find the last register used on firstOffsetGroup.
auto bgl = ToBackend(GetBindGroupLayout(firstOffsetGroup));
uint32_t maxRegister = 0;
for (uint32_t shaderRegister :
ToBackend(GetBindGroupLayout(firstOffsetGroup))->GetBindingOffsets()) {
for (BindingIndex bindingIndex{0}; bindingIndex < bgl->GetBindingCount();
++bindingIndex) {
uint32_t shaderRegister = bgl->GetShaderRegister(bindingIndex);
if (shaderRegister > maxRegister) {
maxRegister = shaderRegister;
}

5
src/dawn_native/d3d12/ShaderModuleD3D12.cpp
View File

@ -203,16 +203,15 @@ namespace dawn_native { namespace d3d12 {
// with the correct registers assigned to each interface variable.
for (BindGroupIndex group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
const BindGroupLayout* bgl = ToBackend(layout->GetBindGroupLayout(group));
const auto& bindingOffsets = bgl->GetBindingOffsets();
const auto& groupBindingInfo = moduleBindingInfo[group];
for (const auto& it : groupBindingInfo) {
BindingNumber binding = it.first;
auto const& bindingInfo = it.second;
BindingIndex bindingIndex = bgl->GetBindingIndex(binding);
uint32_t bindingOffset = bindingOffsets[bindingIndex];
BindingPoint srcBindingPoint{static_cast<uint32_t>(group),
static_cast<uint32_t>(binding)};
BindingPoint dstBindingPoint{static_cast<uint32_t>(group), bindingOffset};
BindingPoint dstBindingPoint{static_cast<uint32_t>(group),
bgl->GetShaderRegister(bindingIndex)};
if (srcBindingPoint != dstBindingPoint) {
bindingPoints.emplace(srcBindingPoint, dstBindingPoint);
}

86
src/tests/end2end/BindGroupTests.cpp
View File

@ -945,6 +945,92 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) {
EXPECT_BUFFER_U32_RANGE_EQ(values.data(), outputBuffer, 0, values.size());
}
// Test that ensures that backends do not remap bindings such that dynamic and non-dynamic bindings
// conflict. This can happen if the backend treats dynamic bindings separately from non-dynamic
// bindings.
TEST_P(BindGroupTests, DynamicAndNonDynamicBindingsDoNotConflictAfterRemapping) {
auto RunTestWith = [&](bool dynamicBufferFirst) {
uint32_t dynamicBufferBindingNumber = dynamicBufferFirst ? 0 : 1;
uint32_t bufferBindingNumber = dynamicBufferFirst ? 1 : 0;
std::array<uint32_t, 1> offsets{kMinUniformBufferOffsetAlignment};
std::array<uint32_t, 2> values = {21, 67};
// Create three buffers large enough to by offset by the largest offset.
wgpu::BufferDescriptor bufferDescriptor;
bufferDescriptor.size = 2 * kMinUniformBufferOffsetAlignment + sizeof(uint32_t);
bufferDescriptor.usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst;
wgpu::Buffer dynamicBuffer = device.CreateBuffer(&bufferDescriptor);
wgpu::Buffer buffer = device.CreateBuffer(&bufferDescriptor);
// Populate the values
queue.WriteBuffer(dynamicBuffer, kMinUniformBufferOffsetAlignment,
&values[dynamicBufferBindingNumber], sizeof(uint32_t));
queue.WriteBuffer(buffer, 0, &values[bufferBindingNumber], sizeof(uint32_t));
wgpu::Buffer outputBuffer = utils::CreateBufferFromData(
device, wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::Storage, {0, 0});
// Create a bind group layout which uses a single dynamic uniform buffer.
wgpu::BindGroupLayout bgl = utils::MakeBindGroupLayout(
device,
{
{dynamicBufferBindingNumber, wgpu::ShaderStage::Compute,
wgpu::BufferBindingType::Uniform, true},
{bufferBindingNumber, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Uniform},
{2, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage},
});
wgpu::BindGroup bindGroup = utils::MakeBindGroup(
device, bgl,
{
{dynamicBufferBindingNumber, dynamicBuffer, 0, sizeof(uint32_t)},
{bufferBindingNumber, buffer, 0, sizeof(uint32_t)},
{2, outputBuffer, 0, 2 * sizeof(uint32_t)},
});
wgpu::ComputePipelineDescriptor pipelineDescriptor;
pipelineDescriptor.compute.module = utils::CreateShaderModule(device, R"(
[[block]] struct Buffer {
value : u32;
};
[[block]] struct OutputBuffer {
value : vec2<u32>;
};
[[group(0), binding(0)]] var<uniform> buffer0 : Buffer;
[[group(0), binding(1)]] var<uniform> buffer1 : Buffer;
[[group(0), binding(2)]] var<storage, read_write> outputBuffer : OutputBuffer;
[[stage(compute), workgroup_size(1)]] fn main() {
outputBuffer.value = vec2<u32>(buffer0.value, buffer1.value);
})");
pipelineDescriptor.compute.entryPoint = "main";
pipelineDescriptor.layout = utils::MakeBasicPipelineLayout(device, &bgl);
wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDescriptor);
wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder();
wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass();
computePassEncoder.SetPipeline(pipeline);
computePassEncoder.SetBindGroup(0, bindGroup, offsets.size(), offsets.data());
computePassEncoder.Dispatch(1);
computePassEncoder.EndPass();
wgpu::CommandBuffer commands = commandEncoder.Finish();
queue.Submit(1, &commands);
EXPECT_BUFFER_U32_RANGE_EQ(values.data(), outputBuffer, 0, values.size());
};
// Run the test with the dynamic buffer in index 0 and with the non-dynamic buffer in index 1,
// and vice versa. This should cause a conflict at index 0, if the binding remapping is too
// aggressive.
RunTestWith(true);
RunTestWith(false);
}
// Test that visibility of bindings in BindGroupLayout can be none
// This test passes by not asserting or crashing.
TEST_P(BindGroupTests, BindGroupLayoutVisibilityCanBeNone) {