Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 1 | // Copyright 2019 The Dawn Authors |
| 2 | // |
| 3 | // Licensed under the Apache License, Version 2.0 (the "License"); |
| 4 | // you may not use this file except in compliance with the License. |
| 5 | // You may obtain a copy of the License at |
| 6 | // |
| 7 | // http://www.apache.org/licenses/LICENSE-2.0 |
| 8 | // |
| 9 | // Unless required by applicable law or agreed to in writing, software |
| 10 | // distributed under the License is distributed on an "AS IS" BASIS, |
| 11 | // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| 12 | // See the License for the specific language governing permissions and |
| 13 | // limitations under the License. |
| 14 | |
| 15 | #include "common/Assert.h" |
| 16 | #include "common/Constants.h" |
| 17 | #include "common/Math.h" |
| 18 | #include "tests/DawnTest.h" |
| 19 | #include "utils/ComboRenderPipelineDescriptor.h" |
| 20 | #include "utils/WGPUHelpers.h" |
| 21 | |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 22 | class GpuMemorySyncTests : public DawnTest { |
| 23 | protected: |
| 24 | wgpu::Buffer CreateBuffer() { |
| 25 | wgpu::BufferDescriptor srcDesc; |
| 26 | srcDesc.size = 4; |
| 27 | srcDesc.usage = |
| 28 | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Storage; |
| 29 | wgpu::Buffer buffer = device.CreateBuffer(&srcDesc); |
| 30 | |
| 31 | int myData = 0; |
Corentin Wallez | 47a3341 | 2020-06-02 09:24:39 +0000 | [diff] [blame] | 32 | queue.WriteBuffer(buffer, 0, &myData, sizeof(myData)); |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 33 | return buffer; |
| 34 | } |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 35 | |
| 36 | std::tuple<wgpu::ComputePipeline, wgpu::BindGroup> CreatePipelineAndBindGroupForCompute( |
| 37 | const wgpu::Buffer& buffer) { |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 38 | wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 39 | [[block]] struct Data { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 40 | a : i32; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 41 | }; |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 42 | [[group(0), binding(0)]] var<storage, read_write> data : Data; |
Sarah | 2a57db7 | 2021-06-23 19:19:06 +0000 | [diff] [blame] | 43 | [[stage(compute), workgroup_size(1)]] fn main() { |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 44 | data.a = data.a + 1; |
| 45 | })"); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 46 | |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 47 | wgpu::ComputePipelineDescriptor cpDesc; |
Brandon Jones | 0d50a2c | 2021-06-09 18:07:32 +0000 | [diff] [blame] | 48 | cpDesc.compute.module = csModule; |
| 49 | cpDesc.compute.entryPoint = "main"; |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 50 | wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cpDesc); |
| 51 | |
| 52 | wgpu::BindGroup bindGroup = |
Yunchao He | 4326a8a | 2019-12-09 19:17:22 +0000 | [diff] [blame] | 53 | utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}}); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 54 | return std::make_tuple(pipeline, bindGroup); |
| 55 | } |
| 56 | |
| 57 | std::tuple<wgpu::RenderPipeline, wgpu::BindGroup> CreatePipelineAndBindGroupForRender( |
| 58 | const wgpu::Buffer& buffer, |
| 59 | wgpu::TextureFormat colorFormat) { |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 60 | wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 61 | [[stage(vertex)]] fn main() -> [[builtin(position)]] vec4<f32> { |
| 62 | return vec4<f32>(0.0, 0.0, 0.0, 1.0); |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 63 | })"); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 64 | |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 65 | wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 66 | [[block]] struct Data { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 67 | i : i32; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 68 | }; |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 69 | [[group(0), binding(0)]] var<storage, read_write> data : Data; |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 70 | [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 71 | data.i = data.i + 1; |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 72 | return vec4<f32>(f32(data.i) / 255.0, 0.0, 0.0, 1.0); |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 73 | })"); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 74 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 75 | utils::ComboRenderPipelineDescriptor rpDesc; |
Brandon Jones | bff9d3a | 2021-03-18 02:54:27 +0000 | [diff] [blame] | 76 | rpDesc.vertex.module = vsModule; |
| 77 | rpDesc.cFragment.module = fsModule; |
| 78 | rpDesc.primitive.topology = wgpu::PrimitiveTopology::PointList; |
| 79 | rpDesc.cTargets[0].format = colorFormat; |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 80 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 81 | wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 82 | |
| 83 | wgpu::BindGroup bindGroup = |
Yunchao He | 4326a8a | 2019-12-09 19:17:22 +0000 | [diff] [blame] | 84 | utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}}); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 85 | return std::make_tuple(pipeline, bindGroup); |
| 86 | } |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 87 | }; |
| 88 | |
| 89 | // Clear storage buffer with zero. Then read data, add one, and write the result to storage buffer |
| 90 | // in compute pass. Iterate this read-add-write steps per compute pass a few time. The successive |
| 91 | // iteration reads the result in buffer from last iteration, which makes the iterations a data |
| 92 | // dependency chain. The test verifies that data in buffer among iterations in compute passes is |
| 93 | // correctly synchronized. |
| 94 | TEST_P(GpuMemorySyncTests, ComputePass) { |
| 95 | // Create pipeline, bind group, and buffer for compute pass. |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 96 | wgpu::Buffer buffer = CreateBuffer(); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 97 | wgpu::ComputePipeline compute; |
| 98 | wgpu::BindGroup bindGroup; |
| 99 | std::tie(compute, bindGroup) = CreatePipelineAndBindGroupForCompute(buffer); |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 100 | wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| 101 | |
| 102 | // Iterate the read-add-write operations in compute pass a few times. |
| 103 | int iteration = 3; |
| 104 | for (int i = 0; i < iteration; ++i) { |
| 105 | wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); |
| 106 | pass.SetPipeline(compute); |
| 107 | pass.SetBindGroup(0, bindGroup); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 108 | pass.Dispatch(1); |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 109 | pass.EndPass(); |
| 110 | } |
| 111 | |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 112 | wgpu::CommandBuffer commands = encoder.Finish(); |
| 113 | queue.Submit(1, &commands); |
| 114 | |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 115 | // Verify the result. |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 116 | EXPECT_BUFFER_U32_EQ(iteration, buffer, 0); |
| 117 | } |
| 118 | |
| 119 | // Clear storage buffer with zero. Then read data, add one, and write the result to storage buffer |
| 120 | // in render pass. Iterate this read-add-write steps per render pass a few time. The successive |
| 121 | // iteration reads the result in buffer from last iteration, which makes the iterations a data |
| 122 | // dependency chain. In addition, color output by fragment shader depends on the data in storage |
| 123 | // buffer, so we can check color in render target to verify that data in buffer among iterations in |
| 124 | // render passes is correctly synchronized. |
| 125 | TEST_P(GpuMemorySyncTests, RenderPass) { |
| 126 | // Create pipeline, bind group, and buffer for render pass. |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 127 | wgpu::Buffer buffer = CreateBuffer(); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 128 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); |
| 129 | wgpu::RenderPipeline render; |
| 130 | wgpu::BindGroup bindGroup; |
| 131 | std::tie(render, bindGroup) = |
| 132 | CreatePipelineAndBindGroupForRender(buffer, renderPass.colorFormat); |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 133 | wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| 134 | |
| 135 | // Iterate the read-add-write operations in render pass a few times. |
| 136 | int iteration = 3; |
| 137 | for (int i = 0; i < iteration; ++i) { |
| 138 | wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); |
| 139 | pass.SetPipeline(render); |
| 140 | pass.SetBindGroup(0, bindGroup); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 141 | pass.Draw(1); |
Yunchao He | 02dd733 | 2019-11-20 00:05:20 +0000 | [diff] [blame] | 142 | pass.EndPass(); |
| 143 | } |
| 144 | |
| 145 | wgpu::CommandBuffer commands = encoder.Finish(); |
| 146 | queue.Submit(1, &commands); |
| 147 | |
| 148 | // Verify the result. |
| 149 | EXPECT_PIXEL_RGBA8_EQ(RGBA8(iteration, 0, 0, 255), renderPass.color, 0, 0); |
| 150 | } |
| 151 | |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 152 | // Write into a storage buffer in a render pass. Then read that data in a compute |
| 153 | // pass. And verify the data flow is correctly synchronized. |
| 154 | TEST_P(GpuMemorySyncTests, RenderPassToComputePass) { |
| 155 | // Create pipeline, bind group, and buffer for render pass and compute pass. |
| 156 | wgpu::Buffer buffer = CreateBuffer(); |
| 157 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); |
| 158 | wgpu::RenderPipeline render; |
| 159 | wgpu::BindGroup bindGroup0; |
| 160 | std::tie(render, bindGroup0) = |
| 161 | CreatePipelineAndBindGroupForRender(buffer, renderPass.colorFormat); |
| 162 | |
| 163 | wgpu::ComputePipeline compute; |
| 164 | wgpu::BindGroup bindGroup1; |
| 165 | std::tie(compute, bindGroup1) = CreatePipelineAndBindGroupForCompute(buffer); |
| 166 | |
| 167 | wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| 168 | |
| 169 | // Write data into a storage buffer in render pass. |
| 170 | wgpu::RenderPassEncoder pass0 = encoder.BeginRenderPass(&renderPass.renderPassInfo); |
| 171 | pass0.SetPipeline(render); |
| 172 | pass0.SetBindGroup(0, bindGroup0); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 173 | pass0.Draw(1); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 174 | pass0.EndPass(); |
| 175 | |
Yunchao He | d28b578 | 2019-12-19 18:50:18 +0000 | [diff] [blame] | 176 | // Read that data in compute pass. |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 177 | wgpu::ComputePassEncoder pass1 = encoder.BeginComputePass(); |
| 178 | pass1.SetPipeline(compute); |
| 179 | pass1.SetBindGroup(0, bindGroup1); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 180 | pass1.Dispatch(1); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 181 | pass1.EndPass(); |
| 182 | |
| 183 | wgpu::CommandBuffer commands = encoder.Finish(); |
| 184 | queue.Submit(1, &commands); |
| 185 | |
| 186 | // Verify the result. |
| 187 | EXPECT_BUFFER_U32_EQ(2, buffer, 0); |
| 188 | } |
| 189 | |
| 190 | // Write into a storage buffer in a compute pass. Then read that data in a render |
| 191 | // pass. And verify the data flow is correctly synchronized. |
| 192 | TEST_P(GpuMemorySyncTests, ComputePassToRenderPass) { |
| 193 | // Create pipeline, bind group, and buffer for compute pass and render pass. |
| 194 | wgpu::Buffer buffer = CreateBuffer(); |
| 195 | wgpu::ComputePipeline compute; |
| 196 | wgpu::BindGroup bindGroup1; |
| 197 | std::tie(compute, bindGroup1) = CreatePipelineAndBindGroupForCompute(buffer); |
| 198 | |
| 199 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); |
| 200 | wgpu::RenderPipeline render; |
| 201 | wgpu::BindGroup bindGroup0; |
| 202 | std::tie(render, bindGroup0) = |
| 203 | CreatePipelineAndBindGroupForRender(buffer, renderPass.colorFormat); |
| 204 | |
| 205 | wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| 206 | |
Yunchao He | d28b578 | 2019-12-19 18:50:18 +0000 | [diff] [blame] | 207 | // Write data into a storage buffer in compute pass. |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 208 | wgpu::ComputePassEncoder pass0 = encoder.BeginComputePass(); |
| 209 | pass0.SetPipeline(compute); |
| 210 | pass0.SetBindGroup(0, bindGroup1); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 211 | pass0.Dispatch(1); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 212 | pass0.EndPass(); |
| 213 | |
| 214 | // Read that data in render pass. |
| 215 | wgpu::RenderPassEncoder pass1 = encoder.BeginRenderPass(&renderPass.renderPassInfo); |
| 216 | pass1.SetPipeline(render); |
| 217 | pass1.SetBindGroup(0, bindGroup0); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 218 | pass1.Draw(1); |
Yunchao He | 733842c | 2019-11-23 00:20:53 +0000 | [diff] [blame] | 219 | pass1.EndPass(); |
| 220 | |
| 221 | wgpu::CommandBuffer commands = encoder.Finish(); |
| 222 | queue.Submit(1, &commands); |
| 223 | |
| 224 | // Verify the result. |
| 225 | EXPECT_PIXEL_RGBA8_EQ(RGBA8(2, 0, 0, 255), renderPass.color, 0, 0); |
| 226 | } |
| 227 | |
Kai Ninomiya | 2afea0c | 2020-07-10 20:33:08 +0000 | [diff] [blame] | 228 | DAWN_INSTANTIATE_TEST(GpuMemorySyncTests, |
| 229 | D3D12Backend(), |
| 230 | MetalBackend(), |
| 231 | OpenGLBackend(), |
Stephen White | f31b78e | 2020-12-04 15:59:29 +0000 | [diff] [blame] | 232 | OpenGLESBackend(), |
Kai Ninomiya | 2afea0c | 2020-07-10 20:33:08 +0000 | [diff] [blame] | 233 | VulkanBackend()); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 234 | |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 235 | class StorageToUniformSyncTests : public DawnTest { |
| 236 | protected: |
| 237 | void CreateBuffer() { |
| 238 | wgpu::BufferDescriptor bufferDesc; |
| 239 | bufferDesc.size = sizeof(float); |
| 240 | bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::Uniform; |
| 241 | mBuffer = device.CreateBuffer(&bufferDesc); |
| 242 | } |
| 243 | |
| 244 | std::tuple<wgpu::ComputePipeline, wgpu::BindGroup> CreatePipelineAndBindGroupForCompute() { |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 245 | wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 246 | [[block]] struct Data { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 247 | a : f32; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 248 | }; |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 249 | [[group(0), binding(0)]] var<storage, read_write> data : Data; |
Sarah | 2a57db7 | 2021-06-23 19:19:06 +0000 | [diff] [blame] | 250 | [[stage(compute), workgroup_size(1)]] fn main() { |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 251 | data.a = 1.0; |
| 252 | })"); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 253 | |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 254 | wgpu::ComputePipelineDescriptor cpDesc; |
Brandon Jones | 0d50a2c | 2021-06-09 18:07:32 +0000 | [diff] [blame] | 255 | cpDesc.compute.module = csModule; |
| 256 | cpDesc.compute.entryPoint = "main"; |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 257 | wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cpDesc); |
| 258 | |
| 259 | wgpu::BindGroup bindGroup = |
Yunchao He | 4326a8a | 2019-12-09 19:17:22 +0000 | [diff] [blame] | 260 | utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, mBuffer}}); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 261 | return std::make_tuple(pipeline, bindGroup); |
| 262 | } |
| 263 | |
| 264 | std::tuple<wgpu::RenderPipeline, wgpu::BindGroup> CreatePipelineAndBindGroupForRender( |
| 265 | wgpu::TextureFormat colorFormat) { |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 266 | wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 267 | [[stage(vertex)]] fn main() -> [[builtin(position)]] vec4<f32> { |
| 268 | return vec4<f32>(0.0, 0.0, 0.0, 1.0); |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 269 | })"); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 270 | |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 271 | wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 272 | [[block]] struct Contents { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 273 | color : f32; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 274 | }; |
dan sinclair | 0f9c2d7 | 2021-01-19 14:18:51 +0000 | [diff] [blame] | 275 | [[group(0), binding(0)]] var<uniform> contents : Contents; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 276 | |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 277 | [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { |
| 278 | return vec4<f32>(contents.color, 0.0, 0.0, 1.0); |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 279 | })"); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 280 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 281 | utils::ComboRenderPipelineDescriptor rpDesc; |
Brandon Jones | bff9d3a | 2021-03-18 02:54:27 +0000 | [diff] [blame] | 282 | rpDesc.vertex.module = vsModule; |
| 283 | rpDesc.cFragment.module = fsModule; |
| 284 | rpDesc.primitive.topology = wgpu::PrimitiveTopology::PointList; |
| 285 | rpDesc.cTargets[0].format = colorFormat; |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 286 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 287 | wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 288 | |
| 289 | wgpu::BindGroup bindGroup = |
Yunchao He | 4326a8a | 2019-12-09 19:17:22 +0000 | [diff] [blame] | 290 | utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, mBuffer}}); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 291 | return std::make_tuple(pipeline, bindGroup); |
| 292 | } |
| 293 | |
| 294 | wgpu::Buffer mBuffer; |
| 295 | }; |
| 296 | |
| 297 | // Write into a storage buffer in compute pass in a command buffer. Then read that data in a render |
| 298 | // pass. The two passes use the same command buffer. |
| 299 | TEST_P(StorageToUniformSyncTests, ReadAfterWriteWithSameCommandBuffer) { |
| 300 | // Create pipeline, bind group, and buffer for compute pass and render pass. |
| 301 | CreateBuffer(); |
| 302 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); |
| 303 | wgpu::ComputePipeline compute; |
| 304 | wgpu::BindGroup computeBindGroup; |
| 305 | std::tie(compute, computeBindGroup) = CreatePipelineAndBindGroupForCompute(); |
| 306 | wgpu::RenderPipeline render; |
| 307 | wgpu::BindGroup renderBindGroup; |
| 308 | std::tie(render, renderBindGroup) = CreatePipelineAndBindGroupForRender(renderPass.colorFormat); |
| 309 | |
| 310 | // Write data into a storage buffer in compute pass. |
| 311 | wgpu::CommandEncoder encoder0 = device.CreateCommandEncoder(); |
| 312 | wgpu::ComputePassEncoder pass0 = encoder0.BeginComputePass(); |
| 313 | pass0.SetPipeline(compute); |
| 314 | pass0.SetBindGroup(0, computeBindGroup); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 315 | pass0.Dispatch(1); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 316 | pass0.EndPass(); |
| 317 | |
| 318 | // Read that data in render pass. |
| 319 | wgpu::RenderPassEncoder pass1 = encoder0.BeginRenderPass(&renderPass.renderPassInfo); |
| 320 | pass1.SetPipeline(render); |
| 321 | pass1.SetBindGroup(0, renderBindGroup); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 322 | pass1.Draw(1); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 323 | pass1.EndPass(); |
| 324 | |
| 325 | wgpu::CommandBuffer commands = encoder0.Finish(); |
| 326 | queue.Submit(1, &commands); |
| 327 | |
| 328 | // Verify the rendering result. |
Yunchao He | 0c02f54 | 2019-11-19 17:57:30 +0000 | [diff] [blame] | 329 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kRed, renderPass.color, 0, 0); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 330 | } |
| 331 | |
| 332 | // Write into a storage buffer in compute pass in a command buffer. Then read that data in a render |
| 333 | // pass. The two passes use the different command buffers. The command buffers are submitted to the |
| 334 | // queue in one shot. |
| 335 | TEST_P(StorageToUniformSyncTests, ReadAfterWriteWithDifferentCommandBuffers) { |
| 336 | // Create pipeline, bind group, and buffer for compute pass and render pass. |
| 337 | CreateBuffer(); |
| 338 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); |
| 339 | wgpu::ComputePipeline compute; |
| 340 | wgpu::BindGroup computeBindGroup; |
| 341 | std::tie(compute, computeBindGroup) = CreatePipelineAndBindGroupForCompute(); |
| 342 | wgpu::RenderPipeline render; |
| 343 | wgpu::BindGroup renderBindGroup; |
| 344 | std::tie(render, renderBindGroup) = CreatePipelineAndBindGroupForRender(renderPass.colorFormat); |
| 345 | |
| 346 | // Write data into a storage buffer in compute pass. |
| 347 | wgpu::CommandBuffer cb[2]; |
| 348 | wgpu::CommandEncoder encoder0 = device.CreateCommandEncoder(); |
| 349 | wgpu::ComputePassEncoder pass0 = encoder0.BeginComputePass(); |
| 350 | pass0.SetPipeline(compute); |
| 351 | pass0.SetBindGroup(0, computeBindGroup); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 352 | pass0.Dispatch(1); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 353 | pass0.EndPass(); |
| 354 | cb[0] = encoder0.Finish(); |
| 355 | |
| 356 | // Read that data in render pass. |
| 357 | wgpu::CommandEncoder encoder1 = device.CreateCommandEncoder(); |
| 358 | wgpu::RenderPassEncoder pass1 = encoder1.BeginRenderPass(&renderPass.renderPassInfo); |
| 359 | pass1.SetPipeline(render); |
| 360 | pass1.SetBindGroup(0, renderBindGroup); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 361 | pass1.Draw(1); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 362 | pass1.EndPass(); |
| 363 | |
| 364 | cb[1] = encoder1.Finish(); |
| 365 | queue.Submit(2, cb); |
| 366 | |
| 367 | // Verify the rendering result. |
Yunchao He | 0c02f54 | 2019-11-19 17:57:30 +0000 | [diff] [blame] | 368 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kRed, renderPass.color, 0, 0); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 369 | } |
| 370 | |
| 371 | // Write into a storage buffer in compute pass in a command buffer. Then read that data in a render |
| 372 | // pass. The two passes use the different command buffers. The command buffers are submitted to the |
| 373 | // queue separately. |
| 374 | TEST_P(StorageToUniformSyncTests, ReadAfterWriteWithDifferentQueueSubmits) { |
| 375 | // Create pipeline, bind group, and buffer for compute pass and render pass. |
| 376 | CreateBuffer(); |
| 377 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); |
| 378 | wgpu::ComputePipeline compute; |
| 379 | wgpu::BindGroup computeBindGroup; |
| 380 | std::tie(compute, computeBindGroup) = CreatePipelineAndBindGroupForCompute(); |
| 381 | wgpu::RenderPipeline render; |
| 382 | wgpu::BindGroup renderBindGroup; |
| 383 | std::tie(render, renderBindGroup) = CreatePipelineAndBindGroupForRender(renderPass.colorFormat); |
| 384 | |
| 385 | // Write data into a storage buffer in compute pass. |
| 386 | wgpu::CommandBuffer cb[2]; |
| 387 | wgpu::CommandEncoder encoder0 = device.CreateCommandEncoder(); |
| 388 | wgpu::ComputePassEncoder pass0 = encoder0.BeginComputePass(); |
| 389 | pass0.SetPipeline(compute); |
| 390 | pass0.SetBindGroup(0, computeBindGroup); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 391 | pass0.Dispatch(1); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 392 | pass0.EndPass(); |
| 393 | cb[0] = encoder0.Finish(); |
| 394 | queue.Submit(1, &cb[0]); |
| 395 | |
| 396 | // Read that data in render pass. |
| 397 | wgpu::CommandEncoder encoder1 = device.CreateCommandEncoder(); |
| 398 | wgpu::RenderPassEncoder pass1 = encoder1.BeginRenderPass(&renderPass.renderPassInfo); |
| 399 | pass1.SetPipeline(render); |
| 400 | pass1.SetBindGroup(0, renderBindGroup); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 401 | pass1.Draw(1); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 402 | pass1.EndPass(); |
| 403 | |
| 404 | cb[1] = encoder1.Finish(); |
| 405 | queue.Submit(1, &cb[1]); |
| 406 | |
| 407 | // Verify the rendering result. |
Yunchao He | 0c02f54 | 2019-11-19 17:57:30 +0000 | [diff] [blame] | 408 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kRed, renderPass.color, 0, 0); |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 409 | } |
| 410 | |
Yunchao He | 612a63a | 2019-11-18 04:28:24 +0000 | [diff] [blame] | 411 | DAWN_INSTANTIATE_TEST(StorageToUniformSyncTests, |
Austin Eng | 6c1d646 | 2020-02-25 16:23:17 +0000 | [diff] [blame] | 412 | D3D12Backend(), |
| 413 | MetalBackend(), |
| 414 | OpenGLBackend(), |
Stephen White | f31b78e | 2020-12-04 15:59:29 +0000 | [diff] [blame] | 415 | OpenGLESBackend(), |
Austin Eng | 6c1d646 | 2020-02-25 16:23:17 +0000 | [diff] [blame] | 416 | VulkanBackend()); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 417 | |
| 418 | constexpr int kRTSize = 8; |
| 419 | constexpr int kVertexBufferStride = 4 * sizeof(float); |
| 420 | |
| 421 | class MultipleWriteThenMultipleReadTests : public DawnTest { |
| 422 | protected: |
| 423 | wgpu::Buffer CreateZeroedBuffer(uint64_t size, wgpu::BufferUsage usage) { |
| 424 | wgpu::BufferDescriptor srcDesc; |
| 425 | srcDesc.size = size; |
| 426 | srcDesc.usage = usage; |
| 427 | wgpu::Buffer buffer = device.CreateBuffer(&srcDesc); |
| 428 | |
| 429 | std::vector<uint8_t> zeros(size, 0); |
Corentin Wallez | 47a3341 | 2020-06-02 09:24:39 +0000 | [diff] [blame] | 430 | queue.WriteBuffer(buffer, 0, zeros.data(), size); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 431 | |
| 432 | return buffer; |
| 433 | } |
| 434 | }; |
| 435 | |
| 436 | // Write into a few storage buffers in compute pass. Then read that data in a render pass. The |
Yunchao He | ce8bf12 | 2019-12-05 21:18:12 +0000 | [diff] [blame] | 437 | // readonly buffers in render pass include vertex buffer, index buffer, uniform buffer, and readonly |
| 438 | // storage buffer. Data to be read in all of these buffers in render pass depend on the write |
| 439 | // operation in compute pass. |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 440 | TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) { |
| 441 | // Create pipeline, bind group, and different buffers for compute pass. |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 442 | wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 443 | [[block]] struct VBContents { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 444 | pos : array<vec4<f32>, 4>; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 445 | }; |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 446 | [[group(0), binding(0)]] var<storage, read_write> vbContents : VBContents; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 447 | |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 448 | [[block]] struct IBContents { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 449 | indices : array<vec4<i32>, 2>; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 450 | }; |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 451 | [[group(0), binding(1)]] var<storage, read_write> ibContents : IBContents; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 452 | |
Ben Clayton | 4773e8d | 2021-07-13 15:21:07 +0000 | [diff] [blame] | 453 | [[block]] struct ColorContents { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 454 | color : f32; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 455 | }; |
Ben Clayton | 4773e8d | 2021-07-13 15:21:07 +0000 | [diff] [blame] | 456 | [[group(0), binding(2)]] var<storage, read_write> uniformContents : ColorContents; |
| 457 | [[group(0), binding(3)]] var<storage, read_write> storageContents : ColorContents; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 458 | |
Sarah | 2a57db7 | 2021-06-23 19:19:06 +0000 | [diff] [blame] | 459 | [[stage(compute), workgroup_size(1)]] fn main() { |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 460 | vbContents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0); |
| 461 | vbContents.pos[1] = vec4<f32>(1.0, 1.0, 0.0, 1.0); |
| 462 | vbContents.pos[2] = vec4<f32>(1.0, -1.0, 0.0, 1.0); |
| 463 | vbContents.pos[3] = vec4<f32>(-1.0, -1.0, 0.0, 1.0); |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 464 | let dummy : i32 = 0; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 465 | ibContents.indices[0] = vec4<i32>(0, 1, 2, 0); |
| 466 | ibContents.indices[1] = vec4<i32>(2, 3, dummy, dummy); |
| 467 | uniformContents.color = 1.0; |
| 468 | storageContents.color = 1.0; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 469 | })"); |
| 470 | |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 471 | wgpu::ComputePipelineDescriptor cpDesc; |
Brandon Jones | 0d50a2c | 2021-06-09 18:07:32 +0000 | [diff] [blame] | 472 | cpDesc.compute.module = csModule; |
| 473 | cpDesc.compute.entryPoint = "main"; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 474 | wgpu::ComputePipeline cp = device.CreateComputePipeline(&cpDesc); |
| 475 | wgpu::Buffer vertexBuffer = CreateZeroedBuffer( |
| 476 | kVertexBufferStride * 4, |
| 477 | wgpu::BufferUsage::Vertex | wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst); |
| 478 | wgpu::Buffer indexBuffer = CreateZeroedBuffer( |
| 479 | sizeof(int) * 4 * 2, |
| 480 | wgpu::BufferUsage::Index | wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst); |
Yunchao He | ce8bf12 | 2019-12-05 21:18:12 +0000 | [diff] [blame] | 481 | wgpu::Buffer uniformBuffer = |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 482 | CreateZeroedBuffer(sizeof(float), wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Storage | |
| 483 | wgpu::BufferUsage::CopyDst); |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 484 | wgpu::Buffer storageBuffer = |
Yunchao He | ce8bf12 | 2019-12-05 21:18:12 +0000 | [diff] [blame] | 485 | CreateZeroedBuffer(sizeof(float), wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 486 | |
Yunchao He | 4326a8a | 2019-12-09 19:17:22 +0000 | [diff] [blame] | 487 | wgpu::BindGroup bindGroup0 = utils::MakeBindGroup( |
| 488 | device, cp.GetBindGroupLayout(0), |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 489 | {{0, vertexBuffer}, {1, indexBuffer}, {2, uniformBuffer}, {3, storageBuffer}}); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 490 | // Write data into storage buffers in compute pass. |
| 491 | wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| 492 | wgpu::ComputePassEncoder pass0 = encoder.BeginComputePass(); |
| 493 | pass0.SetPipeline(cp); |
| 494 | pass0.SetBindGroup(0, bindGroup0); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 495 | pass0.Dispatch(1); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 496 | pass0.EndPass(); |
| 497 | |
| 498 | // Create pipeline, bind group, and reuse buffers in render pass. |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 499 | wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 500 | [[stage(vertex)]] |
| 501 | fn main([[location(0)]] pos : vec4<f32>) -> [[builtin(position)]] vec4<f32> { |
| 502 | return pos; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 503 | })"); |
| 504 | |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 505 | wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 506 | [[block]] struct Buf { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 507 | color : f32; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 508 | }; |
| 509 | |
dan sinclair | 0f9c2d7 | 2021-01-19 14:18:51 +0000 | [diff] [blame] | 510 | [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf; |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 511 | [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 512 | |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 513 | [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { |
| 514 | return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 515 | })"); |
| 516 | |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 517 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize); |
| 518 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 519 | utils::ComboRenderPipelineDescriptor rpDesc; |
Brandon Jones | bff9d3a | 2021-03-18 02:54:27 +0000 | [diff] [blame] | 520 | rpDesc.vertex.module = vsModule; |
| 521 | rpDesc.cFragment.module = fsModule; |
| 522 | rpDesc.primitive.topology = wgpu::PrimitiveTopology::TriangleList; |
| 523 | rpDesc.vertex.bufferCount = 1; |
| 524 | rpDesc.cBuffers[0].arrayStride = kVertexBufferStride; |
| 525 | rpDesc.cBuffers[0].attributeCount = 1; |
| 526 | rpDesc.cAttributes[0].format = wgpu::VertexFormat::Float32x4; |
| 527 | rpDesc.cTargets[0].format = renderPass.colorFormat; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 528 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 529 | wgpu::RenderPipeline rp = device.CreateRenderPipeline(&rpDesc); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 530 | |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 531 | wgpu::BindGroup bindGroup1 = utils::MakeBindGroup(device, rp.GetBindGroupLayout(0), |
| 532 | {{0, uniformBuffer}, {1, storageBuffer}}); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 533 | |
| 534 | // Read data in buffers in render pass. |
| 535 | wgpu::RenderPassEncoder pass1 = encoder.BeginRenderPass(&renderPass.renderPassInfo); |
| 536 | pass1.SetPipeline(rp); |
| 537 | pass1.SetVertexBuffer(0, vertexBuffer); |
Corentin Wallez | 5fad85b | 2020-11-25 08:54:14 +0000 | [diff] [blame] | 538 | pass1.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 539 | pass1.SetBindGroup(0, bindGroup1); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 540 | pass1.DrawIndexed(6); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 541 | pass1.EndPass(); |
| 542 | |
| 543 | wgpu::CommandBuffer commandBuffer = encoder.Finish(); |
| 544 | queue.Submit(1, &commandBuffer); |
| 545 | |
| 546 | // Verify the rendering result. |
Yunchao He | 4eb40c1 | 2021-03-31 22:15:53 +0000 | [diff] [blame] | 547 | uint32_t min = 1, max = kRTSize - 3; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 548 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, min, min); |
| 549 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, max, min); |
| 550 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, min, max); |
| 551 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, max, max); |
| 552 | } |
| 553 | |
| 554 | // Write into a storage buffer in compute pass. Then read that data in buffer in a render pass. The |
Yunchao He | ce8bf12 | 2019-12-05 21:18:12 +0000 | [diff] [blame] | 555 | // buffer is composed of vertices, indices, uniforms and readonly storage. Data to be read in the |
| 556 | // buffer in render pass depend on the write operation in compute pass. |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 557 | TEST_P(MultipleWriteThenMultipleReadTests, OneBuffer) { |
Stephen White | 032500b | 2021-01-26 15:01:18 +0000 | [diff] [blame] | 558 | // TODO(crbug.com/dawn/646): diagnose and fix this OpenGL ES failure. |
| 559 | // "Push constant block cannot be expressed as neither std430 nor std140. ES-targets do not |
| 560 | // support GL_ARB_enhanced_layouts." |
Jiawei Shao | 44fc6e3 | 2021-05-26 01:04:32 +0000 | [diff] [blame] | 561 | DAWN_SUPPRESS_TEST_IF(IsOpenGLES()); |
Stephen White | 032500b | 2021-01-26 15:01:18 +0000 | [diff] [blame] | 562 | |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 563 | // Create pipeline, bind group, and a complex buffer for compute pass. |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 564 | wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 565 | [[block]] struct Contents { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 566 | [[align(256)]] pos : array<vec4<f32>, 4>; |
| 567 | [[align(256)]] indices : array<vec4<i32>, 2>; |
| 568 | [[align(256)]] color0 : f32; |
| 569 | [[align(256)]] color1 : f32; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 570 | }; |
| 571 | |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 572 | [[group(0), binding(0)]] var<storage, read_write> contents : Contents; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 573 | |
Sarah | 2a57db7 | 2021-06-23 19:19:06 +0000 | [diff] [blame] | 574 | [[stage(compute), workgroup_size(1)]] fn main() { |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 575 | contents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0); |
| 576 | contents.pos[1] = vec4<f32>(1.0, 1.0, 0.0, 1.0); |
| 577 | contents.pos[2] = vec4<f32>(1.0, -1.0, 0.0, 1.0); |
| 578 | contents.pos[3] = vec4<f32>(-1.0, -1.0, 0.0, 1.0); |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 579 | let dummy : i32 = 0; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 580 | contents.indices[0] = vec4<i32>(0, 1, 2, 0); |
| 581 | contents.indices[1] = vec4<i32>(2, 3, dummy, dummy); |
| 582 | contents.color0 = 1.0; |
| 583 | contents.color1 = 1.0; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 584 | })"); |
| 585 | |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 586 | wgpu::ComputePipelineDescriptor cpDesc; |
Brandon Jones | 0d50a2c | 2021-06-09 18:07:32 +0000 | [diff] [blame] | 587 | cpDesc.compute.module = csModule; |
| 588 | cpDesc.compute.entryPoint = "main"; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 589 | wgpu::ComputePipeline cp = device.CreateComputePipeline(&cpDesc); |
| 590 | struct Data { |
| 591 | float pos[4][4]; |
| 592 | char padding0[256 - sizeof(float) * 16]; |
| 593 | int indices[2][4]; |
| 594 | char padding1[256 - sizeof(int) * 8]; |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 595 | float color0; |
| 596 | char padding2[256 - sizeof(float)]; |
| 597 | float color1; |
James Price | a0b31e0 | 2021-06-28 08:38:28 +0000 | [diff] [blame] | 598 | char padding3[256 - sizeof(float)]; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 599 | }; |
| 600 | wgpu::Buffer buffer = CreateZeroedBuffer( |
| 601 | sizeof(Data), wgpu::BufferUsage::Vertex | wgpu::BufferUsage::Index | |
| 602 | wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Storage | |
| 603 | wgpu::BufferUsage::CopyDst); |
Yunchao He | 4326a8a | 2019-12-09 19:17:22 +0000 | [diff] [blame] | 604 | wgpu::BindGroup bindGroup0 = |
| 605 | utils::MakeBindGroup(device, cp.GetBindGroupLayout(0), {{0, buffer}}); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 606 | |
| 607 | // Write various data (vertices, indices, and uniforms) into the buffer in compute pass. |
| 608 | wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| 609 | wgpu::ComputePassEncoder pass0 = encoder.BeginComputePass(); |
| 610 | pass0.SetPipeline(cp); |
| 611 | pass0.SetBindGroup(0, bindGroup0); |
Corentin Wallez | 3da19b8 | 2020-03-31 16:23:35 +0000 | [diff] [blame] | 612 | pass0.Dispatch(1); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 613 | pass0.EndPass(); |
| 614 | |
| 615 | // Create pipeline, bind group, and reuse the buffer in render pass. |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 616 | wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 617 | [[stage(vertex)]] |
| 618 | fn main([[location(0)]] pos : vec4<f32>) -> [[builtin(position)]] vec4<f32> { |
| 619 | return pos; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 620 | })"); |
| 621 | |
Corentin Wallez | 7aec4ae | 2021-03-24 15:55:32 +0000 | [diff] [blame] | 622 | wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"( |
Austin Eng | b3ab21e | 2020-12-23 19:42:00 +0000 | [diff] [blame] | 623 | [[block]] struct Buf { |
Ben Clayton | c568684 | 2021-03-17 09:48:19 +0000 | [diff] [blame] | 624 | color : f32; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 625 | }; |
dan sinclair | 0f9c2d7 | 2021-01-19 14:18:51 +0000 | [diff] [blame] | 626 | [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf; |
Ben Clayton | 15eba9a | 2021-06-08 15:36:44 +0000 | [diff] [blame] | 627 | [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 628 | |
Brandon Jones | e87ea2b | 2021-04-14 17:05:07 +0000 | [diff] [blame] | 629 | [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { |
| 630 | return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 631 | })"); |
| 632 | |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 633 | utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize); |
| 634 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 635 | utils::ComboRenderPipelineDescriptor rpDesc; |
Brandon Jones | bff9d3a | 2021-03-18 02:54:27 +0000 | [diff] [blame] | 636 | rpDesc.vertex.module = vsModule; |
| 637 | rpDesc.cFragment.module = fsModule; |
| 638 | rpDesc.primitive.topology = wgpu::PrimitiveTopology::TriangleList; |
| 639 | rpDesc.vertex.bufferCount = 1; |
| 640 | rpDesc.cBuffers[0].arrayStride = kVertexBufferStride; |
| 641 | rpDesc.cBuffers[0].attributeCount = 1; |
| 642 | rpDesc.cAttributes[0].format = wgpu::VertexFormat::Float32x4; |
| 643 | rpDesc.cTargets[0].format = renderPass.colorFormat; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 644 | |
Brandon Jones | 41c87d9 | 2021-05-21 05:01:38 +0000 | [diff] [blame] | 645 | wgpu::RenderPipeline rp = device.CreateRenderPipeline(&rpDesc); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 646 | |
| 647 | wgpu::BindGroup bindGroup1 = |
Yunchao He | 4326a8a | 2019-12-09 19:17:22 +0000 | [diff] [blame] | 648 | utils::MakeBindGroup(device, rp.GetBindGroupLayout(0), |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 649 | {{0, buffer, offsetof(Data, color0), sizeof(float)}, |
| 650 | {1, buffer, offsetof(Data, color1), sizeof(float)}}); |
| 651 | |
| 652 | // Read various data in the buffer in render pass. |
| 653 | wgpu::RenderPassEncoder pass1 = encoder.BeginRenderPass(&renderPass.renderPassInfo); |
| 654 | pass1.SetPipeline(rp); |
| 655 | pass1.SetVertexBuffer(0, buffer); |
Corentin Wallez | 5fad85b | 2020-11-25 08:54:14 +0000 | [diff] [blame] | 656 | pass1.SetIndexBuffer(buffer, wgpu::IndexFormat::Uint32, offsetof(Data, indices)); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 657 | pass1.SetBindGroup(0, bindGroup1); |
Corentin Wallez | 67b1ad7 | 2020-03-31 16:21:35 +0000 | [diff] [blame] | 658 | pass1.DrawIndexed(6); |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 659 | pass1.EndPass(); |
| 660 | |
| 661 | wgpu::CommandBuffer commandBuffer = encoder.Finish(); |
| 662 | queue.Submit(1, &commandBuffer); |
| 663 | |
| 664 | // Verify the rendering result. |
Yunchao He | 4eb40c1 | 2021-03-31 22:15:53 +0000 | [diff] [blame] | 665 | uint32_t min = 1, max = kRTSize - 3; |
Yunchao He | 40b10e4 | 2019-11-28 18:55:45 +0000 | [diff] [blame] | 666 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, min, min); |
| 667 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, max, min); |
| 668 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, min, max); |
| 669 | EXPECT_PIXEL_RGBA8_EQ(RGBA8::kYellow, renderPass.color, max, max); |
| 670 | } |
| 671 | |
Yunchao He | ef8dee9 | 2019-12-09 21:35:38 +0000 | [diff] [blame] | 672 | DAWN_INSTANTIATE_TEST(MultipleWriteThenMultipleReadTests, |
Austin Eng | 6c1d646 | 2020-02-25 16:23:17 +0000 | [diff] [blame] | 673 | D3D12Backend(), |
| 674 | MetalBackend(), |
| 675 | OpenGLBackend(), |
Stephen White | f31b78e | 2020-12-04 15:59:29 +0000 | [diff] [blame] | 676 | OpenGLESBackend(), |
Austin Eng | 6c1d646 | 2020-02-25 16:23:17 +0000 | [diff] [blame] | 677 | VulkanBackend()); |