diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp index 5bf019cc32..d473d89ae6 100644 --- a/src/dawn/native/ShaderModule.cpp +++ b/src/dawn/native/ShaderModule.cpp @@ -696,7 +696,7 @@ ResultOrError> 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; diff --git a/src/dawn/tests/end2end/BindGroupTests.cpp b/src/dawn/tests/end2end/BindGroupTests.cpp index fc78112638..1bc8871dc1 100644 --- a/src/dawn/tests/end2end/BindGroupTests.cpp +++ b/src/dawn/tests/end2end/BindGroupTests.cpp @@ -1045,17 +1045,13 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) { value : u32 } - struct OutputBuffer { - value : vec3 - } - @group(0) @binding(2) var buffer2 : Buffer; @group(0) @binding(3) var buffer3 : Buffer; @group(0) @binding(0) var buffer0 : Buffer; - @group(0) @binding(4) var outputBuffer : OutputBuffer; + @group(0) @binding(4) var outputBuffer : vec3; @compute @workgroup_size(1) fn main() { - outputBuffer.value = vec3(buffer0.value, buffer2.value, buffer3.value); + outputBuffer = vec3(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 1cdf75971c..472d54d123 100644 --- a/src/dawn/tests/end2end/ComputeDispatchTests.cpp +++ b/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 { - workGroups : vec3 - } - - @group(0) @binding(0) var output : OutputBuf; + @group(0) @binding(0) var output : vec3; @compute @workgroup_size(1, 1, 1) fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3, @builtin(num_workgroups) dispatch : vec3) { if (dispatch.x == 0u || dispatch.y == 0u || dispatch.z == 0u) { - output.workGroups = vec3(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu); + output = vec3(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu); return; } if (all(GlobalInvocationID == dispatch - vec3(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 { - expectedDispatch : vec3 - } - struct OutputBuf { - workGroups : vec3 - } - - @group(0) @binding(0) var input : InputBuf; - @group(0) @binding(1) var output : OutputBuf; + @group(0) @binding(0) var input : vec3; + @group(0) @binding(1) var output : vec3; @compute @workgroup_size(1, 1, 1) fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3) { - let dispatch : vec3 = input.expectedDispatch; + let dispatch : vec3 = input; if (dispatch.x == 0u || dispatch.y == 0u || dispatch.z == 0u) { - output.workGroups = vec3(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu); + output = vec3(0xFFFFFFFFu, 0xFFFFFFFFu, 0xFFFFFFFFu); return; } if (all(GlobalInvocationID == dispatch - vec3(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 79cacd9be1..bf9ccdc61d 100644 --- a/src/dawn/tests/end2end/ComputeLayoutMemoryBufferTests.cpp +++ b/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(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 407ae7c141..88bb571b8f 100644 --- a/src/dawn/tests/end2end/MultisampledRenderingTests.cpp +++ b/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 kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color + std::array 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 kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color + std::array 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 kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color + std::array 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 kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color + std::array 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 kUniformData = {kGreen.r, kGreen.g, kGreen.b, kGreen.a, // Color + std::array 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 kUniformData = {kRed.r, kRed.g, kRed.b, kRed.a, // color + std::array 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 335f73b8c5..4dd9715826 100644 --- a/src/dawn/tests/perf_tests/DrawCallPerf.cpp +++ b/src/dawn/tests/perf_tests/DrawCallPerf.cpp @@ -41,26 +41,17 @@ constexpr char kVertexShader[] = R"( })"; constexpr char kFragmentShaderA[] = R"( - struct Uniforms { - color : vec3 - } - @group(0) @binding(0) var uniforms : Uniforms; + @group(0) @binding(0) var color : vec3; @fragment fn main() -> @location(0) vec4 { - return vec4(uniforms.color * (1.0 / 5000.0), 1.0); + return vec4(color * (1.0 / 5000.0), 1.0); })"; constexpr char kFragmentShaderB[] = R"( - struct Constants { - color : vec3 - } - struct Uniforms { - color : vec3 - } - @group(0) @binding(0) var constants : Constants; - @group(1) @binding(0) var uniforms : Uniforms; + @group(0) @binding(0) var constant_color : vec3; + @group(1) @binding(0) var uniform_color : vec3; @fragment fn main() -> @location(0) vec4 { - return vec4((constants.color + uniforms.color) * (1.0 / 5000.0), 1.0); + return vec4((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 dbd07c2a97..b09d0c9801 100644 --- a/src/dawn/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp +++ b/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, 3>,", "u32", "d[0][0]", 120}, {0, 3, "e : array>,", "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 bindings = {{0, 0, "", "", "", 12}, {0, 1, "", "", "", 12}}; + + auto MakeShader = [](const char* stageAttributes) { + std::ostringstream ostream; + ostream << "@group(0) @binding(0) var buffer : vec3;\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& 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)); + } + }); +}