1 // Copyright 2021 The Tint 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 "src/transform/single_entry_point.h"
16
17 #include <utility>
18
19 #include "src/transform/test_helper.h"
20
21 namespace tint {
22 namespace transform {
23 namespace {
24
25 using SingleEntryPointTest = TransformTest;
26
TEST_F(SingleEntryPointTest,Error_MissingTransformData)27 TEST_F(SingleEntryPointTest, Error_MissingTransformData) {
28 auto* src = "";
29
30 auto* expect =
31 "error: missing transform data for tint::transform::SingleEntryPoint";
32
33 auto got = Run<SingleEntryPoint>(src);
34
35 EXPECT_EQ(expect, str(got));
36 }
37
TEST_F(SingleEntryPointTest,Error_NoEntryPoints)38 TEST_F(SingleEntryPointTest, Error_NoEntryPoints) {
39 auto* src = "";
40
41 auto* expect = "error: entry point 'main' not found";
42
43 DataMap data;
44 data.Add<SingleEntryPoint::Config>("main");
45 auto got = Run<SingleEntryPoint>(src, data);
46
47 EXPECT_EQ(expect, str(got));
48 }
49
TEST_F(SingleEntryPointTest,Error_InvalidEntryPoint)50 TEST_F(SingleEntryPointTest, Error_InvalidEntryPoint) {
51 auto* src = R"(
52 [[stage(vertex)]]
53 fn main() -> [[builtin(position)]] vec4<f32> {
54 return vec4<f32>();
55 }
56 )";
57
58 auto* expect = "error: entry point '_' not found";
59
60 SingleEntryPoint::Config cfg("_");
61
62 DataMap data;
63 data.Add<SingleEntryPoint::Config>(cfg);
64 auto got = Run<SingleEntryPoint>(src, data);
65
66 EXPECT_EQ(expect, str(got));
67 }
68
TEST_F(SingleEntryPointTest,Error_NotAnEntryPoint)69 TEST_F(SingleEntryPointTest, Error_NotAnEntryPoint) {
70 auto* src = R"(
71 fn foo() {}
72
73 [[stage(fragment)]]
74 fn main() {}
75 )";
76
77 auto* expect = "error: entry point 'foo' not found";
78
79 SingleEntryPoint::Config cfg("foo");
80
81 DataMap data;
82 data.Add<SingleEntryPoint::Config>(cfg);
83 auto got = Run<SingleEntryPoint>(src, data);
84
85 EXPECT_EQ(expect, str(got));
86 }
87
TEST_F(SingleEntryPointTest,SingleEntryPoint)88 TEST_F(SingleEntryPointTest, SingleEntryPoint) {
89 auto* src = R"(
90 [[stage(compute), workgroup_size(1)]]
91 fn main() {
92 }
93 )";
94
95 SingleEntryPoint::Config cfg("main");
96
97 DataMap data;
98 data.Add<SingleEntryPoint::Config>(cfg);
99 auto got = Run<SingleEntryPoint>(src, data);
100
101 EXPECT_EQ(src, str(got));
102 }
103
TEST_F(SingleEntryPointTest,MultipleEntryPoints)104 TEST_F(SingleEntryPointTest, MultipleEntryPoints) {
105 auto* src = R"(
106 [[stage(vertex)]]
107 fn vert_main() -> [[builtin(position)]] vec4<f32> {
108 return vec4<f32>();
109 }
110
111 [[stage(fragment)]]
112 fn frag_main() {
113 }
114
115 [[stage(compute), workgroup_size(1)]]
116 fn comp_main1() {
117 }
118
119 [[stage(compute), workgroup_size(1)]]
120 fn comp_main2() {
121 }
122 )";
123
124 auto* expect = R"(
125 [[stage(compute), workgroup_size(1)]]
126 fn comp_main1() {
127 }
128 )";
129
130 SingleEntryPoint::Config cfg("comp_main1");
131
132 DataMap data;
133 data.Add<SingleEntryPoint::Config>(cfg);
134 auto got = Run<SingleEntryPoint>(src, data);
135
136 EXPECT_EQ(expect, str(got));
137 }
138
TEST_F(SingleEntryPointTest,GlobalVariables)139 TEST_F(SingleEntryPointTest, GlobalVariables) {
140 auto* src = R"(
141 var<private> a : f32;
142
143 var<private> b : f32;
144
145 var<private> c : f32;
146
147 var<private> d : f32;
148
149 [[stage(vertex)]]
150 fn vert_main() -> [[builtin(position)]] vec4<f32> {
151 a = 0.0;
152 return vec4<f32>();
153 }
154
155 [[stage(fragment)]]
156 fn frag_main() {
157 b = 0.0;
158 }
159
160 [[stage(compute), workgroup_size(1)]]
161 fn comp_main1() {
162 c = 0.0;
163 }
164
165 [[stage(compute), workgroup_size(1)]]
166 fn comp_main2() {
167 d = 0.0;
168 }
169 )";
170
171 auto* expect = R"(
172 var<private> c : f32;
173
174 [[stage(compute), workgroup_size(1)]]
175 fn comp_main1() {
176 c = 0.0;
177 }
178 )";
179
180 SingleEntryPoint::Config cfg("comp_main1");
181
182 DataMap data;
183 data.Add<SingleEntryPoint::Config>(cfg);
184 auto got = Run<SingleEntryPoint>(src, data);
185
186 EXPECT_EQ(expect, str(got));
187 }
188
TEST_F(SingleEntryPointTest,GlobalConstants)189 TEST_F(SingleEntryPointTest, GlobalConstants) {
190 auto* src = R"(
191 let a : f32 = 1.0;
192
193 let b : f32 = 1.0;
194
195 let c : f32 = 1.0;
196
197 let d : f32 = 1.0;
198
199 [[stage(vertex)]]
200 fn vert_main() -> [[builtin(position)]] vec4<f32> {
201 let local_a : f32 = a;
202 return vec4<f32>();
203 }
204
205 [[stage(fragment)]]
206 fn frag_main() {
207 let local_b : f32 = b;
208 }
209
210 [[stage(compute), workgroup_size(1)]]
211 fn comp_main1() {
212 let local_c : f32 = c;
213 }
214
215 [[stage(compute), workgroup_size(1)]]
216 fn comp_main2() {
217 let local_d : f32 = d;
218 }
219 )";
220
221 auto* expect = R"(
222 let c : f32 = 1.0;
223
224 [[stage(compute), workgroup_size(1)]]
225 fn comp_main1() {
226 let local_c : f32 = c;
227 }
228 )";
229
230 SingleEntryPoint::Config cfg("comp_main1");
231
232 DataMap data;
233 data.Add<SingleEntryPoint::Config>(cfg);
234 auto got = Run<SingleEntryPoint>(src, data);
235
236 EXPECT_EQ(expect, str(got));
237 }
238
TEST_F(SingleEntryPointTest,WorkgroupSizeLetPreserved)239 TEST_F(SingleEntryPointTest, WorkgroupSizeLetPreserved) {
240 auto* src = R"(
241 let size : i32 = 1;
242
243 [[stage(compute), workgroup_size(size)]]
244 fn main() {
245 }
246 )";
247
248 auto* expect = src;
249
250 SingleEntryPoint::Config cfg("main");
251
252 DataMap data;
253 data.Add<SingleEntryPoint::Config>(cfg);
254 auto got = Run<SingleEntryPoint>(src, data);
255
256 EXPECT_EQ(expect, str(got));
257 }
258
TEST_F(SingleEntryPointTest,OverridableConstants)259 TEST_F(SingleEntryPointTest, OverridableConstants) {
260 auto* src = R"(
261 [[override(1001)]] let c1 : u32 = 1u;
262 [[override]] let c2 : u32 = 1u;
263 [[override(0)]] let c3 : u32 = 1u;
264 [[override(9999)]] let c4 : u32 = 1u;
265
266 [[stage(compute), workgroup_size(1)]]
267 fn comp_main1() {
268 let local_d = c1;
269 }
270
271 [[stage(compute), workgroup_size(1)]]
272 fn comp_main2() {
273 let local_d = c2;
274 }
275
276 [[stage(compute), workgroup_size(1)]]
277 fn comp_main3() {
278 let local_d = c3;
279 }
280
281 [[stage(compute), workgroup_size(1)]]
282 fn comp_main4() {
283 let local_d = c4;
284 }
285
286 [[stage(compute), workgroup_size(1)]]
287 fn comp_main5() {
288 let local_d = 1u;
289 }
290 )";
291
292 {
293 SingleEntryPoint::Config cfg("comp_main1");
294 auto* expect = R"(
295 [[override(1001)]] let c1 : u32 = 1u;
296
297 [[stage(compute), workgroup_size(1)]]
298 fn comp_main1() {
299 let local_d = c1;
300 }
301 )";
302 DataMap data;
303 data.Add<SingleEntryPoint::Config>(cfg);
304 auto got = Run<SingleEntryPoint>(src, data);
305 EXPECT_EQ(expect, str(got));
306 }
307
308 {
309 SingleEntryPoint::Config cfg("comp_main2");
310 // The decorator is replaced with the one with explicit id
311 // And should not be affected by other constants stripped away
312 auto* expect = R"(
313 [[override(1)]] let c2 : u32 = 1u;
314
315 [[stage(compute), workgroup_size(1)]]
316 fn comp_main2() {
317 let local_d = c2;
318 }
319 )";
320 DataMap data;
321 data.Add<SingleEntryPoint::Config>(cfg);
322 auto got = Run<SingleEntryPoint>(src, data);
323 EXPECT_EQ(expect, str(got));
324 }
325
326 {
327 SingleEntryPoint::Config cfg("comp_main3");
328 auto* expect = R"(
329 [[override(0)]] let c3 : u32 = 1u;
330
331 [[stage(compute), workgroup_size(1)]]
332 fn comp_main3() {
333 let local_d = c3;
334 }
335 )";
336 DataMap data;
337 data.Add<SingleEntryPoint::Config>(cfg);
338 auto got = Run<SingleEntryPoint>(src, data);
339 EXPECT_EQ(expect, str(got));
340 }
341
342 {
343 SingleEntryPoint::Config cfg("comp_main4");
344 auto* expect = R"(
345 [[override(9999)]] let c4 : u32 = 1u;
346
347 [[stage(compute), workgroup_size(1)]]
348 fn comp_main4() {
349 let local_d = c4;
350 }
351 )";
352 DataMap data;
353 data.Add<SingleEntryPoint::Config>(cfg);
354 auto got = Run<SingleEntryPoint>(src, data);
355 EXPECT_EQ(expect, str(got));
356 }
357
358 {
359 SingleEntryPoint::Config cfg("comp_main5");
360 auto* expect = R"(
361 [[stage(compute), workgroup_size(1)]]
362 fn comp_main5() {
363 let local_d = 1u;
364 }
365 )";
366 DataMap data;
367 data.Add<SingleEntryPoint::Config>(cfg);
368 auto got = Run<SingleEntryPoint>(src, data);
369 EXPECT_EQ(expect, str(got));
370 }
371 }
372
TEST_F(SingleEntryPointTest,CalledFunctions)373 TEST_F(SingleEntryPointTest, CalledFunctions) {
374 auto* src = R"(
375 fn inner1() {
376 }
377
378 fn inner2() {
379 }
380
381 fn inner_shared() {
382 }
383
384 fn outer1() {
385 inner1();
386 inner_shared();
387 }
388
389 fn outer2() {
390 inner2();
391 inner_shared();
392 }
393
394 [[stage(compute), workgroup_size(1)]]
395 fn comp_main1() {
396 outer1();
397 }
398
399 [[stage(compute), workgroup_size(1)]]
400 fn comp_main2() {
401 outer2();
402 }
403 )";
404
405 auto* expect = R"(
406 fn inner1() {
407 }
408
409 fn inner_shared() {
410 }
411
412 fn outer1() {
413 inner1();
414 inner_shared();
415 }
416
417 [[stage(compute), workgroup_size(1)]]
418 fn comp_main1() {
419 outer1();
420 }
421 )";
422
423 SingleEntryPoint::Config cfg("comp_main1");
424
425 DataMap data;
426 data.Add<SingleEntryPoint::Config>(cfg);
427 auto got = Run<SingleEntryPoint>(src, data);
428
429 EXPECT_EQ(expect, str(got));
430 }
431
TEST_F(SingleEntryPointTest,GlobalsReferencedByCalledFunctions)432 TEST_F(SingleEntryPointTest, GlobalsReferencedByCalledFunctions) {
433 auto* src = R"(
434 var<private> inner1_var : f32;
435
436 var<private> inner2_var : f32;
437
438 var<private> inner_shared_var : f32;
439
440 var<private> outer1_var : f32;
441
442 var<private> outer2_var : f32;
443
444 fn inner1() {
445 inner1_var = 0.0;
446 }
447
448 fn inner2() {
449 inner2_var = 0.0;
450 }
451
452 fn inner_shared() {
453 inner_shared_var = 0.0;
454 }
455
456 fn outer1() {
457 inner1();
458 inner_shared();
459 outer1_var = 0.0;
460 }
461
462 fn outer2() {
463 inner2();
464 inner_shared();
465 outer2_var = 0.0;
466 }
467
468 [[stage(compute), workgroup_size(1)]]
469 fn comp_main1() {
470 outer1();
471 }
472
473 [[stage(compute), workgroup_size(1)]]
474 fn comp_main2() {
475 outer2();
476 }
477 )";
478
479 auto* expect = R"(
480 var<private> inner1_var : f32;
481
482 var<private> inner_shared_var : f32;
483
484 var<private> outer1_var : f32;
485
486 fn inner1() {
487 inner1_var = 0.0;
488 }
489
490 fn inner_shared() {
491 inner_shared_var = 0.0;
492 }
493
494 fn outer1() {
495 inner1();
496 inner_shared();
497 outer1_var = 0.0;
498 }
499
500 [[stage(compute), workgroup_size(1)]]
501 fn comp_main1() {
502 outer1();
503 }
504 )";
505
506 SingleEntryPoint::Config cfg("comp_main1");
507
508 DataMap data;
509 data.Add<SingleEntryPoint::Config>(cfg);
510 auto got = Run<SingleEntryPoint>(src, data);
511
512 EXPECT_EQ(expect, str(got));
513 }
514
515 } // namespace
516 } // namespace transform
517 } // namespace tint
518