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/array_length_from_uniform.h"
16
17 #include <utility>
18
19 #include "src/transform/simplify_pointers.h"
20 #include "src/transform/test_helper.h"
21 #include "src/transform/unshadow.h"
22
23 namespace tint {
24 namespace transform {
25 namespace {
26
27 using ArrayLengthFromUniformTest = TransformTest;
28
TEST_F(ArrayLengthFromUniformTest,Error_MissingTransformData)29 TEST_F(ArrayLengthFromUniformTest, Error_MissingTransformData) {
30 auto* src = "";
31
32 auto* expect =
33 "error: missing transform data for "
34 "tint::transform::ArrayLengthFromUniform";
35
36 auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src);
37
38 EXPECT_EQ(expect, str(got));
39 }
40
TEST_F(ArrayLengthFromUniformTest,Error_MissingSimplifyPointers)41 TEST_F(ArrayLengthFromUniformTest, Error_MissingSimplifyPointers) {
42 auto* src = "";
43
44 auto* expect =
45 "error: tint::transform::ArrayLengthFromUniform depends on "
46 "tint::transform::SimplifyPointers but the dependency was not run";
47
48 auto got = Run<ArrayLengthFromUniform>(src);
49
50 EXPECT_EQ(expect, str(got));
51 }
52
TEST_F(ArrayLengthFromUniformTest,Basic)53 TEST_F(ArrayLengthFromUniformTest, Basic) {
54 auto* src = R"(
55 [[block]]
56 struct SB {
57 x : i32;
58 arr : array<i32>;
59 };
60
61 [[group(0), binding(0)]] var<storage, read> sb : SB;
62
63 [[stage(compute), workgroup_size(1)]]
64 fn main() {
65 var len : u32 = arrayLength(&sb.arr);
66 }
67 )";
68
69 auto* expect = R"(
70 [[block]]
71 struct tint_symbol {
72 buffer_size : array<vec4<u32>, 1u>;
73 };
74
75 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
76
77 [[block]]
78 struct SB {
79 x : i32;
80 arr : array<i32>;
81 };
82
83 [[group(0), binding(0)]] var<storage, read> sb : SB;
84
85 [[stage(compute), workgroup_size(1)]]
86 fn main() {
87 var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
88 }
89 )";
90
91 ArrayLengthFromUniform::Config cfg({0, 30u});
92 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
93
94 DataMap data;
95 data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
96
97 auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
98
99 EXPECT_EQ(expect, str(got));
100 EXPECT_EQ(std::unordered_set<uint32_t>({0}),
101 got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
102 }
103
TEST_F(ArrayLengthFromUniformTest,WithStride)104 TEST_F(ArrayLengthFromUniformTest, WithStride) {
105 auto* src = R"(
106 [[block]]
107 struct SB {
108 x : i32;
109 y : f32;
110 arr : [[stride(64)]] array<i32>;
111 };
112
113 [[group(0), binding(0)]] var<storage, read> sb : SB;
114
115 [[stage(compute), workgroup_size(1)]]
116 fn main() {
117 var len : u32 = arrayLength(&sb.arr);
118 }
119 )";
120
121 auto* expect = R"(
122 [[block]]
123 struct tint_symbol {
124 buffer_size : array<vec4<u32>, 1u>;
125 };
126
127 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
128
129 [[block]]
130 struct SB {
131 x : i32;
132 y : f32;
133 arr : [[stride(64)]] array<i32>;
134 };
135
136 [[group(0), binding(0)]] var<storage, read> sb : SB;
137
138 [[stage(compute), workgroup_size(1)]]
139 fn main() {
140 var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u);
141 }
142 )";
143
144 ArrayLengthFromUniform::Config cfg({0, 30u});
145 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
146
147 DataMap data;
148 data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
149
150 auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
151
152 EXPECT_EQ(expect, str(got));
153 EXPECT_EQ(std::unordered_set<uint32_t>({0}),
154 got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
155 }
156
TEST_F(ArrayLengthFromUniformTest,MultipleStorageBuffers)157 TEST_F(ArrayLengthFromUniformTest, MultipleStorageBuffers) {
158 auto* src = R"(
159 [[block]]
160 struct SB1 {
161 x : i32;
162 arr1 : array<i32>;
163 };
164 [[block]]
165 struct SB2 {
166 x : i32;
167 arr2 : array<vec4<f32>>;
168 };
169 [[block]]
170 struct SB3 {
171 x : i32;
172 arr3 : array<vec4<f32>>;
173 };
174 [[block]]
175 struct SB4 {
176 x : i32;
177 arr4 : array<vec4<f32>>;
178 };
179 [[block]]
180 struct SB5 {
181 x : i32;
182 arr5 : array<vec4<f32>>;
183 };
184
185 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
186 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
187 [[group(2), binding(2)]] var<storage, read> sb3 : SB3;
188 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
189 [[group(4), binding(2)]] var<storage, read> sb5 : SB5;
190
191 [[stage(compute), workgroup_size(1)]]
192 fn main() {
193 var len1 : u32 = arrayLength(&(sb1.arr1));
194 var len2 : u32 = arrayLength(&(sb2.arr2));
195 var len3 : u32 = arrayLength(&(sb3.arr3));
196 var len4 : u32 = arrayLength(&(sb4.arr4));
197 var len5 : u32 = arrayLength(&(sb5.arr5));
198 var x : u32 = (len1 + len2 + len3 + len4 + len5);
199 }
200 )";
201
202 auto* expect = R"(
203 [[block]]
204 struct tint_symbol {
205 buffer_size : array<vec4<u32>, 2u>;
206 };
207
208 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
209
210 [[block]]
211 struct SB1 {
212 x : i32;
213 arr1 : array<i32>;
214 };
215
216 [[block]]
217 struct SB2 {
218 x : i32;
219 arr2 : array<vec4<f32>>;
220 };
221
222 [[block]]
223 struct SB3 {
224 x : i32;
225 arr3 : array<vec4<f32>>;
226 };
227
228 [[block]]
229 struct SB4 {
230 x : i32;
231 arr4 : array<vec4<f32>>;
232 };
233
234 [[block]]
235 struct SB5 {
236 x : i32;
237 arr5 : array<vec4<f32>>;
238 };
239
240 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
241
242 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
243
244 [[group(2), binding(2)]] var<storage, read> sb3 : SB3;
245
246 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
247
248 [[group(4), binding(2)]] var<storage, read> sb5 : SB5;
249
250 [[stage(compute), workgroup_size(1)]]
251 fn main() {
252 var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
253 var len2 : u32 = ((tint_symbol_1.buffer_size[0u][1u] - 16u) / 16u);
254 var len3 : u32 = ((tint_symbol_1.buffer_size[0u][2u] - 16u) / 16u);
255 var len4 : u32 = ((tint_symbol_1.buffer_size[0u][3u] - 16u) / 16u);
256 var len5 : u32 = ((tint_symbol_1.buffer_size[1u][0u] - 16u) / 16u);
257 var x : u32 = ((((len1 + len2) + len3) + len4) + len5);
258 }
259 )";
260
261 ArrayLengthFromUniform::Config cfg({0, 30u});
262 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 2u}, 0);
263 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{1u, 2u}, 1);
264 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{2u, 2u}, 2);
265 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{3u, 2u}, 3);
266 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{4u, 2u}, 4);
267
268 DataMap data;
269 data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
270
271 auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
272
273 EXPECT_EQ(expect, str(got));
274 EXPECT_EQ(std::unordered_set<uint32_t>({0, 1, 2, 3, 4}),
275 got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
276 }
277
TEST_F(ArrayLengthFromUniformTest,MultipleUnusedStorageBuffers)278 TEST_F(ArrayLengthFromUniformTest, MultipleUnusedStorageBuffers) {
279 auto* src = R"(
280 [[block]]
281 struct SB1 {
282 x : i32;
283 arr1 : array<i32>;
284 };
285 [[block]]
286 struct SB2 {
287 x : i32;
288 arr2 : array<vec4<f32>>;
289 };
290 [[block]]
291 struct SB3 {
292 x : i32;
293 arr3 : array<vec4<f32>>;
294 };
295 [[block]]
296 struct SB4 {
297 x : i32;
298 arr4 : array<vec4<f32>>;
299 };
300 [[block]]
301 struct SB5 {
302 x : i32;
303 arr5 : array<vec4<f32>>;
304 };
305
306 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
307 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
308 [[group(2), binding(2)]] var<storage, read> sb3 : SB3;
309 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
310 [[group(4), binding(2)]] var<storage, read> sb5 : SB5;
311
312 [[stage(compute), workgroup_size(1)]]
313 fn main() {
314 var len1 : u32 = arrayLength(&(sb1.arr1));
315 var len3 : u32 = arrayLength(&(sb3.arr3));
316 var x : u32 = (len1 + len3);
317 }
318 )";
319
320 auto* expect = R"(
321 [[block]]
322 struct tint_symbol {
323 buffer_size : array<vec4<u32>, 1u>;
324 };
325
326 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
327
328 [[block]]
329 struct SB1 {
330 x : i32;
331 arr1 : array<i32>;
332 };
333
334 [[block]]
335 struct SB2 {
336 x : i32;
337 arr2 : array<vec4<f32>>;
338 };
339
340 [[block]]
341 struct SB3 {
342 x : i32;
343 arr3 : array<vec4<f32>>;
344 };
345
346 [[block]]
347 struct SB4 {
348 x : i32;
349 arr4 : array<vec4<f32>>;
350 };
351
352 [[block]]
353 struct SB5 {
354 x : i32;
355 arr5 : array<vec4<f32>>;
356 };
357
358 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
359
360 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
361
362 [[group(2), binding(2)]] var<storage, read> sb3 : SB3;
363
364 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
365
366 [[group(4), binding(2)]] var<storage, read> sb5 : SB5;
367
368 [[stage(compute), workgroup_size(1)]]
369 fn main() {
370 var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
371 var len3 : u32 = ((tint_symbol_1.buffer_size[0u][2u] - 16u) / 16u);
372 var x : u32 = (len1 + len3);
373 }
374 )";
375
376 ArrayLengthFromUniform::Config cfg({0, 30u});
377 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 2u}, 0);
378 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{1u, 2u}, 1);
379 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{2u, 2u}, 2);
380 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{3u, 2u}, 3);
381 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{4u, 2u}, 4);
382
383 DataMap data;
384 data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
385
386 auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
387
388 EXPECT_EQ(expect, str(got));
389 EXPECT_EQ(std::unordered_set<uint32_t>({0, 2}),
390 got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
391 }
392
TEST_F(ArrayLengthFromUniformTest,NoArrayLengthCalls)393 TEST_F(ArrayLengthFromUniformTest, NoArrayLengthCalls) {
394 auto* src = R"(
395 [[block]]
396 struct SB {
397 x : i32;
398 arr : array<i32>;
399 };
400
401 [[group(0), binding(0)]] var<storage, read> sb : SB;
402
403 [[stage(compute), workgroup_size(1)]]
404 fn main() {
405 ignore(&(sb.arr));
406 }
407 )";
408
409 ArrayLengthFromUniform::Config cfg({0, 30u});
410 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
411
412 DataMap data;
413 data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
414
415 auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
416
417 EXPECT_EQ(src, str(got));
418 EXPECT_EQ(std::unordered_set<uint32_t>(),
419 got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
420 }
421
TEST_F(ArrayLengthFromUniformTest,MissingBindingPointToIndexMapping)422 TEST_F(ArrayLengthFromUniformTest, MissingBindingPointToIndexMapping) {
423 auto* src = R"(
424 [[block]]
425 struct SB1 {
426 x : i32;
427 arr1 : array<i32>;
428 };
429
430 [[block]]
431 struct SB2 {
432 x : i32;
433 arr2 : array<vec4<f32>>;
434 };
435
436 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
437
438 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
439
440 [[stage(compute), workgroup_size(1)]]
441 fn main() {
442 var len1 : u32 = arrayLength(&(sb1.arr1));
443 var len2 : u32 = arrayLength(&(sb2.arr2));
444 var x : u32 = (len1 + len2);
445 }
446 )";
447
448 auto* expect = R"(
449 [[block]]
450 struct tint_symbol {
451 buffer_size : array<vec4<u32>, 1u>;
452 };
453
454 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
455
456 [[block]]
457 struct SB1 {
458 x : i32;
459 arr1 : array<i32>;
460 };
461
462 [[block]]
463 struct SB2 {
464 x : i32;
465 arr2 : array<vec4<f32>>;
466 };
467
468 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
469
470 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
471
472 [[stage(compute), workgroup_size(1)]]
473 fn main() {
474 var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
475 var len2 : u32 = arrayLength(&(sb2.arr2));
476 var x : u32 = (len1 + len2);
477 }
478 )";
479
480 ArrayLengthFromUniform::Config cfg({0, 30u});
481 cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 2}, 0);
482
483 DataMap data;
484 data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
485
486 auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
487
488 EXPECT_EQ(expect, str(got));
489 EXPECT_EQ(std::unordered_set<uint32_t>({0}),
490 got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
491 }
492
493 } // namespace
494 } // namespace transform
495 } // namespace tint
496