[dawn][native] Reflect dynamic arrays from shader and validate against layout
Bug: 435317394
Change-Id: I5de3b6772f6623a572bbd877576960625674442b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/258697
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp
index 4573220..13ad3f8 100644
--- a/src/dawn/native/ShaderModule.cpp
+++ b/src/dawn/native/ShaderModule.cpp
@@ -420,6 +420,18 @@
DAWN_UNREACHABLE();
}
+ResultOrError<wgpu::DynamicBindingKind> FromArrayResourceType(
+ tint::inspector::RuntimeBindingArrayInfo::ResourceType type) {
+ switch (type) {
+ case tint::inspector::RuntimeBindingArrayInfo::ResourceType::kSampledTexture:
+ return wgpu::DynamicBindingKind::SampledTexture;
+ case tint::inspector::RuntimeBindingArrayInfo::ResourceType::kNone:
+ return DAWN_VALIDATION_ERROR(
+ "Attempted to convert 'None' array resource type from Tint.");
+ }
+ DAWN_UNREACHABLE();
+}
+
// Validation errors, if any, are stored within outputParseResult instead of get returned as
// ErrorData.
MaybeError ParseWGSL(std::unique_ptr<tint::Source::File> file,
@@ -716,6 +728,30 @@
});
}
+MaybeError ValidateCompatibilityOfDynamicBindingArrayWithLayout(
+ DeviceBase* device,
+ const BindGroupLayoutInternalBase* layout,
+ const GroupDynamicBindingArrayInfo& shaderDynamicArray) {
+ DAWN_INVALID_IF(!layout->HasDynamicArray(), "%s doesn't contain a dynamic binding array.",
+ layout);
+
+ DAWN_INVALID_IF(layout->GetAPIDynamicArrayStart() != shaderDynamicArray.start,
+ "@binding for the dynamic array in the shader (%u) doesn't match the start "
+ "(%u) defined in %s.",
+ shaderDynamicArray.start, layout->GetAPIDynamicArrayStart(), layout);
+
+ // If the dynamic binding array is never accessed with any type in the shader, it is valid to
+ // use with any DynamicArrayKind.
+ if (shaderDynamicArray.kind != wgpu::DynamicBindingKind::Undefined) {
+ DAWN_INVALID_IF(shaderDynamicArray.kind != layout->GetDynamicArrayKind(),
+ "Shader dynamic binding array is used with types (of kind %s) incompatible "
+ "with %s's kind of dynamic binding array (%s).",
+ shaderDynamicArray.kind, layout, layout->GetDynamicArrayKind());
+ }
+
+ return {};
+}
+
MaybeError ValidateCompatibilityWithBindGroupLayout(DeviceBase* device,
BindGroupIndex group,
const EntryPointMetadata& entryPoint,
@@ -730,6 +766,15 @@
group, bindingId, layout);
}
+ // Check that the dynamic binding array, if any in the shader, matches the BindGroupLayout.
+ if (entryPoint.dynamicBindingArrays.contains(group)) {
+ DAWN_TRY_CONTEXT(
+ ValidateCompatibilityOfDynamicBindingArrayWithLayout(
+ device, layout, entryPoint.dynamicBindingArrays.at(group)),
+ "validating that the entry-point's dynamic binding array for @group(%u) matches %s",
+ group, layout);
+ }
+
return {};
}
@@ -1175,6 +1220,51 @@
resource.binding, resource.bind_group);
}
+ // Dynamic binding array reflection
+ for (const tint::inspector::RuntimeBindingArrayInfo& array :
+ inspector->GetRuntimeBindingArrayInfo(entryPoint.name)) {
+ BindGroupIndex group(array.group);
+ if (DelayedInvalidIf(group >= kMaxBindGroupsTyped,
+ "The entry-point uses a binding with a group decoration (%u) "
+ "that exceeds maxBindGroups (%u) - 1.",
+ group, kMaxBindGroups)) {
+ continue;
+ }
+
+ BindingNumber binding(array.binding);
+ if (DelayedInvalidIf(
+ binding >= kMaxBindingsPerBindGroupTyped,
+ "Binding number (%u) exceeds the maxBindingsPerBindGroup limit (%u) - 1.",
+ uint32_t(binding), kMaxBindingsPerBindGroup)) {
+ continue;
+ }
+
+ // Check that all the uses of the dynamic binding array have compatible DynamicArrayKind.
+ wgpu::DynamicBindingKind kind = wgpu::DynamicBindingKind::Undefined;
+ for (const auto& type : array.type_info) {
+ wgpu::DynamicBindingKind kindForType;
+ DAWN_TRY_ASSIGN(kindForType, FromArrayResourceType(type.type));
+
+ // This is the first kind that we compute, just store it.
+ if (kind == wgpu::DynamicBindingKind::Undefined) {
+ kind = kindForType;
+ continue;
+ }
+
+ DAWN_INVALID_IF(kindForType != kind,
+ "Dynamic binding array for @group(%u) used with two incompatible kinds "
+ "of types %s vs. %s",
+ group, kind, kindForType);
+ }
+
+ DAWN_INVALID_IF(metadata->dynamicBindingArrays.contains(group),
+ "Duplicate dynamic binding array for group: %u.", group);
+ metadata->dynamicBindingArrays[group] = {{
+ .start = binding,
+ .kind = kind,
+ }};
+ }
+
// Sampler binding point placeholder for non-sampler texture usage. Make it
// ToTint(EntryPointMetadata::nonSamplerBindingPoint), so that we have
// FromTint(tintNonSamplerBindingPoint) == EntryPointMetadata::nonSamplerBindingPoint, and we
@@ -1511,7 +1601,11 @@
for (BindGroupIndex group : ~layout->GetBindGroupLayoutsMask()) {
DAWN_INVALID_IF(entryPoint.bindings[group].size() > 0,
"The entry-point uses bindings in group %u but %s doesn't have a "
- "BindGroupLayout for this index",
+ "BindGroupLayout for this index.",
+ group, layout);
+ DAWN_INVALID_IF(entryPoint.dynamicBindingArrays.contains(group),
+ "The entry-point uses a dynamic binding array in group %u but %s doesn't "
+ "have a BindGroupLayout for this index.",
group, layout);
}
diff --git a/src/dawn/native/ShaderModule.h b/src/dawn/native/ShaderModule.h
index 53ee98b..a70f61a 100644
--- a/src/dawn/native/ShaderModule.h
+++ b/src/dawn/native/ShaderModule.h
@@ -202,6 +202,15 @@
using BindingGroupInfoMap = absl::flat_hash_map<BindingNumber, ShaderBindingInfo>;
using BindingInfoArray = ityp::array<BindGroupIndex, BindingGroupInfoMap, kMaxBindGroups>;
+// Shader metadata that's the equivalent for the dynamic binding arrays in the BGLs.
+#define GROUP_DYNAMIC_BINDING_ARRAY_INFO_MEMBERS(X) \
+ X(BindingNumber, start) \
+ X(wgpu::DynamicBindingKind, kind)
+DAWN_SERIALIZABLE(struct, GroupDynamicBindingArrayInfo, GROUP_DYNAMIC_BINDING_ARRAY_INFO_MEMBERS){};
+#undef GROUP_DYNAMIC_BINDING_ARRAY_INFO_MEMBERS
+
+using DynamicBindingArrayInfo = absl::flat_hash_map<BindGroupIndex, GroupDynamicBindingArrayInfo>;
+
// Define types for the shader reflection data structures in detail namespaces to prevent messing
// up dawn::native namespace. These types can be exposed within EntryPointMetadata if needed.
namespace detail {
@@ -273,6 +282,9 @@
X(std::vector<std::string>, infringedLimitErrors) \
/* bindings[G][B] is the reflection data for the binding defined with @group(G) @binding(B)*/ \
X(BindingInfoArray, bindings) \
+ /* dynamicBindingArray[G] is the reflection data for the dynamic binding array of @group(G)*/ \
+ /* if one is present in the shader module */ \
+ X(DynamicBindingArrayInfo, dynamicBindingArrays) \
/* Contains the reflection information of all sampler and non-sampler texture (storage */ \
/* texture not included) usage in the entry point. For non-sampler usage, */ \
/* nonSamplerBindingPoint is used for sampler slot. */ \
diff --git a/src/dawn/tests/unittests/validation/DynamicBindingArrayValidationTests.cpp b/src/dawn/tests/unittests/validation/DynamicBindingArrayValidationTests.cpp
index fb8c439..0f84540 100644
--- a/src/dawn/tests/unittests/validation/DynamicBindingArrayValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/DynamicBindingArrayValidationTests.cpp
@@ -618,5 +618,256 @@
}
}
+// Test that a shader using a dynamic binding array requires a layout with one.
+TEST_F(DynamicBindingArrayTests, ShaderRequiresLayoutWithDynamicArray) {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module = utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding(0) var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+
+ // Success case, the layout has a dynamic binding array.
+ wgpu::BindGroupLayout bglDynamic =
+ MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bglDynamic);
+ device.CreateComputePipeline(&csDesc);
+
+ // Error case, the layout doesn't have a dynamic binding array.
+ wgpu::BindGroupLayout bglStatic = utils::MakeBindGroupLayout(device, {});
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bglStatic);
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
+
+ // Error case, the layout doesn't have a dynamic binding array (even if there is a similar
+ // looking binding).
+ wgpu::BindGroupLayout bglStaticWithTexture = utils::MakeBindGroupLayout(
+ device, {{0, wgpu::ShaderStage::Compute, wgpu::TextureSampleType::Float}});
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bglStaticWithTexture);
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
+}
+
+// Test that it is valid to have a layout specifying a dynamic binding array with a shader that
+// doesn't have one.
+TEST_F(DynamicBindingArrayTests, ShaderNoDynamicArrayWithLayoutThatHasOne) {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module = utils::CreateShaderModule(device, R"(
+ @compute @workgroup_size(1) fn main() {
+ }
+ )");
+ wgpu::BindGroupLayout bglDynamic =
+ MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bglDynamic);
+ device.CreateComputePipeline(&csDesc);
+}
+
+// Test that the dynamic array start must match between shader and layout.
+TEST_F(DynamicBindingArrayTests, ShaderAndLayoutDynamicArrayStartMatches) {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module = utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding(1) var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+
+ // Success case, start of the array matches.
+ wgpu::BindGroupLayout bgl1 = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture, 1);
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bgl1);
+ device.CreateComputePipeline(&csDesc);
+
+ // Error case, layout start is before the shader's start.
+ wgpu::BindGroupLayout bgl0 = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture, 0);
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bgl0);
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
+
+ // Error case, layout start is after the shader's start.
+ wgpu::BindGroupLayout bgl2 = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture, 2);
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bgl2);
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
+}
+
+// Test that the @binding decoration of the dynamic array must be less than maxBindingsPerBindGroup.
+TEST_F(DynamicBindingArrayTests, ShaderArrayStartLessThanMaxBindingsPerBindGroup) {
+ wgpu::BindGroupLayout bgl =
+ MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture, kMaxBindingsPerBindGroup - 1);
+
+ // Control case, we are just below the limit.
+ {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module =
+ utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding()" + std::to_string(kMaxBindingsPerBindGroup - 1) +
+ R"() var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bgl);
+ device.CreateComputePipeline(&csDesc);
+ }
+
+ // Error case, we are above the limit.
+ {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module =
+ utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding()" + std::to_string(kMaxBindingsPerBindGroup) +
+ R"() var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+ csDesc.layout = utils::MakeBasicPipelineLayout(device, &bgl);
+ // Two errors happen because we cannot create a layout that matches, but check that the
+ // shader's validation about maxBindingsPerBindGroup is the one that's reported.
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc),
+ testing::HasSubstr("maxBindingsPerBindGroup"));
+ }
+}
+
+// Test that the @group decoration of the dynamic array must be less than maxBindGroups.
+TEST_F(DynamicBindingArrayTests, ShaderArrayAtMaxBindGroups) {
+ std::array<wgpu::BindGroupLayout, kMaxBindGroups> bgls;
+ bgls[bgls.size() - 1] = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+
+ wgpu::PipelineLayoutDescriptor plDesc;
+ plDesc.bindGroupLayoutCount = bgls.size();
+ plDesc.bindGroupLayouts = bgls.data();
+ wgpu::PipelineLayout pl = device.CreatePipelineLayout(&plDesc);
+
+ // Control case, we are just below the limit.
+ {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module =
+ utils::CreateShaderModule(device,
+ R"(
+ enable chromium_experimental_dynamic_binding;
+ @group()" + std::to_string(kMaxBindGroups - 1) +
+ R"() @binding(0) var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+ csDesc.layout = pl;
+ device.CreateComputePipeline(&csDesc);
+ }
+
+ // Error case, we are above the limit.
+ {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module =
+ utils::CreateShaderModule(device,
+ R"(
+ enable chromium_experimental_dynamic_binding;
+ @group()" + std::to_string(kMaxBindGroups) +
+ R"() @binding(0) var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+ csDesc.layout = pl;
+ // Two errors happen because we cannot create a layout that matches, but check that the
+ // shader's validation about maxBindingsPerBindGroup is the one that's reported.
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc),
+ testing::HasSubstr("maxBindGroups"));
+ }
+}
+
+// Test that the group for the dynamic binding array must be in the PipelineLayout.
+TEST_F(DynamicBindingArrayTests, ShaderBindingArrayMustHaveGroupInPipelineLayout) {
+ std::array<wgpu::BindGroupLayout, 3> bgls = {
+ nullptr, MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture), nullptr};
+
+ wgpu::PipelineLayoutDescriptor plDesc;
+ plDesc.bindGroupLayoutCount = bgls.size();
+ plDesc.bindGroupLayouts = bgls.data();
+
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.layout = device.CreatePipelineLayout(&plDesc);
+
+ // Control case, the group is in the pipeline layout.
+ csDesc.compute.module = utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(1) @binding(0) var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+ device.CreateComputePipeline(&csDesc);
+
+ // Error case, the group is not in the layout (@group(0) case)
+ csDesc.compute.module = utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding(0) var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
+
+ // Error case, the group is not in the layout (@group(2) case)
+ csDesc.compute.module = utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(2) @binding(0) var a : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ }
+ )");
+ ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
+}
+
+// Test that a shader cannot have two dynamic binding arrays on the same group.
+TEST_F(DynamicBindingArrayTests, ShaderTwoDynamicArraysSameGroupIsAnError) {
+ // Control case, the two dynamic binding arrays are on different groups.
+ utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding(0) var a : binding_array<texture_2d<f32>>;
+ @group(1) @binding(1) var b : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ _ = b[42];
+ }
+ )");
+
+ // Error case, the two dynamic binding arrays are on the same group.
+ ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding(0) var a : binding_array<texture_2d<f32>>;
+ @group(0) @binding(1) var b : binding_array<texture_2d<f32>>;
+
+ @compute @workgroup_size(1) fn main() {
+ _ = a[42];
+ _ = b[42];
+ }
+ )"));
+}
+
+// TODO(https://crbug.com/435317394): Add tests for the DynamicArrayKind. It is not possible to do
+// it at the moment because we cannot reflect DynamicArrayKind::Undefined (would require referencing
+// but not indexing the array) or any value that's not DynamicArrayKind::SampledTexture (no support
+// in Dawn or Tint for other cases). Tests to add after that are:
+// - The kind in the layout must matche the deduced kind for the shader.
+// - Case with a runtime but typed binding_array.
+// - Case with an untyped binding_array.
+// - A shader only referencing but not indexing an untyped binding_array is valid to use with any
+// DynamicArrayKind in the layout.
+// - An error is produced at shader module compilation time if it uses the same binding_array with
+// different DynamicArrayKinds.
+
} // anonymous namespace
} // namespace dawn
