1 // Copyright 2021 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 "tests/DawnTest.h"
16
17 #include "utils/ComboRenderPipelineDescriptor.h"
18 #include "utils/WGPUHelpers.h"
19
20 #include <numeric>
21 #include <vector>
22
23 class ShaderTests : public DawnTest {
24 public:
CreateBuffer(const uint32_t count)25 wgpu::Buffer CreateBuffer(const uint32_t count) {
26 std::vector<uint32_t> data(count, 0);
27 uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(uint32_t));
28 return utils::CreateBufferFromData(device, data.data(), bufferSize,
29 wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc);
30 }
CreateComputePipeline(const std::string & shader,const char * entryPoint,const std::vector<wgpu::ConstantEntry> * constants=nullptr)31 wgpu::ComputePipeline CreateComputePipeline(
32 const std::string& shader,
33 const char* entryPoint,
34 const std::vector<wgpu::ConstantEntry>* constants = nullptr) {
35 wgpu::ComputePipelineDescriptor csDesc;
36 csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
37 csDesc.compute.entryPoint = entryPoint;
38 if (constants) {
39 csDesc.compute.constants = constants->data();
40 csDesc.compute.constantCount = constants->size();
41 }
42 return device.CreateComputePipeline(&csDesc);
43 }
44 };
45
46 // Test that log2 is being properly calculated, base on crbug.com/1046622
TEST_P(ShaderTests,ComputeLog2)47 TEST_P(ShaderTests, ComputeLog2) {
48 uint32_t const kSteps = 19;
49 std::vector<uint32_t> expected{0, 1, 1, 2, 2, 3, 3, 4, 4, 5, 5, 6, 6, 7, 7, 8, 8, 9, 32};
50 wgpu::Buffer buffer = CreateBuffer(kSteps);
51
52 std::string shader = R"(
53 [[block]] struct Buf {
54 data : array<u32, 19>;
55 };
56
57 [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
58
59 [[stage(compute), workgroup_size(1)]] fn main() {
60 let factor : f32 = 1.0001;
61
62 buf.data[0] = u32(log2(1.0 * factor));
63 buf.data[1] = u32(log2(2.0 * factor));
64 buf.data[2] = u32(log2(3.0 * factor));
65 buf.data[3] = u32(log2(4.0 * factor));
66 buf.data[4] = u32(log2(7.0 * factor));
67 buf.data[5] = u32(log2(8.0 * factor));
68 buf.data[6] = u32(log2(15.0 * factor));
69 buf.data[7] = u32(log2(16.0 * factor));
70 buf.data[8] = u32(log2(31.0 * factor));
71 buf.data[9] = u32(log2(32.0 * factor));
72 buf.data[10] = u32(log2(63.0 * factor));
73 buf.data[11] = u32(log2(64.0 * factor));
74 buf.data[12] = u32(log2(127.0 * factor));
75 buf.data[13] = u32(log2(128.0 * factor));
76 buf.data[14] = u32(log2(255.0 * factor));
77 buf.data[15] = u32(log2(256.0 * factor));
78 buf.data[16] = u32(log2(511.0 * factor));
79 buf.data[17] = u32(log2(512.0 * factor));
80 buf.data[18] = u32(log2(4294967295.0 * factor));
81 })";
82
83 wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main");
84
85 wgpu::BindGroup bindGroup =
86 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
87
88 wgpu::CommandBuffer commands;
89 {
90 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
91 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
92 pass.SetPipeline(pipeline);
93 pass.SetBindGroup(0, bindGroup);
94 pass.Dispatch(1);
95 pass.EndPass();
96
97 commands = encoder.Finish();
98 }
99
100 queue.Submit(1, &commands);
101
102 EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kSteps);
103 }
104
TEST_P(ShaderTests,BadWGSL)105 TEST_P(ShaderTests, BadWGSL) {
106 DAWN_TEST_UNSUPPORTED_IF(HasToggleEnabled("skip_validation"));
107
108 std::string shader = R"(
109 I am an invalid shader and should never pass validation!
110 })";
111 ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, shader.c_str()));
112 }
113
114 // Tests that shaders using non-struct function parameters and return values for shader stage I/O
115 // can compile and link successfully.
TEST_P(ShaderTests,WGSLParamIO)116 TEST_P(ShaderTests, WGSLParamIO) {
117 std::string vertexShader = R"(
118 [[stage(vertex)]]
119 fn main([[builtin(vertex_index)]] VertexIndex : u32) -> [[builtin(position)]] vec4<f32> {
120 var pos = array<vec2<f32>, 3>(
121 vec2<f32>(-1.0, 1.0),
122 vec2<f32>( 1.0, 1.0),
123 vec2<f32>( 0.0, -1.0));
124 return vec4<f32>(pos[VertexIndex], 0.0, 1.0);
125 })";
126 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, vertexShader.c_str());
127
128 std::string fragmentShader = R"(
129 [[stage(fragment)]]
130 fn main([[builtin(position)]] fragCoord : vec4<f32>) -> [[location(0)]] vec4<f32> {
131 return vec4<f32>(fragCoord.xy, 0.0, 1.0);
132 })";
133 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, fragmentShader.c_str());
134
135 utils::ComboRenderPipelineDescriptor rpDesc;
136 rpDesc.vertex.module = vsModule;
137 rpDesc.cFragment.module = fsModule;
138 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
139 }
140
141 // Tests that a vertex shader using struct function parameters and return values for shader stage
142 // I/O can compile and link successfully against a fragement shader using compatible non-struct I/O.
TEST_P(ShaderTests,WGSLMixedStructParamIO)143 TEST_P(ShaderTests, WGSLMixedStructParamIO) {
144 std::string vertexShader = R"(
145 struct VertexIn {
146 [[location(0)]] position : vec3<f32>;
147 [[location(1)]] color : vec4<f32>;
148 };
149
150 struct VertexOut {
151 [[location(0)]] color : vec4<f32>;
152 [[builtin(position)]] position : vec4<f32>;
153 };
154
155 [[stage(vertex)]]
156 fn main(input : VertexIn) -> VertexOut {
157 var output : VertexOut;
158 output.position = vec4<f32>(input.position, 1.0);
159 output.color = input.color;
160 return output;
161 })";
162 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, vertexShader.c_str());
163
164 std::string fragmentShader = R"(
165 [[stage(fragment)]]
166 fn main([[location(0)]] color : vec4<f32>) -> [[location(0)]] vec4<f32> {
167 return color;
168 })";
169 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, fragmentShader.c_str());
170
171 utils::ComboRenderPipelineDescriptor rpDesc;
172 rpDesc.vertex.module = vsModule;
173 rpDesc.cFragment.module = fsModule;
174 rpDesc.vertex.bufferCount = 1;
175 rpDesc.cBuffers[0].attributeCount = 2;
176 rpDesc.cBuffers[0].arrayStride = 28;
177 rpDesc.cAttributes[0].shaderLocation = 0;
178 rpDesc.cAttributes[0].format = wgpu::VertexFormat::Float32x3;
179 rpDesc.cAttributes[1].shaderLocation = 1;
180 rpDesc.cAttributes[1].format = wgpu::VertexFormat::Float32x4;
181 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
182 }
183
184 // Tests that shaders using struct function parameters and return values for shader stage I/O
185 // can compile and link successfully.
TEST_P(ShaderTests,WGSLStructIO)186 TEST_P(ShaderTests, WGSLStructIO) {
187 std::string vertexShader = R"(
188 struct VertexIn {
189 [[location(0)]] position : vec3<f32>;
190 [[location(1)]] color : vec4<f32>;
191 };
192
193 struct VertexOut {
194 [[location(0)]] color : vec4<f32>;
195 [[builtin(position)]] position : vec4<f32>;
196 };
197
198 [[stage(vertex)]]
199 fn main(input : VertexIn) -> VertexOut {
200 var output : VertexOut;
201 output.position = vec4<f32>(input.position, 1.0);
202 output.color = input.color;
203 return output;
204 })";
205 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, vertexShader.c_str());
206
207 std::string fragmentShader = R"(
208 struct FragmentIn {
209 [[location(0)]] color : vec4<f32>;
210 [[builtin(position)]] fragCoord : vec4<f32>;
211 };
212
213 [[stage(fragment)]]
214 fn main(input : FragmentIn) -> [[location(0)]] vec4<f32> {
215 return input.color * input.fragCoord;
216 })";
217 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, fragmentShader.c_str());
218
219 utils::ComboRenderPipelineDescriptor rpDesc;
220 rpDesc.vertex.module = vsModule;
221 rpDesc.cFragment.module = fsModule;
222 rpDesc.vertex.bufferCount = 1;
223 rpDesc.cBuffers[0].attributeCount = 2;
224 rpDesc.cBuffers[0].arrayStride = 28;
225 rpDesc.cAttributes[0].shaderLocation = 0;
226 rpDesc.cAttributes[0].format = wgpu::VertexFormat::Float32x3;
227 rpDesc.cAttributes[1].shaderLocation = 1;
228 rpDesc.cAttributes[1].format = wgpu::VertexFormat::Float32x4;
229 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
230 }
231
232 // Tests that shaders I/O structs that us compatible locations but are not sorted by hand can link.
TEST_P(ShaderTests,WGSLUnsortedStructIO)233 TEST_P(ShaderTests, WGSLUnsortedStructIO) {
234 std::string vertexShader = R"(
235 struct VertexIn {
236 [[location(0)]] position : vec3<f32>;
237 [[location(1)]] color : vec4<f32>;
238 };
239
240 struct VertexOut {
241 [[builtin(position)]] position : vec4<f32>;
242 [[location(0)]] color : vec4<f32>;
243 };
244
245 [[stage(vertex)]]
246 fn main(input : VertexIn) -> VertexOut {
247 var output : VertexOut;
248 output.position = vec4<f32>(input.position, 1.0);
249 output.color = input.color;
250 return output;
251 })";
252 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, vertexShader.c_str());
253
254 std::string fragmentShader = R"(
255 struct FragmentIn {
256 [[location(0)]] color : vec4<f32>;
257 [[builtin(position)]] fragCoord : vec4<f32>;
258 };
259
260 [[stage(fragment)]]
261 fn main(input : FragmentIn) -> [[location(0)]] vec4<f32> {
262 return input.color * input.fragCoord;
263 })";
264 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, fragmentShader.c_str());
265
266 utils::ComboRenderPipelineDescriptor rpDesc;
267 rpDesc.vertex.module = vsModule;
268 rpDesc.cFragment.module = fsModule;
269 rpDesc.vertex.bufferCount = 1;
270 rpDesc.cBuffers[0].attributeCount = 2;
271 rpDesc.cBuffers[0].arrayStride = 28;
272 rpDesc.cAttributes[0].shaderLocation = 0;
273 rpDesc.cAttributes[0].format = wgpu::VertexFormat::Float32x3;
274 rpDesc.cAttributes[1].shaderLocation = 1;
275 rpDesc.cAttributes[1].format = wgpu::VertexFormat::Float32x4;
276 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
277 }
278
279 // Tests that shaders I/O structs can be shared between vertex and fragment shaders.
TEST_P(ShaderTests,WGSLSharedStructIO)280 TEST_P(ShaderTests, WGSLSharedStructIO) {
281 std::string shader = R"(
282 struct VertexIn {
283 [[location(0)]] position : vec3<f32>;
284 [[location(1)]] color : vec4<f32>;
285 };
286
287 struct VertexOut {
288 [[location(0)]] color : vec4<f32>;
289 [[builtin(position)]] position : vec4<f32>;
290 };
291
292 [[stage(vertex)]]
293 fn vertexMain(input : VertexIn) -> VertexOut {
294 var output : VertexOut;
295 output.position = vec4<f32>(input.position, 1.0);
296 output.color = input.color;
297 return output;
298 }
299
300 [[stage(fragment)]]
301 fn fragmentMain(input : VertexOut) -> [[location(0)]] vec4<f32> {
302 return input.color;
303 })";
304 wgpu::ShaderModule shaderModule = utils::CreateShaderModule(device, shader.c_str());
305
306 utils::ComboRenderPipelineDescriptor rpDesc;
307 rpDesc.vertex.module = shaderModule;
308 rpDesc.vertex.entryPoint = "vertexMain";
309 rpDesc.cFragment.module = shaderModule;
310 rpDesc.cFragment.entryPoint = "fragmentMain";
311 rpDesc.vertex.bufferCount = 1;
312 rpDesc.cBuffers[0].attributeCount = 2;
313 rpDesc.cBuffers[0].arrayStride = 28;
314 rpDesc.cAttributes[0].shaderLocation = 0;
315 rpDesc.cAttributes[0].format = wgpu::VertexFormat::Float32x3;
316 rpDesc.cAttributes[1].shaderLocation = 1;
317 rpDesc.cAttributes[1].format = wgpu::VertexFormat::Float32x4;
318 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc);
319 }
320
321 // This is a regression test for an issue caused by the FirstIndexOffset transfrom being done before
322 // the BindingRemapper, causing an intermediate AST to be invalid (and fail the overall
323 // compilation).
TEST_P(ShaderTests,FirstIndexOffsetRegisterConflictInHLSLTransforms)324 TEST_P(ShaderTests, FirstIndexOffsetRegisterConflictInHLSLTransforms) {
325 // TODO(crbug.com/dawn/658): Crashes on bots because there are two entrypoints in the shader.
326 DAWN_SUPPRESS_TEST_IF(IsOpenGL() || IsOpenGLES());
327
328 const char* shader = R"(
329 // Dumped WGSL:
330
331 struct Inputs {
332 [[location(1)]] attrib1 : u32;
333 // The extra register added to handle base_vertex for vertex_index conflicts with [1]
334 [[builtin(vertex_index)]] vertexIndex: u32;
335 };
336
337 // [1] a binding point that conflicts with the regitster
338 [[block]] struct S1 { data : array<vec4<u32>, 20>; };
339 [[group(0), binding(1)]] var<uniform> providedData1 : S1;
340
341 [[stage(vertex)]] fn vsMain(input : Inputs) -> [[builtin(position)]] vec4<f32> {
342 _ = providedData1.data[input.vertexIndex][0];
343 return vec4<f32>();
344 }
345
346 [[stage(fragment)]] fn fsMain() -> [[location(0)]] vec4<f32> {
347 return vec4<f32>();
348 }
349 )";
350 auto module = utils::CreateShaderModule(device, shader);
351
352 utils::ComboRenderPipelineDescriptor rpDesc;
353 rpDesc.vertex.module = module;
354 rpDesc.vertex.entryPoint = "vsMain";
355 rpDesc.cFragment.module = module;
356 rpDesc.cFragment.entryPoint = "fsMain";
357 rpDesc.vertex.bufferCount = 1;
358 rpDesc.cBuffers[0].attributeCount = 1;
359 rpDesc.cBuffers[0].arrayStride = 16;
360 rpDesc.cAttributes[0].shaderLocation = 1;
361 rpDesc.cAttributes[0].format = wgpu::VertexFormat::Uint8x2;
362 device.CreateRenderPipeline(&rpDesc);
363 }
364
365 // Test that WGSL built-in variable [[sample_index]] can be used in fragment shaders.
TEST_P(ShaderTests,SampleIndex)366 TEST_P(ShaderTests, SampleIndex) {
367 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
368 [[stage(vertex)]]
369 fn main([[location(0)]] pos : vec4<f32>) -> [[builtin(position)]] vec4<f32> {
370 return pos;
371 })");
372
373 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
374 [[stage(fragment)]] fn main([[builtin(sample_index)]] sampleIndex : u32)
375 -> [[location(0)]] vec4<f32> {
376 return vec4<f32>(f32(sampleIndex), 1.0, 0.0, 1.0);
377 })");
378
379 utils::ComboRenderPipelineDescriptor descriptor;
380 descriptor.vertex.module = vsModule;
381 descriptor.cFragment.module = fsModule;
382 descriptor.primitive.topology = wgpu::PrimitiveTopology::TriangleList;
383 descriptor.vertex.bufferCount = 1;
384 descriptor.cBuffers[0].arrayStride = 4 * sizeof(float);
385 descriptor.cBuffers[0].attributeCount = 1;
386 descriptor.cAttributes[0].format = wgpu::VertexFormat::Float32x4;
387 descriptor.cTargets[0].format = wgpu::TextureFormat::RGBA8Unorm;
388
389 device.CreateRenderPipeline(&descriptor);
390 }
391
392 // Test overridable constants without numeric identifiers
TEST_P(ShaderTests,OverridableConstants)393 TEST_P(ShaderTests, OverridableConstants) {
394 DAWN_TEST_UNSUPPORTED_IF(IsOpenGL());
395 DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
396
397 uint32_t const kCount = 11;
398 std::vector<uint32_t> expected(kCount);
399 std::iota(expected.begin(), expected.end(), 0);
400 wgpu::Buffer buffer = CreateBuffer(kCount);
401
402 std::string shader = R"(
403 [[override]] let c0: bool; // type: bool
404 [[override]] let c1: bool = false; // default override
405 [[override]] let c2: f32; // type: float32
406 [[override]] let c3: f32 = 0.0; // default override
407 [[override]] let c4: f32 = 4.0; // default
408 [[override]] let c5: i32; // type: int32
409 [[override]] let c6: i32 = 0; // default override
410 [[override]] let c7: i32 = 7; // default
411 [[override]] let c8: u32; // type: uint32
412 [[override]] let c9: u32 = 0u; // default override
413 [[override]] let c10: u32 = 10u; // default
414
415 [[block]] struct Buf {
416 data : array<u32, 11>;
417 };
418
419 [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
420
421 [[stage(compute), workgroup_size(1)]] fn main() {
422 buf.data[0] = u32(c0);
423 buf.data[1] = u32(c1);
424 buf.data[2] = u32(c2);
425 buf.data[3] = u32(c3);
426 buf.data[4] = u32(c4);
427 buf.data[5] = u32(c5);
428 buf.data[6] = u32(c6);
429 buf.data[7] = u32(c7);
430 buf.data[8] = u32(c8);
431 buf.data[9] = u32(c9);
432 buf.data[10] = u32(c10);
433 })";
434
435 std::vector<wgpu::ConstantEntry> constants;
436 constants.push_back({nullptr, "c0", 0});
437 constants.push_back({nullptr, "c1", 1});
438 constants.push_back({nullptr, "c2", 2});
439 constants.push_back({nullptr, "c3", 3});
440 // c4 is not assigned, testing default value
441 constants.push_back({nullptr, "c5", 5});
442 constants.push_back({nullptr, "c6", 6});
443 // c7 is not assigned, testing default value
444 constants.push_back({nullptr, "c8", 8});
445 constants.push_back({nullptr, "c9", 9});
446 // c10 is not assigned, testing default value
447
448 wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants);
449
450 wgpu::BindGroup bindGroup =
451 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
452
453 wgpu::CommandBuffer commands;
454 {
455 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
456 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
457 pass.SetPipeline(pipeline);
458 pass.SetBindGroup(0, bindGroup);
459 pass.Dispatch(1);
460 pass.EndPass();
461
462 commands = encoder.Finish();
463 }
464
465 queue.Submit(1, &commands);
466
467 EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kCount);
468 }
469
470 // Test overridable constants with numeric identifiers
TEST_P(ShaderTests,OverridableConstantsNumericIdentifiers)471 TEST_P(ShaderTests, OverridableConstantsNumericIdentifiers) {
472 DAWN_TEST_UNSUPPORTED_IF(IsOpenGL());
473 DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
474
475 uint32_t const kCount = 4;
476 std::vector<uint32_t> expected{1u, 2u, 3u, 0u};
477 wgpu::Buffer buffer = CreateBuffer(kCount);
478
479 std::string shader = R"(
480 [[override(1001)]] let c1: u32; // some big numeric id
481 [[override(1)]] let c2: u32 = 0u; // id == 1 might collide with some generated constant id
482 [[override(1003)]] let c3: u32 = 3u; // default
483 [[override(1004)]] let c4: u32; // default unspecified
484
485 [[block]] struct Buf {
486 data : array<u32, 4>;
487 };
488
489 [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
490
491 [[stage(compute), workgroup_size(1)]] fn main() {
492 buf.data[0] = c1;
493 buf.data[1] = c2;
494 buf.data[2] = c3;
495 buf.data[3] = c4;
496 })";
497
498 std::vector<wgpu::ConstantEntry> constants;
499 constants.push_back({nullptr, "1001", 1});
500 constants.push_back({nullptr, "1", 2});
501 // c3 is not assigned, testing default value
502 constants.push_back({nullptr, "1004", 0});
503
504 wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants);
505
506 wgpu::BindGroup bindGroup =
507 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
508
509 wgpu::CommandBuffer commands;
510 {
511 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
512 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
513 pass.SetPipeline(pipeline);
514 pass.SetBindGroup(0, bindGroup);
515 pass.Dispatch(1);
516 pass.EndPass();
517
518 commands = encoder.Finish();
519 }
520
521 queue.Submit(1, &commands);
522
523 EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kCount);
524 }
525
526 // Test overridable constants precision
527 // D3D12 HLSL shader uses defines so we want float number to have enough precision
TEST_P(ShaderTests,OverridableConstantsPrecision)528 TEST_P(ShaderTests, OverridableConstantsPrecision) {
529 DAWN_TEST_UNSUPPORTED_IF(IsOpenGL());
530 DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
531
532 uint32_t const kCount = 2;
533 float const kValue1 = 3.14159;
534 float const kValue2 = 3.141592653589793238;
535 std::vector<float> expected{kValue1, kValue2};
536 wgpu::Buffer buffer = CreateBuffer(kCount);
537
538 std::string shader = R"(
539 [[override(1001)]] let c1: f32;
540 [[override(1002)]] let c2: f32;
541
542 [[block]] struct Buf {
543 data : array<f32, 2>;
544 };
545
546 [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
547
548 [[stage(compute), workgroup_size(1)]] fn main() {
549 buf.data[0] = c1;
550 buf.data[1] = c2;
551 })";
552
553 std::vector<wgpu::ConstantEntry> constants;
554 constants.push_back({nullptr, "1001", kValue1});
555 constants.push_back({nullptr, "1002", kValue2});
556 wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants);
557
558 wgpu::BindGroup bindGroup =
559 utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}});
560
561 wgpu::CommandBuffer commands;
562 {
563 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
564 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
565 pass.SetPipeline(pipeline);
566 pass.SetBindGroup(0, bindGroup);
567 pass.Dispatch(1);
568 pass.EndPass();
569
570 commands = encoder.Finish();
571 }
572
573 queue.Submit(1, &commands);
574
575 EXPECT_BUFFER_FLOAT_RANGE_EQ(expected.data(), buffer, 0, kCount);
576 }
577
578 // Test overridable constants for different entry points
TEST_P(ShaderTests,OverridableConstantsMultipleEntryPoints)579 TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) {
580 DAWN_TEST_UNSUPPORTED_IF(IsOpenGL());
581 DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
582
583 uint32_t const kCount = 1;
584 std::vector<uint32_t> expected1{1u};
585 std::vector<uint32_t> expected2{2u};
586 std::vector<uint32_t> expected3{3u};
587
588 wgpu::Buffer buffer1 = CreateBuffer(kCount);
589 wgpu::Buffer buffer2 = CreateBuffer(kCount);
590 wgpu::Buffer buffer3 = CreateBuffer(kCount);
591
592 std::string shader = R"(
593 [[override(1001)]] let c1: u32;
594 [[override(1002)]] let c2: u32;
595
596 [[block]] struct Buf {
597 data : array<u32, 1>;
598 };
599
600 [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
601
602 [[stage(compute), workgroup_size(1)]] fn main1() {
603 buf.data[0] = c1;
604 }
605
606 [[stage(compute), workgroup_size(1)]] fn main2() {
607 buf.data[0] = c2;
608 }
609
610 [[stage(compute), workgroup_size(1)]] fn main3() {
611 buf.data[0] = 3u;
612 }
613 )";
614
615 std::vector<wgpu::ConstantEntry> constants1;
616 constants1.push_back({nullptr, "1001", 1});
617 std::vector<wgpu::ConstantEntry> constants2;
618 constants2.push_back({nullptr, "1002", 2});
619
620 wgpu::ShaderModule shaderModule = utils::CreateShaderModule(device, shader.c_str());
621
622 wgpu::ComputePipelineDescriptor csDesc1;
623 csDesc1.compute.module = shaderModule;
624 csDesc1.compute.entryPoint = "main1";
625 csDesc1.compute.constants = constants1.data();
626 csDesc1.compute.constantCount = constants1.size();
627 wgpu::ComputePipeline pipeline1 = device.CreateComputePipeline(&csDesc1);
628
629 wgpu::ComputePipelineDescriptor csDesc2;
630 csDesc2.compute.module = shaderModule;
631 csDesc2.compute.entryPoint = "main2";
632 csDesc2.compute.constants = constants2.data();
633 csDesc2.compute.constantCount = constants2.size();
634 wgpu::ComputePipeline pipeline2 = device.CreateComputePipeline(&csDesc2);
635
636 wgpu::ComputePipelineDescriptor csDesc3;
637 csDesc3.compute.module = shaderModule;
638 csDesc3.compute.entryPoint = "main3";
639 wgpu::ComputePipeline pipeline3 = device.CreateComputePipeline(&csDesc3);
640
641 wgpu::BindGroup bindGroup1 =
642 utils::MakeBindGroup(device, pipeline1.GetBindGroupLayout(0), {{0, buffer1}});
643 wgpu::BindGroup bindGroup2 =
644 utils::MakeBindGroup(device, pipeline2.GetBindGroupLayout(0), {{0, buffer2}});
645 wgpu::BindGroup bindGroup3 =
646 utils::MakeBindGroup(device, pipeline3.GetBindGroupLayout(0), {{0, buffer3}});
647
648 wgpu::CommandBuffer commands;
649 {
650 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
651 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
652 pass.SetPipeline(pipeline1);
653 pass.SetBindGroup(0, bindGroup1);
654 pass.Dispatch(1);
655
656 pass.SetPipeline(pipeline2);
657 pass.SetBindGroup(0, bindGroup2);
658 pass.Dispatch(1);
659
660 pass.SetPipeline(pipeline3);
661 pass.SetBindGroup(0, bindGroup3);
662 pass.Dispatch(1);
663
664 pass.EndPass();
665
666 commands = encoder.Finish();
667 }
668
669 queue.Submit(1, &commands);
670
671 EXPECT_BUFFER_U32_RANGE_EQ(expected1.data(), buffer1, 0, kCount);
672 EXPECT_BUFFER_U32_RANGE_EQ(expected2.data(), buffer2, 0, kCount);
673 EXPECT_BUFFER_U32_RANGE_EQ(expected3.data(), buffer3, 0, kCount);
674 }
675
676 // Test overridable constants with render pipeline
677 // Draw a triangle covering the render target, with vertex position and color values from
678 // overridable constants
TEST_P(ShaderTests,OverridableConstantsRenderPipeline)679 TEST_P(ShaderTests, OverridableConstantsRenderPipeline) {
680 DAWN_TEST_UNSUPPORTED_IF(IsOpenGL());
681 DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
682
683 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
684 [[override(1111)]] let xright: f32;
685 [[override(2222)]] let ytop: f32;
686 [[stage(vertex)]]
687 fn main([[builtin(vertex_index)]] VertexIndex : u32)
688 -> [[builtin(position)]] vec4<f32> {
689 var pos = array<vec2<f32>, 3>(
690 vec2<f32>(-1.0, ytop),
691 vec2<f32>(-1.0, -ytop),
692 vec2<f32>(xright, 0.0));
693
694 return vec4<f32>(pos[VertexIndex], 0.0, 1.0);
695 })");
696
697 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
698 [[override(1000)]] let intensity: f32 = 0.0;
699 [[stage(fragment)]] fn main()
700 -> [[location(0)]] vec4<f32> {
701 return vec4<f32>(intensity, intensity, intensity, 1.0);
702 })");
703
704 utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1);
705
706 utils::ComboRenderPipelineDescriptor descriptor;
707 descriptor.vertex.module = vsModule;
708 descriptor.cFragment.module = fsModule;
709 descriptor.primitive.topology = wgpu::PrimitiveTopology::TriangleList;
710 descriptor.cTargets[0].format = renderPass.colorFormat;
711
712 std::vector<wgpu::ConstantEntry> vertexConstants;
713 vertexConstants.push_back({nullptr, "1111", 3.0}); // x right
714 vertexConstants.push_back({nullptr, "2222", 3.0}); // y top
715 descriptor.vertex.constants = vertexConstants.data();
716 descriptor.vertex.constantCount = vertexConstants.size();
717 std::vector<wgpu::ConstantEntry> fragmentConstants;
718 fragmentConstants.push_back({nullptr, "1000", 1.0}); // color intensity
719 descriptor.cFragment.constants = fragmentConstants.data();
720 descriptor.cFragment.constantCount = fragmentConstants.size();
721
722 wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&descriptor);
723
724 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
725 wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo);
726 pass.SetPipeline(pipeline);
727 pass.Draw(3);
728 pass.EndPass();
729 wgpu::CommandBuffer commands = encoder.Finish();
730 queue.Submit(1, &commands);
731
732 EXPECT_PIXEL_RGBA8_EQ(RGBA8(255, 255, 255, 255), renderPass.color, 0, 0);
733 }
734
735 // TODO(tint:1155): Test overridable constants used for workgroup size
736
737 DAWN_INSTANTIATE_TEST(ShaderTests,
738 D3D12Backend(),
739 MetalBackend(),
740 OpenGLBackend(),
741 OpenGLESBackend(),
742 VulkanBackend());
743