1 // Copyright 2020 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/ast/stage_decoration.h"
16 #include "src/ast/struct_block_decoration.h"
17 #include "src/ast/variable_decl_statement.h"
18 #include "src/ast/workgroup_decoration.h"
19 #include "src/writer/wgsl/test_helper.h"
20
21 namespace tint {
22 namespace writer {
23 namespace wgsl {
24 namespace {
25
26 using WgslGeneratorImplTest = TestHelper;
27
TEST_F(WgslGeneratorImplTest,Emit_Function)28 TEST_F(WgslGeneratorImplTest, Emit_Function) {
29 auto* func = Func("my_func", ast::VariableList{}, ty.void_(),
30 ast::StatementList{
31 Return(),
32 },
33 ast::DecorationList{});
34
35 GeneratorImpl& gen = Build();
36
37 gen.increment_indent();
38
39 ASSERT_TRUE(gen.EmitFunction(func));
40 EXPECT_EQ(gen.result(), R"( fn my_func() {
41 return;
42 }
43 )");
44 }
45
TEST_F(WgslGeneratorImplTest,Emit_Function_WithParams)46 TEST_F(WgslGeneratorImplTest, Emit_Function_WithParams) {
47 auto* func = Func(
48 "my_func", ast::VariableList{Param("a", ty.f32()), Param("b", ty.i32())},
49 ty.void_(),
50 ast::StatementList{
51 Return(),
52 },
53 ast::DecorationList{});
54
55 GeneratorImpl& gen = Build();
56
57 gen.increment_indent();
58
59 ASSERT_TRUE(gen.EmitFunction(func));
60 EXPECT_EQ(gen.result(), R"( fn my_func(a : f32, b : i32) {
61 return;
62 }
63 )");
64 }
65
TEST_F(WgslGeneratorImplTest,Emit_Function_WithDecoration_WorkgroupSize)66 TEST_F(WgslGeneratorImplTest, Emit_Function_WithDecoration_WorkgroupSize) {
67 auto* func = Func("my_func", ast::VariableList{}, ty.void_(),
68 ast::StatementList{Return()},
69 ast::DecorationList{
70 Stage(ast::PipelineStage::kCompute),
71 WorkgroupSize(2, 4, 6),
72 });
73
74 GeneratorImpl& gen = Build();
75
76 gen.increment_indent();
77
78 ASSERT_TRUE(gen.EmitFunction(func));
79 EXPECT_EQ(gen.result(), R"( [[stage(compute), workgroup_size(2, 4, 6)]]
80 fn my_func() {
81 return;
82 }
83 )");
84 }
85
TEST_F(WgslGeneratorImplTest,Emit_Function_WithDecoration_WorkgroupSize_WithIdent)86 TEST_F(WgslGeneratorImplTest,
87 Emit_Function_WithDecoration_WorkgroupSize_WithIdent) {
88 GlobalConst("height", ty.i32(), Expr(2));
89 auto* func = Func("my_func", ast::VariableList{}, ty.void_(),
90 ast::StatementList{Return()},
91 ast::DecorationList{
92 Stage(ast::PipelineStage::kCompute),
93 WorkgroupSize(2, "height"),
94 });
95
96 GeneratorImpl& gen = Build();
97
98 gen.increment_indent();
99
100 ASSERT_TRUE(gen.EmitFunction(func));
101 EXPECT_EQ(gen.result(), R"( [[stage(compute), workgroup_size(2, height)]]
102 fn my_func() {
103 return;
104 }
105 )");
106 }
107
TEST_F(WgslGeneratorImplTest,Emit_Function_EntryPoint_Parameters)108 TEST_F(WgslGeneratorImplTest, Emit_Function_EntryPoint_Parameters) {
109 auto* vec4 = ty.vec4<f32>();
110 auto* coord = Param("coord", vec4, {Builtin(ast::Builtin::kPosition)});
111 auto* loc1 = Param("loc1", ty.f32(), {Location(1u)});
112 auto* func = Func("frag_main", ast::VariableList{coord, loc1}, ty.void_(),
113 ast::StatementList{},
114 ast::DecorationList{
115 Stage(ast::PipelineStage::kFragment),
116 });
117
118 GeneratorImpl& gen = Build();
119
120 gen.increment_indent();
121
122 ASSERT_TRUE(gen.EmitFunction(func));
123 EXPECT_EQ(gen.result(), R"( [[stage(fragment)]]
124 fn frag_main([[builtin(position)]] coord : vec4<f32>, [[location(1)]] loc1 : f32) {
125 }
126 )");
127 }
128
TEST_F(WgslGeneratorImplTest,Emit_Function_EntryPoint_ReturnValue)129 TEST_F(WgslGeneratorImplTest, Emit_Function_EntryPoint_ReturnValue) {
130 auto* func = Func("frag_main", ast::VariableList{}, ty.f32(),
131 ast::StatementList{
132 Return(1.f),
133 },
134 ast::DecorationList{
135 Stage(ast::PipelineStage::kFragment),
136 },
137 ast::DecorationList{
138 Location(1u),
139 });
140
141 GeneratorImpl& gen = Build();
142
143 gen.increment_indent();
144
145 ASSERT_TRUE(gen.EmitFunction(func));
146 EXPECT_EQ(gen.result(), R"( [[stage(fragment)]]
147 fn frag_main() -> [[location(1)]] f32 {
148 return 1.0;
149 }
150 )");
151 }
152
153 // https://crbug.com/tint/297
TEST_F(WgslGeneratorImplTest,Emit_Function_Multiple_EntryPoint_With_Same_ModuleVar)154 TEST_F(WgslGeneratorImplTest,
155 Emit_Function_Multiple_EntryPoint_With_Same_ModuleVar) {
156 // [[block]] struct Data {
157 // d : f32;
158 // };
159 // [[binding(0), group(0)]] var<storage> data : Data;
160 //
161 // [[stage(compute), workgroup_size(1)]]
162 // fn a() {
163 // return;
164 // }
165 //
166 // [[stage(compute), workgroup_size(1)]]
167 // fn b() {
168 // return;
169 // }
170
171 auto* s = Structure("Data", {Member("d", ty.f32())},
172 {create<ast::StructBlockDecoration>()});
173
174 Global("data", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
175 ast::DecorationList{
176 create<ast::BindingDecoration>(0),
177 create<ast::GroupDecoration>(0),
178 });
179
180 {
181 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
182 MemberAccessor("data", "d"));
183
184 Func("a", ast::VariableList{}, ty.void_(),
185 ast::StatementList{
186 Decl(var),
187 Return(),
188 },
189 ast::DecorationList{
190 Stage(ast::PipelineStage::kCompute),
191 WorkgroupSize(1),
192 });
193 }
194
195 {
196 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
197 MemberAccessor("data", "d"));
198
199 Func("b", ast::VariableList{}, ty.void_(),
200 ast::StatementList{
201 Decl(var),
202 Return(),
203 },
204 ast::DecorationList{
205 Stage(ast::PipelineStage::kCompute),
206 WorkgroupSize(1),
207 });
208 }
209
210 GeneratorImpl& gen = Build();
211
212 ASSERT_TRUE(gen.Generate()) << gen.error();
213 EXPECT_EQ(gen.result(), R"([[block]]
214 struct Data {
215 d : f32;
216 };
217
218 [[binding(0), group(0)]] var<storage, read_write> data : Data;
219
220 [[stage(compute), workgroup_size(1)]]
221 fn a() {
222 var v : f32 = data.d;
223 return;
224 }
225
226 [[stage(compute), workgroup_size(1)]]
227 fn b() {
228 var v : f32 = data.d;
229 return;
230 }
231 )");
232 }
233
234 } // namespace
235 } // namespace wgsl
236 } // namespace writer
237 } // namespace tint
238