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>
diff --git a/src/dawn_native/d3d12/BindGroupD3D12.cpp b/src/dawn_native/d3d12/BindGroupD3D12.cpp
index a9f2731..4832cdd 100644
--- a/src/dawn_native/d3d12/BindGroupD3D12.cpp
+++ b/src/dawn_native/d3d12/BindGroupD3D12.cpp
@@ -40,7 +40,7 @@
mCPUViewAllocation = viewAllocation;
- const auto& bindingOffsets = bgl->GetBindingOffsets();
+ const auto& descriptorHeapOffsets = bgl->GetDescriptorHeapOffsets();
ID3D12Device* d3d12Device = device->GetD3D12Device();
@@ -74,8 +74,8 @@
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 @@
d3d12Device->CreateUnorderedAccessView(
resource, nullptr, &desc,
viewAllocation.OffsetFrom(viewSizeIncrement,
- bindingOffsets[bindingIndex]));
+ descriptorHeapOffsets[bindingIndex]));
break;
}
case wgpu::BufferBindingType::ReadOnlyStorage: {
@@ -118,7 +118,7 @@
d3d12Device->CreateShaderResourceView(
resource, &desc,
viewAllocation.OffsetFrom(viewSizeIncrement,
- bindingOffsets[bindingIndex]));
+ descriptorHeapOffsets[bindingIndex]));
break;
}
case wgpu::BufferBindingType::Undefined:
@@ -142,7 +142,8 @@
d3d12Device->CreateShaderResourceView(
resource, &srv,
- viewAllocation.OffsetFrom(viewSizeIncrement, bindingOffsets[bindingIndex]));
+ viewAllocation.OffsetFrom(viewSizeIncrement,
+ descriptorHeapOffsets[bindingIndex]));
break;
}
@@ -165,7 +166,7 @@
d3d12Device->CreateShaderResourceView(
resource, &srv,
viewAllocation.OffsetFrom(viewSizeIncrement,
- bindingOffsets[bindingIndex]));
+ descriptorHeapOffsets[bindingIndex]));
break;
}
@@ -174,7 +175,7 @@
d3d12Device->CreateUnorderedAccessView(
resource, nullptr, &uav,
viewAllocation.OffsetFrom(viewSizeIncrement,
- bindingOffsets[bindingIndex]));
+ descriptorHeapOffsets[bindingIndex]));
break;
}
@@ -201,7 +202,8 @@
d3d12Device->CreateShaderResourceView(
resource, &srv,
- viewAllocation.OffsetFrom(viewSizeIncrement, bindingOffsets[bindingIndex]));
+ viewAllocation.OffsetFrom(viewSizeIncrement,
+ descriptorHeapOffsets[bindingIndex]));
break;
}
diff --git a/src/dawn_native/d3d12/BindGroupLayoutD3D12.cpp b/src/dawn_native/d3d12/BindGroupLayoutD3D12.cpp
index 28719de..30b5a54 100644
--- a/src/dawn_native/d3d12/BindGroupLayoutD3D12.cpp
+++ b/src/dawn_native/d3d12/BindGroupLayoutD3D12.cpp
@@ -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 @@
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 @@
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
diff --git a/src/dawn_native/d3d12/BindGroupLayoutD3D12.h b/src/dawn_native/d3d12/BindGroupLayoutD3D12.h
index 5035666..eb7b313 100644
--- a/src/dawn_native/d3d12/BindGroupLayoutD3D12.h
+++ b/src/dawn_native/d3d12/BindGroupLayoutD3D12.h
@@ -28,6 +28,13 @@
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 @@
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;
diff --git a/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp b/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp
index f239fe5..448ec22 100644
--- a/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp
+++ b/src/dawn_native/d3d12/PipelineLayoutD3D12.cpp
@@ -70,9 +70,15 @@
// 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 @@
// 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 @@
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 @@
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 @@
// 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 @@
}
}
+ // 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;
}
diff --git a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
index eac2a4f..2ecb5fc 100644
--- a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
@@ -203,16 +203,15 @@
// 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);
}
diff --git a/src/tests/end2end/BindGroupTests.cpp b/src/tests/end2end/BindGroupTests.cpp
index 90ecbed..81eeb50 100644
--- a/src/tests/end2end/BindGroupTests.cpp
+++ b/src/tests/end2end/BindGroupTests.cpp
@@ -945,6 +945,92 @@
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) {
