blob: a6fe1418d49b26fba37dd3361e6a0a224a63a1fc [file] [log] [blame]
Yunchao He612a63a2019-11-18 04:28:24 +00001// 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 He02dd7332019-11-20 00:05:20 +000022class 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 Wallez47a33412020-06-02 09:24:39 +000032 queue.WriteBuffer(buffer, 0, &myData, sizeof(myData));
Yunchao He02dd7332019-11-20 00:05:20 +000033 return buffer;
34 }
Yunchao He733842c2019-11-23 00:20:53 +000035
36 std::tuple<wgpu::ComputePipeline, wgpu::BindGroup> CreatePipelineAndBindGroupForCompute(
37 const wgpu::Buffer& buffer) {
Corentin Wallez7aec4ae2021-03-24 15:55:32 +000038 wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +000039 [[block]] struct Data {
Ben Claytonc5686842021-03-17 09:48:19 +000040 a : i32;
Austin Engb3ab21e2020-12-23 19:42:00 +000041 };
Ben Clayton15eba9a2021-06-08 15:36:44 +000042 [[group(0), binding(0)]] var<storage, read_write> data : Data;
Sarah2a57db72021-06-23 19:19:06 +000043 [[stage(compute), workgroup_size(1)]] fn main() {
Austin Engb3ab21e2020-12-23 19:42:00 +000044 data.a = data.a + 1;
45 })");
Yunchao He733842c2019-11-23 00:20:53 +000046
Yunchao He733842c2019-11-23 00:20:53 +000047 wgpu::ComputePipelineDescriptor cpDesc;
Brandon Jones0d50a2c2021-06-09 18:07:32 +000048 cpDesc.compute.module = csModule;
49 cpDesc.compute.entryPoint = "main";
Yunchao He733842c2019-11-23 00:20:53 +000050 wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cpDesc);
51
52 wgpu::BindGroup bindGroup =
Yunchao He4326a8a2019-12-09 19:17:22 +000053 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
Yunchao He733842c2019-11-23 00:20:53 +000054 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 Wallez7aec4ae2021-03-24 15:55:32 +000060 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
Brandon Jonese87ea2b2021-04-14 17:05:07 +000061 [[stage(vertex)]] fn main() -> [[builtin(position)]] vec4<f32> {
62 return vec4<f32>(0.0, 0.0, 0.0, 1.0);
Austin Engb3ab21e2020-12-23 19:42:00 +000063 })");
Yunchao He733842c2019-11-23 00:20:53 +000064
Corentin Wallez7aec4ae2021-03-24 15:55:32 +000065 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +000066 [[block]] struct Data {
Ben Claytonc5686842021-03-17 09:48:19 +000067 i : i32;
Austin Engb3ab21e2020-12-23 19:42:00 +000068 };
Ben Clayton15eba9a2021-06-08 15:36:44 +000069 [[group(0), binding(0)]] var<storage, read_write> data : Data;
Brandon Jonese87ea2b2021-04-14 17:05:07 +000070 [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
Austin Engb3ab21e2020-12-23 19:42:00 +000071 data.i = data.i + 1;
Brandon Jonese87ea2b2021-04-14 17:05:07 +000072 return vec4<f32>(f32(data.i) / 255.0, 0.0, 0.0, 1.0);
Austin Engb3ab21e2020-12-23 19:42:00 +000073 })");
Yunchao He733842c2019-11-23 00:20:53 +000074
Brandon Jones41c87d92021-05-21 05:01:38 +000075 utils::ComboRenderPipelineDescriptor rpDesc;
Brandon Jonesbff9d3a2021-03-18 02:54:27 +000076 rpDesc.vertex.module = vsModule;
77 rpDesc.cFragment.module = fsModule;
78 rpDesc.primitive.topology = wgpu::PrimitiveTopology::PointList;
79 rpDesc.cTargets[0].format = colorFormat;
Yunchao He733842c2019-11-23 00:20:53 +000080
Brandon Jones41c87d92021-05-21 05:01:38 +000081 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
Yunchao He733842c2019-11-23 00:20:53 +000082
83 wgpu::BindGroup bindGroup =
Yunchao He4326a8a2019-12-09 19:17:22 +000084 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
Yunchao He733842c2019-11-23 00:20:53 +000085 return std::make_tuple(pipeline, bindGroup);
86 }
Yunchao He02dd7332019-11-20 00:05:20 +000087};
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.
94TEST_P(GpuMemorySyncTests, ComputePass) {
95 // Create pipeline, bind group, and buffer for compute pass.
Yunchao He02dd7332019-11-20 00:05:20 +000096 wgpu::Buffer buffer = CreateBuffer();
Yunchao He733842c2019-11-23 00:20:53 +000097 wgpu::ComputePipeline compute;
98 wgpu::BindGroup bindGroup;
99 std::tie(compute, bindGroup) = CreatePipelineAndBindGroupForCompute(buffer);
Yunchao He02dd7332019-11-20 00:05:20 +0000100 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 Wallez3da19b82020-03-31 16:23:35 +0000108 pass.Dispatch(1);
Yunchao He02dd7332019-11-20 00:05:20 +0000109 pass.EndPass();
110 }
111
Yunchao He02dd7332019-11-20 00:05:20 +0000112 wgpu::CommandBuffer commands = encoder.Finish();
113 queue.Submit(1, &commands);
114
Yunchao He733842c2019-11-23 00:20:53 +0000115 // Verify the result.
Yunchao He02dd7332019-11-20 00:05:20 +0000116 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.
125TEST_P(GpuMemorySyncTests, RenderPass) {
126 // Create pipeline, bind group, and buffer for render pass.
Yunchao He02dd7332019-11-20 00:05:20 +0000127 wgpu::Buffer buffer = CreateBuffer();
Yunchao He733842c2019-11-23 00:20:53 +0000128 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 He02dd7332019-11-20 00:05:20 +0000133 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 Wallez67b1ad72020-03-31 16:21:35 +0000141 pass.Draw(1);
Yunchao He02dd7332019-11-20 00:05:20 +0000142 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 He733842c2019-11-23 00:20:53 +0000152// 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.
154TEST_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 Wallez67b1ad72020-03-31 16:21:35 +0000173 pass0.Draw(1);
Yunchao He733842c2019-11-23 00:20:53 +0000174 pass0.EndPass();
175
Yunchao Hed28b5782019-12-19 18:50:18 +0000176 // Read that data in compute pass.
Yunchao He733842c2019-11-23 00:20:53 +0000177 wgpu::ComputePassEncoder pass1 = encoder.BeginComputePass();
178 pass1.SetPipeline(compute);
179 pass1.SetBindGroup(0, bindGroup1);
Corentin Wallez3da19b82020-03-31 16:23:35 +0000180 pass1.Dispatch(1);
Yunchao He733842c2019-11-23 00:20:53 +0000181 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.
192TEST_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 Hed28b5782019-12-19 18:50:18 +0000207 // Write data into a storage buffer in compute pass.
Yunchao He733842c2019-11-23 00:20:53 +0000208 wgpu::ComputePassEncoder pass0 = encoder.BeginComputePass();
209 pass0.SetPipeline(compute);
210 pass0.SetBindGroup(0, bindGroup1);
Corentin Wallez3da19b82020-03-31 16:23:35 +0000211 pass0.Dispatch(1);
Yunchao He733842c2019-11-23 00:20:53 +0000212 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 Wallez67b1ad72020-03-31 16:21:35 +0000218 pass1.Draw(1);
Yunchao He733842c2019-11-23 00:20:53 +0000219 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 Ninomiya2afea0c2020-07-10 20:33:08 +0000228DAWN_INSTANTIATE_TEST(GpuMemorySyncTests,
229 D3D12Backend(),
230 MetalBackend(),
231 OpenGLBackend(),
Stephen Whitef31b78e2020-12-04 15:59:29 +0000232 OpenGLESBackend(),
Kai Ninomiya2afea0c2020-07-10 20:33:08 +0000233 VulkanBackend());
Yunchao He40b10e42019-11-28 18:55:45 +0000234
Yunchao He612a63a2019-11-18 04:28:24 +0000235class 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 Wallez7aec4ae2021-03-24 15:55:32 +0000245 wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +0000246 [[block]] struct Data {
Ben Claytonc5686842021-03-17 09:48:19 +0000247 a : f32;
Austin Engb3ab21e2020-12-23 19:42:00 +0000248 };
Ben Clayton15eba9a2021-06-08 15:36:44 +0000249 [[group(0), binding(0)]] var<storage, read_write> data : Data;
Sarah2a57db72021-06-23 19:19:06 +0000250 [[stage(compute), workgroup_size(1)]] fn main() {
Austin Engb3ab21e2020-12-23 19:42:00 +0000251 data.a = 1.0;
252 })");
Yunchao He612a63a2019-11-18 04:28:24 +0000253
Yunchao He612a63a2019-11-18 04:28:24 +0000254 wgpu::ComputePipelineDescriptor cpDesc;
Brandon Jones0d50a2c2021-06-09 18:07:32 +0000255 cpDesc.compute.module = csModule;
256 cpDesc.compute.entryPoint = "main";
Yunchao He612a63a2019-11-18 04:28:24 +0000257 wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cpDesc);
258
259 wgpu::BindGroup bindGroup =
Yunchao He4326a8a2019-12-09 19:17:22 +0000260 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, mBuffer}});
Yunchao He612a63a2019-11-18 04:28:24 +0000261 return std::make_tuple(pipeline, bindGroup);
262 }
263
264 std::tuple<wgpu::RenderPipeline, wgpu::BindGroup> CreatePipelineAndBindGroupForRender(
265 wgpu::TextureFormat colorFormat) {
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000266 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
Brandon Jonese87ea2b2021-04-14 17:05:07 +0000267 [[stage(vertex)]] fn main() -> [[builtin(position)]] vec4<f32> {
268 return vec4<f32>(0.0, 0.0, 0.0, 1.0);
Austin Engb3ab21e2020-12-23 19:42:00 +0000269 })");
Yunchao He612a63a2019-11-18 04:28:24 +0000270
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000271 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +0000272 [[block]] struct Contents {
Ben Claytonc5686842021-03-17 09:48:19 +0000273 color : f32;
Austin Engb3ab21e2020-12-23 19:42:00 +0000274 };
dan sinclair0f9c2d72021-01-19 14:18:51 +0000275 [[group(0), binding(0)]] var<uniform> contents : Contents;
Austin Engb3ab21e2020-12-23 19:42:00 +0000276
Brandon Jonese87ea2b2021-04-14 17:05:07 +0000277 [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
278 return vec4<f32>(contents.color, 0.0, 0.0, 1.0);
Austin Engb3ab21e2020-12-23 19:42:00 +0000279 })");
Yunchao He612a63a2019-11-18 04:28:24 +0000280
Brandon Jones41c87d92021-05-21 05:01:38 +0000281 utils::ComboRenderPipelineDescriptor rpDesc;
Brandon Jonesbff9d3a2021-03-18 02:54:27 +0000282 rpDesc.vertex.module = vsModule;
283 rpDesc.cFragment.module = fsModule;
284 rpDesc.primitive.topology = wgpu::PrimitiveTopology::PointList;
285 rpDesc.cTargets[0].format = colorFormat;
Yunchao He612a63a2019-11-18 04:28:24 +0000286
Brandon Jones41c87d92021-05-21 05:01:38 +0000287 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
Yunchao He612a63a2019-11-18 04:28:24 +0000288
289 wgpu::BindGroup bindGroup =
Yunchao He4326a8a2019-12-09 19:17:22 +0000290 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, mBuffer}});
Yunchao He612a63a2019-11-18 04:28:24 +0000291 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.
299TEST_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 Wallez3da19b82020-03-31 16:23:35 +0000315 pass0.Dispatch(1);
Yunchao He612a63a2019-11-18 04:28:24 +0000316 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 Wallez67b1ad72020-03-31 16:21:35 +0000322 pass1.Draw(1);
Yunchao He612a63a2019-11-18 04:28:24 +0000323 pass1.EndPass();
324
325 wgpu::CommandBuffer commands = encoder0.Finish();
326 queue.Submit(1, &commands);
327
328 // Verify the rendering result.
Yunchao He0c02f542019-11-19 17:57:30 +0000329 EXPECT_PIXEL_RGBA8_EQ(RGBA8::kRed, renderPass.color, 0, 0);
Yunchao He612a63a2019-11-18 04:28:24 +0000330}
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.
335TEST_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 Wallez3da19b82020-03-31 16:23:35 +0000352 pass0.Dispatch(1);
Yunchao He612a63a2019-11-18 04:28:24 +0000353 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 Wallez67b1ad72020-03-31 16:21:35 +0000361 pass1.Draw(1);
Yunchao He612a63a2019-11-18 04:28:24 +0000362 pass1.EndPass();
363
364 cb[1] = encoder1.Finish();
365 queue.Submit(2, cb);
366
367 // Verify the rendering result.
Yunchao He0c02f542019-11-19 17:57:30 +0000368 EXPECT_PIXEL_RGBA8_EQ(RGBA8::kRed, renderPass.color, 0, 0);
Yunchao He612a63a2019-11-18 04:28:24 +0000369}
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.
374TEST_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 Wallez3da19b82020-03-31 16:23:35 +0000391 pass0.Dispatch(1);
Yunchao He612a63a2019-11-18 04:28:24 +0000392 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 Wallez67b1ad72020-03-31 16:21:35 +0000401 pass1.Draw(1);
Yunchao He612a63a2019-11-18 04:28:24 +0000402 pass1.EndPass();
403
404 cb[1] = encoder1.Finish();
405 queue.Submit(1, &cb[1]);
406
407 // Verify the rendering result.
Yunchao He0c02f542019-11-19 17:57:30 +0000408 EXPECT_PIXEL_RGBA8_EQ(RGBA8::kRed, renderPass.color, 0, 0);
Yunchao He612a63a2019-11-18 04:28:24 +0000409}
410
Yunchao He612a63a2019-11-18 04:28:24 +0000411DAWN_INSTANTIATE_TEST(StorageToUniformSyncTests,
Austin Eng6c1d6462020-02-25 16:23:17 +0000412 D3D12Backend(),
413 MetalBackend(),
414 OpenGLBackend(),
Stephen Whitef31b78e2020-12-04 15:59:29 +0000415 OpenGLESBackend(),
Austin Eng6c1d6462020-02-25 16:23:17 +0000416 VulkanBackend());
Yunchao He40b10e42019-11-28 18:55:45 +0000417
418constexpr int kRTSize = 8;
419constexpr int kVertexBufferStride = 4 * sizeof(float);
420
421class 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 Wallez47a33412020-06-02 09:24:39 +0000430 queue.WriteBuffer(buffer, 0, zeros.data(), size);
Yunchao He40b10e42019-11-28 18:55:45 +0000431
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 Hece8bf122019-12-05 21:18:12 +0000437// 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 He40b10e42019-11-28 18:55:45 +0000440TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) {
441 // Create pipeline, bind group, and different buffers for compute pass.
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000442 wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +0000443 [[block]] struct VBContents {
Ben Claytonc5686842021-03-17 09:48:19 +0000444 pos : array<vec4<f32>, 4>;
Yunchao He40b10e42019-11-28 18:55:45 +0000445 };
Ben Clayton15eba9a2021-06-08 15:36:44 +0000446 [[group(0), binding(0)]] var<storage, read_write> vbContents : VBContents;
Yunchao He40b10e42019-11-28 18:55:45 +0000447
Austin Engb3ab21e2020-12-23 19:42:00 +0000448 [[block]] struct IBContents {
Ben Claytonc5686842021-03-17 09:48:19 +0000449 indices : array<vec4<i32>, 2>;
Yunchao He40b10e42019-11-28 18:55:45 +0000450 };
Ben Clayton15eba9a2021-06-08 15:36:44 +0000451 [[group(0), binding(1)]] var<storage, read_write> ibContents : IBContents;
Yunchao He40b10e42019-11-28 18:55:45 +0000452
Ben Clayton4773e8d2021-07-13 15:21:07 +0000453 [[block]] struct ColorContents {
Ben Claytonc5686842021-03-17 09:48:19 +0000454 color : f32;
Yunchao He40b10e42019-11-28 18:55:45 +0000455 };
Ben Clayton4773e8d2021-07-13 15:21:07 +0000456 [[group(0), binding(2)]] var<storage, read_write> uniformContents : ColorContents;
457 [[group(0), binding(3)]] var<storage, read_write> storageContents : ColorContents;
Yunchao He40b10e42019-11-28 18:55:45 +0000458
Sarah2a57db72021-06-23 19:19:06 +0000459 [[stage(compute), workgroup_size(1)]] fn main() {
Austin Engb3ab21e2020-12-23 19:42:00 +0000460 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 Jonese87ea2b2021-04-14 17:05:07 +0000464 let dummy : i32 = 0;
Austin Engb3ab21e2020-12-23 19:42:00 +0000465 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 He40b10e42019-11-28 18:55:45 +0000469 })");
470
Yunchao He40b10e42019-11-28 18:55:45 +0000471 wgpu::ComputePipelineDescriptor cpDesc;
Brandon Jones0d50a2c2021-06-09 18:07:32 +0000472 cpDesc.compute.module = csModule;
473 cpDesc.compute.entryPoint = "main";
Yunchao He40b10e42019-11-28 18:55:45 +0000474 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 Hece8bf122019-12-05 21:18:12 +0000481 wgpu::Buffer uniformBuffer =
Yunchao He40b10e42019-11-28 18:55:45 +0000482 CreateZeroedBuffer(sizeof(float), wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Storage |
483 wgpu::BufferUsage::CopyDst);
Austin Engb3ab21e2020-12-23 19:42:00 +0000484 wgpu::Buffer storageBuffer =
Yunchao Hece8bf122019-12-05 21:18:12 +0000485 CreateZeroedBuffer(sizeof(float), wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst);
Yunchao He40b10e42019-11-28 18:55:45 +0000486
Yunchao He4326a8a2019-12-09 19:17:22 +0000487 wgpu::BindGroup bindGroup0 = utils::MakeBindGroup(
488 device, cp.GetBindGroupLayout(0),
Austin Engb3ab21e2020-12-23 19:42:00 +0000489 {{0, vertexBuffer}, {1, indexBuffer}, {2, uniformBuffer}, {3, storageBuffer}});
Yunchao He40b10e42019-11-28 18:55:45 +0000490 // 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 Wallez3da19b82020-03-31 16:23:35 +0000495 pass0.Dispatch(1);
Yunchao He40b10e42019-11-28 18:55:45 +0000496 pass0.EndPass();
497
498 // Create pipeline, bind group, and reuse buffers in render pass.
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000499 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
Brandon Jonese87ea2b2021-04-14 17:05:07 +0000500 [[stage(vertex)]]
501 fn main([[location(0)]] pos : vec4<f32>) -> [[builtin(position)]] vec4<f32> {
502 return pos;
Yunchao He40b10e42019-11-28 18:55:45 +0000503 })");
504
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000505 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +0000506 [[block]] struct Buf {
Ben Claytonc5686842021-03-17 09:48:19 +0000507 color : f32;
Yunchao He40b10e42019-11-28 18:55:45 +0000508 };
509
dan sinclair0f9c2d72021-01-19 14:18:51 +0000510 [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf;
Ben Clayton15eba9a2021-06-08 15:36:44 +0000511 [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf;
Yunchao He40b10e42019-11-28 18:55:45 +0000512
Brandon Jonese87ea2b2021-04-14 17:05:07 +0000513 [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
514 return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0);
Yunchao He40b10e42019-11-28 18:55:45 +0000515 })");
516
Yunchao He40b10e42019-11-28 18:55:45 +0000517 utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize);
518
Brandon Jones41c87d92021-05-21 05:01:38 +0000519 utils::ComboRenderPipelineDescriptor rpDesc;
Brandon Jonesbff9d3a2021-03-18 02:54:27 +0000520 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 He40b10e42019-11-28 18:55:45 +0000528
Brandon Jones41c87d92021-05-21 05:01:38 +0000529 wgpu::RenderPipeline rp = device.CreateRenderPipeline(&rpDesc);
Yunchao He40b10e42019-11-28 18:55:45 +0000530
Austin Engb3ab21e2020-12-23 19:42:00 +0000531 wgpu::BindGroup bindGroup1 = utils::MakeBindGroup(device, rp.GetBindGroupLayout(0),
532 {{0, uniformBuffer}, {1, storageBuffer}});
Yunchao He40b10e42019-11-28 18:55:45 +0000533
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 Wallez5fad85b2020-11-25 08:54:14 +0000538 pass1.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0);
Yunchao He40b10e42019-11-28 18:55:45 +0000539 pass1.SetBindGroup(0, bindGroup1);
Corentin Wallez67b1ad72020-03-31 16:21:35 +0000540 pass1.DrawIndexed(6);
Yunchao He40b10e42019-11-28 18:55:45 +0000541 pass1.EndPass();
542
543 wgpu::CommandBuffer commandBuffer = encoder.Finish();
544 queue.Submit(1, &commandBuffer);
545
546 // Verify the rendering result.
Yunchao He4eb40c12021-03-31 22:15:53 +0000547 uint32_t min = 1, max = kRTSize - 3;
Yunchao He40b10e42019-11-28 18:55:45 +0000548 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 Hece8bf122019-12-05 21:18:12 +0000555// 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 He40b10e42019-11-28 18:55:45 +0000557TEST_P(MultipleWriteThenMultipleReadTests, OneBuffer) {
Stephen White032500b2021-01-26 15:01:18 +0000558 // 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 Shao44fc6e32021-05-26 01:04:32 +0000561 DAWN_SUPPRESS_TEST_IF(IsOpenGLES());
Stephen White032500b2021-01-26 15:01:18 +0000562
Yunchao He40b10e42019-11-28 18:55:45 +0000563 // Create pipeline, bind group, and a complex buffer for compute pass.
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000564 wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +0000565 [[block]] struct Contents {
Ben Claytonc5686842021-03-17 09:48:19 +0000566 [[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 He40b10e42019-11-28 18:55:45 +0000570 };
571
Ben Clayton15eba9a2021-06-08 15:36:44 +0000572 [[group(0), binding(0)]] var<storage, read_write> contents : Contents;
Austin Engb3ab21e2020-12-23 19:42:00 +0000573
Sarah2a57db72021-06-23 19:19:06 +0000574 [[stage(compute), workgroup_size(1)]] fn main() {
Austin Engb3ab21e2020-12-23 19:42:00 +0000575 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 Jonese87ea2b2021-04-14 17:05:07 +0000579 let dummy : i32 = 0;
Austin Engb3ab21e2020-12-23 19:42:00 +0000580 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 He40b10e42019-11-28 18:55:45 +0000584 })");
585
Yunchao He40b10e42019-11-28 18:55:45 +0000586 wgpu::ComputePipelineDescriptor cpDesc;
Brandon Jones0d50a2c2021-06-09 18:07:32 +0000587 cpDesc.compute.module = csModule;
588 cpDesc.compute.entryPoint = "main";
Yunchao He40b10e42019-11-28 18:55:45 +0000589 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 Engb3ab21e2020-12-23 19:42:00 +0000595 float color0;
596 char padding2[256 - sizeof(float)];
597 float color1;
James Pricea0b31e02021-06-28 08:38:28 +0000598 char padding3[256 - sizeof(float)];
Yunchao He40b10e42019-11-28 18:55:45 +0000599 };
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 He4326a8a2019-12-09 19:17:22 +0000604 wgpu::BindGroup bindGroup0 =
605 utils::MakeBindGroup(device, cp.GetBindGroupLayout(0), {{0, buffer}});
Yunchao He40b10e42019-11-28 18:55:45 +0000606
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 Wallez3da19b82020-03-31 16:23:35 +0000612 pass0.Dispatch(1);
Yunchao He40b10e42019-11-28 18:55:45 +0000613 pass0.EndPass();
614
615 // Create pipeline, bind group, and reuse the buffer in render pass.
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000616 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
Brandon Jonese87ea2b2021-04-14 17:05:07 +0000617 [[stage(vertex)]]
618 fn main([[location(0)]] pos : vec4<f32>) -> [[builtin(position)]] vec4<f32> {
619 return pos;
Yunchao He40b10e42019-11-28 18:55:45 +0000620 })");
621
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000622 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
Austin Engb3ab21e2020-12-23 19:42:00 +0000623 [[block]] struct Buf {
Ben Claytonc5686842021-03-17 09:48:19 +0000624 color : f32;
Yunchao He40b10e42019-11-28 18:55:45 +0000625 };
dan sinclair0f9c2d72021-01-19 14:18:51 +0000626 [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf;
Ben Clayton15eba9a2021-06-08 15:36:44 +0000627 [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf;
Yunchao He40b10e42019-11-28 18:55:45 +0000628
Brandon Jonese87ea2b2021-04-14 17:05:07 +0000629 [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
630 return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0);
Yunchao He40b10e42019-11-28 18:55:45 +0000631 })");
632
Yunchao He40b10e42019-11-28 18:55:45 +0000633 utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize);
634
Brandon Jones41c87d92021-05-21 05:01:38 +0000635 utils::ComboRenderPipelineDescriptor rpDesc;
Brandon Jonesbff9d3a2021-03-18 02:54:27 +0000636 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 He40b10e42019-11-28 18:55:45 +0000644
Brandon Jones41c87d92021-05-21 05:01:38 +0000645 wgpu::RenderPipeline rp = device.CreateRenderPipeline(&rpDesc);
Yunchao He40b10e42019-11-28 18:55:45 +0000646
647 wgpu::BindGroup bindGroup1 =
Yunchao He4326a8a2019-12-09 19:17:22 +0000648 utils::MakeBindGroup(device, rp.GetBindGroupLayout(0),
Yunchao He40b10e42019-11-28 18:55:45 +0000649 {{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 Wallez5fad85b2020-11-25 08:54:14 +0000656 pass1.SetIndexBuffer(buffer, wgpu::IndexFormat::Uint32, offsetof(Data, indices));
Yunchao He40b10e42019-11-28 18:55:45 +0000657 pass1.SetBindGroup(0, bindGroup1);
Corentin Wallez67b1ad72020-03-31 16:21:35 +0000658 pass1.DrawIndexed(6);
Yunchao He40b10e42019-11-28 18:55:45 +0000659 pass1.EndPass();
660
661 wgpu::CommandBuffer commandBuffer = encoder.Finish();
662 queue.Submit(1, &commandBuffer);
663
664 // Verify the rendering result.
Yunchao He4eb40c12021-03-31 22:15:53 +0000665 uint32_t min = 1, max = kRTSize - 3;
Yunchao He40b10e42019-11-28 18:55:45 +0000666 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 Heef8dee92019-12-09 21:35:38 +0000672DAWN_INSTANTIATE_TEST(MultipleWriteThenMultipleReadTests,
Austin Eng6c1d6462020-02-25 16:23:17 +0000673 D3D12Backend(),
674 MetalBackend(),
675 OpenGLBackend(),
Stephen Whitef31b78e2020-12-04 15:59:29 +0000676 OpenGLESBackend(),
Austin Eng6c1d6462020-02-25 16:23:17 +0000677 VulkanBackend());