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/writer/spirv/spv_dump.h"
18 #include "src/writer/spirv/test_helper.h"
19
20 namespace tint {
21 namespace writer {
22 namespace spirv {
23 namespace {
24
25 using BuilderTest = TestHelper;
26
TEST_F(BuilderTest,Function_Empty)27 TEST_F(BuilderTest, Function_Empty) {
28 Func("a_func", {}, ty.void_(), ast::StatementList{}, ast::DecorationList{});
29
30 spirv::Builder& b = Build();
31
32 auto* func = program->AST().Functions()[0];
33 ASSERT_TRUE(b.GenerateFunction(func));
34 EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "a_func"
35 %2 = OpTypeVoid
36 %1 = OpTypeFunction %2
37 %3 = OpFunction %2 None %1
38 %4 = OpLabel
39 OpReturn
40 OpFunctionEnd
41 )");
42 }
43
TEST_F(BuilderTest,Function_Terminator_Return)44 TEST_F(BuilderTest, Function_Terminator_Return) {
45 Func("a_func", {}, ty.void_(),
46 ast::StatementList{
47 Return(),
48 },
49 ast::DecorationList{});
50
51 spirv::Builder& b = Build();
52
53 auto* func = program->AST().Functions()[0];
54 ASSERT_TRUE(b.GenerateFunction(func));
55 EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "a_func"
56 %2 = OpTypeVoid
57 %1 = OpTypeFunction %2
58 %3 = OpFunction %2 None %1
59 %4 = OpLabel
60 OpReturn
61 OpFunctionEnd
62 )");
63 }
64
TEST_F(BuilderTest,Function_Terminator_ReturnValue)65 TEST_F(BuilderTest, Function_Terminator_ReturnValue) {
66 Global("a", ty.f32(), ast::StorageClass::kPrivate);
67
68 Func("a_func", {}, ty.f32(), ast::StatementList{Return("a")},
69 ast::DecorationList{});
70
71 spirv::Builder& b = Build();
72
73 auto* var_a = program->AST().GlobalVariables()[0];
74 auto* func = program->AST().Functions()[0];
75
76 ASSERT_TRUE(b.GenerateGlobalVariable(var_a)) << b.error();
77 ASSERT_TRUE(b.GenerateFunction(func)) << b.error();
78 EXPECT_EQ(DumpBuilder(b), R"(OpName %1 "a"
79 OpName %6 "a_func"
80 %3 = OpTypeFloat 32
81 %2 = OpTypePointer Private %3
82 %4 = OpConstantNull %3
83 %1 = OpVariable %2 Private %4
84 %5 = OpTypeFunction %3
85 %6 = OpFunction %3 None %5
86 %7 = OpLabel
87 %8 = OpLoad %3 %1
88 OpReturnValue %8
89 OpFunctionEnd
90 )");
91 }
92
TEST_F(BuilderTest,Function_Terminator_Discard)93 TEST_F(BuilderTest, Function_Terminator_Discard) {
94 Func("a_func", {}, ty.void_(),
95 ast::StatementList{
96 create<ast::DiscardStatement>(),
97 },
98 ast::DecorationList{});
99
100 spirv::Builder& b = Build();
101
102 auto* func = program->AST().Functions()[0];
103 ASSERT_TRUE(b.GenerateFunction(func));
104 EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "a_func"
105 %2 = OpTypeVoid
106 %1 = OpTypeFunction %2
107 %3 = OpFunction %2 None %1
108 %4 = OpLabel
109 OpKill
110 OpFunctionEnd
111 )");
112 }
113
TEST_F(BuilderTest,Function_WithParams)114 TEST_F(BuilderTest, Function_WithParams) {
115 ast::VariableList params = {Param("a", ty.f32()), Param("b", ty.i32())};
116
117 Func("a_func", params, ty.f32(), ast::StatementList{Return("a")},
118 ast::DecorationList{});
119
120 spirv::Builder& b = Build();
121
122 auto* func = program->AST().Functions()[0];
123 ASSERT_TRUE(b.GenerateFunction(func));
124 EXPECT_EQ(DumpBuilder(b), R"(OpName %4 "a_func"
125 OpName %5 "a"
126 OpName %6 "b"
127 %2 = OpTypeFloat 32
128 %3 = OpTypeInt 32 1
129 %1 = OpTypeFunction %2 %2 %3
130 %4 = OpFunction %2 None %1
131 %5 = OpFunctionParameter %2
132 %6 = OpFunctionParameter %3
133 %7 = OpLabel
134 OpReturnValue %5
135 OpFunctionEnd
136 )") << DumpBuilder(b);
137 }
138
TEST_F(BuilderTest,Function_WithBody)139 TEST_F(BuilderTest, Function_WithBody) {
140 Func("a_func", {}, ty.void_(),
141 ast::StatementList{
142 Return(),
143 },
144 ast::DecorationList{});
145
146 spirv::Builder& b = Build();
147
148 auto* func = program->AST().Functions()[0];
149 ASSERT_TRUE(b.GenerateFunction(func));
150 EXPECT_EQ(DumpBuilder(b), R"(OpName %3 "a_func"
151 %2 = OpTypeVoid
152 %1 = OpTypeFunction %2
153 %3 = OpFunction %2 None %1
154 %4 = OpLabel
155 OpReturn
156 OpFunctionEnd
157 )");
158 }
159
TEST_F(BuilderTest,FunctionType)160 TEST_F(BuilderTest, FunctionType) {
161 Func("a_func", {}, ty.void_(), ast::StatementList{}, ast::DecorationList{});
162
163 spirv::Builder& b = Build();
164
165 auto* func = program->AST().Functions()[0];
166 ASSERT_TRUE(b.GenerateFunction(func));
167 EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeVoid
168 %1 = OpTypeFunction %2
169 )");
170 }
171
TEST_F(BuilderTest,FunctionType_DeDuplicate)172 TEST_F(BuilderTest, FunctionType_DeDuplicate) {
173 auto* func1 = Func("a_func", {}, ty.void_(), ast::StatementList{},
174 ast::DecorationList{});
175 auto* func2 = Func("b_func", {}, ty.void_(), ast::StatementList{},
176 ast::DecorationList{});
177
178 spirv::Builder& b = Build();
179
180 ASSERT_TRUE(b.GenerateFunction(func1));
181 ASSERT_TRUE(b.GenerateFunction(func2));
182 EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeVoid
183 %1 = OpTypeFunction %2
184 )");
185 }
186
187 // https://crbug.com/tint/297
TEST_F(BuilderTest,Emit_Multiple_EntryPoint_With_Same_ModuleVar)188 TEST_F(BuilderTest, Emit_Multiple_EntryPoint_With_Same_ModuleVar) {
189 // [[block]] struct Data {
190 // d : f32;
191 // };
192 // [[binding(0), group(0)]] var<storage> data : Data;
193 //
194 // [[stage(compute), workgroup_size(1)]]
195 // fn a() {
196 // return;
197 // }
198 //
199 // [[stage(compute), workgroup_size(1)]]
200 // fn b() {
201 // return;
202 // }
203
204 auto* s = Structure("Data", {Member("d", ty.f32())},
205 {create<ast::StructBlockDecoration>()});
206
207 Global("data", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
208 ast::DecorationList{
209 create<ast::BindingDecoration>(0),
210 create<ast::GroupDecoration>(0),
211 });
212
213 {
214 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
215 MemberAccessor("data", "d"));
216
217 Func("a", ast::VariableList{}, ty.void_(),
218 ast::StatementList{
219 Decl(var),
220 Return(),
221 },
222 ast::DecorationList{Stage(ast::PipelineStage::kCompute),
223 WorkgroupSize(1)});
224 }
225
226 {
227 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
228 MemberAccessor("data", "d"));
229
230 Func("b", ast::VariableList{}, ty.void_(),
231 ast::StatementList{
232 Decl(var),
233 Return(),
234 },
235 ast::DecorationList{Stage(ast::PipelineStage::kCompute),
236 WorkgroupSize(1)});
237 }
238
239 spirv::Builder& b = Build();
240
241 ASSERT_TRUE(b.Build());
242 EXPECT_EQ(DumpBuilder(b), R"(OpCapability Shader
243 OpMemoryModel Logical GLSL450
244 OpEntryPoint GLCompute %7 "a"
245 OpEntryPoint GLCompute %17 "b"
246 OpExecutionMode %7 LocalSize 1 1 1
247 OpExecutionMode %17 LocalSize 1 1 1
248 OpName %3 "Data"
249 OpMemberName %3 0 "d"
250 OpName %1 "data"
251 OpName %7 "a"
252 OpName %14 "v"
253 OpName %17 "b"
254 OpName %21 "v"
255 OpDecorate %3 Block
256 OpMemberDecorate %3 0 Offset 0
257 OpDecorate %1 Binding 0
258 OpDecorate %1 DescriptorSet 0
259 %4 = OpTypeFloat 32
260 %3 = OpTypeStruct %4
261 %2 = OpTypePointer StorageBuffer %3
262 %1 = OpVariable %2 StorageBuffer
263 %6 = OpTypeVoid
264 %5 = OpTypeFunction %6
265 %9 = OpTypeInt 32 0
266 %10 = OpConstant %9 0
267 %11 = OpTypePointer StorageBuffer %4
268 %15 = OpTypePointer Function %4
269 %16 = OpConstantNull %4
270 %7 = OpFunction %6 None %5
271 %8 = OpLabel
272 %14 = OpVariable %15 Function %16
273 %12 = OpAccessChain %11 %1 %10
274 %13 = OpLoad %4 %12
275 OpStore %14 %13
276 OpReturn
277 OpFunctionEnd
278 %17 = OpFunction %6 None %5
279 %18 = OpLabel
280 %21 = OpVariable %15 Function %16
281 %19 = OpAccessChain %11 %1 %10
282 %20 = OpLoad %4 %19
283 OpStore %21 %20
284 OpReturn
285 OpFunctionEnd
286 )");
287 }
288
289 } // namespace
290 } // namespace spirv
291 } // namespace writer
292 } // namespace tint
293