• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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