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>

src/dawn/native/ShaderModule.cpp

@@ -696,7 +696,7 @@ ResultOrError<std::unique_ptr<EntryPointMetadata>> ReflectEntryPointUsingTint(
 
         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;

src/dawn/tests/end2end/BindGroupTests.cpp

@@ -1045,17 +1045,13 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) {
             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);

src/dawn/tests/end2end/ComputeDispatchTests.cpp

@@ -28,22 +28,18 @@ class ComputeDispatchTests : public DawnTest {
         // 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 {
+            @group(0) @binding(0) var<storage, read_write> output : vec3<u32>;
-                workGroups : vec3<u32>
-            }
-
-            @group(0) @binding(0) var<storage, read_write> output : OutputBuf;
 
             @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 @@ class ComputeDispatchTests : public DawnTest {
 
         // Test the use of the compute pipelines without using @num_workgroups
         wgpu::ShaderModule moduleWithoutNumWorkgroups = utils::CreateShaderModule(device, R"(
-            struct InputBuf {
+            @group(0) @binding(0) var<uniform> input : vec3<u32>;
-                expectedDispatch : vec3<u32>
+            @group(0) @binding(1) var<storage, read_write> output : 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;
 
             @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;

src/dawn/tests/end2end/ComputeLayoutMemoryBufferTests.cpp

@@ -640,7 +640,7 @@ fn main() {
 
     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

src/dawn/tests/end2end/MultisampledRenderingTests.cpp

@@ -329,7 +329,7 @@ TEST_P(MultisampledRenderingTest, MultisampledRenderingWithDepthTest) {
         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 @@ TEST_P(MultisampledRenderingTest, MultisampledRenderingWithDepthTest) {
             {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 @@ TEST_P(MultisampledRenderingTest, MultisampledRenderingWithDepthTestAndSampleMas
         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 @@ TEST_P(MultisampledRenderingTest, MultisampledRenderingWithDepthTestAndSampleMas
             {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 @@ TEST_P(MultisampledRenderingTest, MultisampledRenderingWithDepthTestAndAlphaToCo
         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 @@ TEST_P(MultisampledRenderingTest, MultisampledRenderingWithDepthTestAndAlphaToCo
             {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(),

src/dawn/tests/perf_tests/DrawCallPerf.cpp

@@ -41,26 +41,17 @@ constexpr char kVertexShader[] = R"(
         })";
 
 constexpr char kFragmentShaderA[] = R"(
-        struct Uniforms {
+        @group(0) @binding(0) var<uniform> color : vec3<f32>;
-            color : vec3<f32>
-        }
-        @group(0) @binding(0) var<uniform> uniforms : Uniforms;
         @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 {
+        @group(0) @binding(0) var<uniform> constant_color : vec3<f32>;
-            color : vec3<f32>
+        @group(1) @binding(0) var<uniform> uniform_color : vec3<f32>;
-        }
-        struct Uniforms {
-            color : vec3<f32>
-        }
-        @group(0) @binding(0) var<uniform> constants : Constants;
-        @group(1) @binding(0) var<uniform> uniforms : Uniforms;
 
         @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 {

src/dawn/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp

@@ -571,7 +571,7 @@ TEST_F(MinBufferSizeDefaultLayoutTests, MultipleBindGroups) {
 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 @@ TEST_F(MinBufferSizeDefaultLayoutTests, RenderPassConsidersBothStages) {
 
     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));
+        }
+    });
+}