Include padding bytes in minBindingSize validation
The "minimum buffer binding size" for a buffer binding variable with
type `T` is `SizeOf(T)`, which includes trailing padding bytes for
structures.
Update several tests that were not creating large enough buffers. Add
a new test for validating the size of a buffer with a non-struct vec3
type, which should still be 12 bytes.
Fixed: tint:1377
Change-Id: Iddbc22c561a67b6aa6659d7ddf78b1b12b230930
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111582
Reviewed-by: Austin Eng <enga@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp
index 5bf019c..d473d89 100644
--- a/src/dawn/native/ShaderModule.cpp
+++ b/src/dawn/native/ShaderModule.cpp
@@ -696,7 +696,7 @@
switch (info.bindingType) {
case BindingInfoType::Buffer:
- info.buffer.minBindingSize = resource.size_no_padding;
+ info.buffer.minBindingSize = resource.size;
DAWN_TRY_ASSIGN(info.buffer.type,
TintResourceTypeToBufferBindingType(resource.resource_type));
break;
diff --git a/src/dawn/tests/end2end/BindGroupTests.cpp b/src/dawn/tests/end2end/BindGroupTests.cpp
index fc78112..1bc8871 100644
--- a/src/dawn/tests/end2end/BindGroupTests.cpp
+++ b/src/dawn/tests/end2end/BindGroupTests.cpp
@@ -1045,17 +1045,13 @@
value : u32
}
- struct OutputBuffer {
- value : vec3<u32>
- }
-
@group(0) @binding(2) var<uniform> buffer2 : Buffer;
@group(0) @binding(3) var<storage, read> buffer3 : Buffer;
@group(0) @binding(0) var<storage, read> buffer0 : Buffer;
- @group(0) @binding(4) var<storage, read_write> outputBuffer : OutputBuffer;
+ @group(0) @binding(4) var<storage, read_write> outputBuffer : vec3<u32>;
@compute @workgroup_size(1) fn main() {
- outputBuffer.value = vec3<u32>(buffer0.value, buffer2.value, buffer3.value);
+ outputBuffer = vec3<u32>(buffer0.value, buffer2.value, buffer3.value);
})");
pipelineDescriptor.compute.entryPoint = "main";
pipelineDescriptor.layout = utils::MakeBasicPipelineLayout(device, &bgl);
diff --git a/src/dawn/tests/end2end/ComputeDispatchTests.cpp b/src/dawn/tests/end2end/ComputeDispatchTests.cpp
index 1cdf759..472d54d 100644
--- a/src/dawn/tests/end2end/ComputeDispatchTests.cpp
+++ b/src/dawn/tests/end2end/ComputeDispatchTests.cpp
@@ -28,22 +28,18 @@
// Write workgroup number into the output buffer if we saw the biggest dispatch
// To make sure the dispatch was not called, write maximum u32 value for 0 dispatches
wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
- struct OutputBuf {
- workGroups : vec3<u32>
- }
-
- @group(0) @binding(0) var<storage, read_write> output : OutputBuf;
+ @group(0) @binding(0) var<storage, read_write> output : vec3<u32>;
@compute @workgroup_size(1, 1, 1)
fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>,
@builtin(num_workgroups) dispatch : vec3<u32>) {
if (dispatch.x == 0u || dispatch.y == 0u || dispatch.z == 0u) {
- output.workGroups = vec3<u32>(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu);
+ output = vec3<u32>(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu);
return;
}
if (all(GlobalInvocationID == dispatch - vec3<u32>(1u, 1u, 1u))) {
- output.workGroups = dispatch;
+ output = dispatch;
}
})");
@@ -54,27 +50,20 @@
// Test the use of the compute pipelines without using @num_workgroups
wgpu::ShaderModule moduleWithoutNumWorkgroups = utils::CreateShaderModule(device, R"(
- struct InputBuf {
- expectedDispatch : vec3<u32>
- }
- struct OutputBuf {
- workGroups : vec3<u32>
- }
-
- @group(0) @binding(0) var<uniform> input : InputBuf;
- @group(0) @binding(1) var<storage, read_write> output : OutputBuf;
+ @group(0) @binding(0) var<uniform> input : vec3<u32>;
+ @group(0) @binding(1) var<storage, read_write> output : vec3<u32>;
@compute @workgroup_size(1, 1, 1)
fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
- let dispatch : vec3<u32> = input.expectedDispatch;
+ let dispatch : vec3<u32> = input;
if (dispatch.x == 0u || dispatch.y == 0u || dispatch.z == 0u) {
- output.workGroups = vec3<u32>(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu);
+ output = vec3<u32>(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu);
return;
}
if (all(GlobalInvocationID == dispatch - vec3<u32>(1u, 1u, 1u))) {
- output.workGroups = dispatch;
+ output = dispatch;
}
})");
csDesc.compute.module = moduleWithoutNumWorkgroups;
diff --git a/src/dawn/tests/end2end/ComputeLayoutMemoryBufferTests.cpp b/src/dawn/tests/end2end/ComputeLayoutMemoryBufferTests.cpp
index 79cacd9..bf9ccdc 100644
--- a/src/dawn/tests/end2end/ComputeLayoutMemoryBufferTests.cpp
+++ b/src/dawn/tests/end2end/ComputeLayoutMemoryBufferTests.cpp
@@ -640,7 +640,7 @@
MemoryDataBuilder expectedDataBuilder; // The expected data to be copied by the shader
expectedDataBuilder.AddSubBuilder(field.GetDataBuilder());
- expectedDataBuilder.AlignTo(4); // Storage buffer size must be a multiple of 4
+ expectedDataBuilder.AlignTo(std::max<size_t>(field.GetAlign(), 4u));
// Expectation and input buffer have identical data bytes but different padding bytes.
// Initializes the dst buffer with data bytes different from input and expectation, and padding
diff --git a/src/dawn/tests/end2end/MultisampledRenderingTests.cpp b/src/dawn/tests/end2end/MultisampledRenderingTests.cpp
index 407ae7c..88bb571 100644
--- a/src/dawn/tests/end2end/MultisampledRenderingTests.cpp
+++ b/src/dawn/tests/end2end/MultisampledRenderingTests.cpp
@@ -329,7 +329,7 @@
utils::ComboRenderPassDescriptor renderPass =
CreateComboRenderPassDescriptorForTest({mMultisampledColorView}, {mResolveView},
wgpu::LoadOp::Clear, wgpu::LoadOp::Clear, true);
- std::array<float, 5> kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color
+ std::array<float, 8> kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color
0.2f}; // depth
constexpr uint32_t kSize = sizeof(kUniformData);
EncodeRenderPassForTest(commandEncoder, renderPass, pipeline, kUniformData.data(), kSize);
@@ -343,7 +343,7 @@
{mMultisampledColorView}, {mResolveView}, wgpu::LoadOp::Load, wgpu::LoadOp::Load,
kTestDepth);
- std::array<float, 5> kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color
+ std::array<float, 8> kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color
0.5f}; // depth
constexpr uint32_t kSize = sizeof(kUniformData);
EncodeRenderPassForTest(commandEncoder, renderPass, pipeline, kUniformData.data(), kSize);
@@ -714,7 +714,7 @@
utils::ComboRenderPassDescriptor renderPass =
CreateComboRenderPassDescriptorForTest({mMultisampledColorView}, {mResolveView},
wgpu::LoadOp::Clear, wgpu::LoadOp::Clear, true);
- std::array<float, 5> kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color
+ std::array<float, 8> kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color
0.2f}; // depth
constexpr uint32_t kSize = sizeof(kUniformData);
EncodeRenderPassForTest(commandEncoder, renderPass, pipelineGreen, kUniformData.data(),
@@ -729,7 +729,7 @@
{mMultisampledColorView}, {mResolveView}, wgpu::LoadOp::Load, wgpu::LoadOp::Load,
kTestDepth);
- std::array<float, 5> kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color
+ std::array<float, 8> kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color
0.5f}; // depth
constexpr uint32_t kSize = sizeof(kUniformData);
EncodeRenderPassForTest(commandEncoder, renderPass, pipelineRed, kUniformData.data(),
@@ -1003,7 +1003,7 @@
utils::ComboRenderPassDescriptor renderPass = CreateComboRenderPassDescriptorForTest(
{mMultisampledColorView}, {mResolveView}, wgpu::LoadOp::Clear, wgpu::LoadOp::Clear,
kTestDepth);
- std::array<float, 5> kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color
+ std::array<float, 8> kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color
0.2f}; // depth
constexpr uint32_t kSize = sizeof(kUniformData);
EncodeRenderPassForTest(commandEncoder, renderPass, pipelineGreen, kUniformData.data(),
@@ -1018,7 +1018,7 @@
{mMultisampledColorView}, {mResolveView}, wgpu::LoadOp::Load, wgpu::LoadOp::Load,
kTestDepth);
- std::array<float, 5> kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color
+ std::array<float, 8> kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color
0.5f}; // depth
constexpr uint32_t kSize = sizeof(kUniformData);
EncodeRenderPassForTest(commandEncoder, renderPass, pipelineRed, kUniformData.data(),
diff --git a/src/dawn/tests/perf_tests/DrawCallPerf.cpp b/src/dawn/tests/perf_tests/DrawCallPerf.cpp
index 335f73b..4dd9715 100644
--- a/src/dawn/tests/perf_tests/DrawCallPerf.cpp
+++ b/src/dawn/tests/perf_tests/DrawCallPerf.cpp
@@ -41,26 +41,17 @@
})";
constexpr char kFragmentShaderA[] = R"(
- struct Uniforms {
- color : vec3<f32>
- }
- @group(0) @binding(0) var<uniform> uniforms : Uniforms;
+ @group(0) @binding(0) var<uniform> color : vec3<f32>;
@fragment fn main() -> @location(0) vec4<f32> {
- return vec4<f32>(uniforms.color * (1.0 / 5000.0), 1.0);
+ return vec4<f32>(color * (1.0 / 5000.0), 1.0);
})";
constexpr char kFragmentShaderB[] = R"(
- struct Constants {
- color : vec3<f32>
- }
- struct Uniforms {
- color : vec3<f32>
- }
- @group(0) @binding(0) var<uniform> constants : Constants;
- @group(1) @binding(0) var<uniform> uniforms : Uniforms;
+ @group(0) @binding(0) var<uniform> constant_color : vec3<f32>;
+ @group(1) @binding(0) var<uniform> uniform_color : vec3<f32>;
@fragment fn main() -> @location(0) vec4<f32> {
- return vec4<f32>((constants.color + uniforms.color) * (1.0 / 5000.0), 1.0);
+ return vec4<f32>((constant_color + uniform_color) * (1.0 / 5000.0), 1.0);
})";
enum class Pipeline {
diff --git a/src/dawn/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp b/src/dawn/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp
index dbd07c2..b09d0c9 100644
--- a/src/dawn/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp
@@ -571,7 +571,7 @@
TEST_F(MinBufferSizeDefaultLayoutTests, NonDefaultLayout) {
CheckShaderBindingSizeReflection(
{{{0, 0, "@size(256) a : u32, b : u32,", "u32", "a", 260},
- {0, 1, "c : u32, @align(16) d : u32,", "u32", "c", 20},
+ {0, 1, "c : u32, @align(16) d : u32,", "u32", "c", 32},
{0, 2, "d : array<array<u32, 10>, 3>,", "u32", "d[0][0]", 120},
{0, 3, "e : array<array<u32, 10>>,", "u32", "e[0][0]", 40}}});
}
@@ -593,3 +593,31 @@
CheckLayoutBindingSizeValidation(renderLayout, {{0, 0, "", "", "", 8}, {0, 1, "", "", "", 16}});
}
+
+// Make sure that buffers with non-struct vec3 types do not include padding in the min buffer size.
+TEST_F(MinBufferSizePipelineCreationTests, NonStructVec3) {
+ std::vector<BindingDescriptor> bindings = {{0, 0, "", "", "", 12}, {0, 1, "", "", "", 12}};
+
+ auto MakeShader = [](const char* stageAttributes) {
+ std::ostringstream ostream;
+ ostream << "@group(0) @binding(0) var<storage, read_write> buffer : vec3<u32>;\n";
+ ostream << stageAttributes << " fn main() { buffer = vec3(42, 0, 7); }\n";
+ return ostream.str();
+ };
+ std::string computeShader = MakeShader("@compute @workgroup_size(1)");
+ std::string fragShader = MakeShader("@fragment");
+ std::string vertexShader = CreateVertexShaderWithBindings({});
+
+ CheckSizeBounds({12}, [&](const std::vector<uint64_t>& sizes, bool expectation) {
+ wgpu::BindGroupLayout layout = utils::MakeBindGroupLayout(
+ device, {{0, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment,
+ wgpu::BufferBindingType::Storage, false, sizes[0]}});
+ if (expectation) {
+ CreateRenderPipeline({layout}, vertexShader, fragShader);
+ CreateComputePipeline({layout}, computeShader);
+ } else {
+ ASSERT_DEVICE_ERROR(CreateRenderPipeline({layout}, vertexShader, fragShader));
+ ASSERT_DEVICE_ERROR(CreateComputePipeline({layout}, computeShader));
+ }
+ });
+}
