blob: ded0afa596d1e17ea7b22a9299ed43b23b2035ac [file] [log] [blame]
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001// 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/tint/resolver/resolver.h"
16
17#include <algorithm>
18#include <cmath>
19#include <iomanip>
20#include <limits>
21#include <utility>
22
23#include "src/tint/ast/alias.h"
24#include "src/tint/ast/array.h"
25#include "src/tint/ast/assignment_statement.h"
dan sinclair5361d9e2022-08-31 13:39:48 +000026#include "src/tint/ast/attribute.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000027#include "src/tint/ast/bitcast_expression.h"
28#include "src/tint/ast/break_statement.h"
29#include "src/tint/ast/call_statement.h"
30#include "src/tint/ast/continue_statement.h"
31#include "src/tint/ast/depth_texture.h"
32#include "src/tint/ast/disable_validation_attribute.h"
33#include "src/tint/ast/discard_statement.h"
34#include "src/tint/ast/fallthrough_statement.h"
35#include "src/tint/ast/for_loop_statement.h"
36#include "src/tint/ast/id_attribute.h"
37#include "src/tint/ast/if_statement.h"
38#include "src/tint/ast/internal_attribute.h"
39#include "src/tint/ast/interpolate_attribute.h"
40#include "src/tint/ast/loop_statement.h"
41#include "src/tint/ast/matrix.h"
42#include "src/tint/ast/pointer.h"
43#include "src/tint/ast/return_statement.h"
44#include "src/tint/ast/sampled_texture.h"
45#include "src/tint/ast/sampler.h"
46#include "src/tint/ast/storage_texture.h"
47#include "src/tint/ast/switch_statement.h"
48#include "src/tint/ast/traverse_expressions.h"
49#include "src/tint/ast/type_name.h"
50#include "src/tint/ast/unary_op_expression.h"
51#include "src/tint/ast/variable_decl_statement.h"
52#include "src/tint/ast/vector.h"
dan sinclair49d1a2d2022-06-16 12:01:27 +000053#include "src/tint/ast/while_statement.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000054#include "src/tint/ast/workgroup_attribute.h"
James Pricebe656f72022-05-11 22:05:15 +000055#include "src/tint/resolver/uniformity.h"
Ben Clayton65228372022-05-20 17:18:50 +000056#include "src/tint/sem/abstract_float.h"
57#include "src/tint/sem/abstract_int.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000058#include "src/tint/sem/array.h"
Ben Clayton01004b72022-04-28 18:49:04 +000059#include "src/tint/sem/atomic.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000060#include "src/tint/sem/call.h"
Ben Clayton01004b72022-04-28 18:49:04 +000061#include "src/tint/sem/depth_multisampled_texture.h"
62#include "src/tint/sem/depth_texture.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000063#include "src/tint/sem/for_loop_statement.h"
64#include "src/tint/sem/function.h"
65#include "src/tint/sem/if_statement.h"
Antonio Maioranodfeaf2902022-06-24 20:34:00 +000066#include "src/tint/sem/index_accessor_expression.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000067#include "src/tint/sem/loop_statement.h"
Ben Clayton932418e2022-05-24 21:42:03 +000068#include "src/tint/sem/materialize.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000069#include "src/tint/sem/member_accessor_expression.h"
70#include "src/tint/sem/module.h"
Ben Clayton01004b72022-04-28 18:49:04 +000071#include "src/tint/sem/multisampled_texture.h"
72#include "src/tint/sem/pointer.h"
73#include "src/tint/sem/reference.h"
74#include "src/tint/sem/sampled_texture.h"
75#include "src/tint/sem/sampler.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000076#include "src/tint/sem/statement.h"
Ben Clayton01004b72022-04-28 18:49:04 +000077#include "src/tint/sem/storage_texture.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000078#include "src/tint/sem/struct.h"
79#include "src/tint/sem/switch_statement.h"
80#include "src/tint/sem/type_constructor.h"
81#include "src/tint/sem/type_conversion.h"
82#include "src/tint/sem/variable.h"
dan sinclair49d1a2d2022-06-16 12:01:27 +000083#include "src/tint/sem/while_statement.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000084#include "src/tint/utils/defer.h"
85#include "src/tint/utils/math.h"
86#include "src/tint/utils/reverse.h"
87#include "src/tint/utils/scoped_assignment.h"
88#include "src/tint/utils/transform.h"
Ben Clayton958a4642022-07-26 07:55:24 +000089#include "src/tint/utils/vector.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000090
dan sinclaird2093792022-04-07 17:45:45 +000091namespace tint::resolver {
Ryan Harrisondbc13af2022-02-21 15:19:07 +000092
Ben Clayton0d757d22022-06-01 20:34:40 +000093Resolver::Resolver(ProgramBuilder* builder)
Ryan Harrisondbc13af2022-02-21 15:19:07 +000094 : builder_(builder),
95 diagnostics_(builder->Diagnostics()),
Ben Clayton65c5c9d2022-07-15 14:14:09 +000096 const_eval_(*builder),
Ben Claytone1325162022-05-06 15:13:01 +000097 intrinsic_table_(IntrinsicTable::Create(*builder)),
dan sinclairf05575b2022-04-21 13:40:16 +000098 sem_(builder, dependencies_),
Ben Clayton0d757d22022-06-01 20:34:40 +000099 validator_(builder, sem_) {}
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000100
101Resolver::~Resolver() = default;
102
103bool Resolver::Resolve() {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000104 if (builder_->Diagnostics().contains_errors()) {
105 return false;
106 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000107
Ben Clayton4a92a3c2022-07-18 20:50:02 +0000108 builder_->Sem().Reserve(builder_->LastAllocatedNodeID());
Ben Claytonb71bcdd2022-07-18 16:47:32 +0000109
Ben Claytonaf473882022-07-21 23:46:15 +0000110 // Pre-allocate the marked bitset with the total number of AST nodes.
111 marked_.Resize(builder_->ASTNodes().Count());
112
dan sinclair41e4d9a2022-05-01 14:40:55 +0000113 if (!DependencyGraph::Build(builder_->AST(), builder_->Symbols(), builder_->Diagnostics(),
114 dependencies_)) {
115 return false;
116 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000117
dan sinclair41e4d9a2022-05-01 14:40:55 +0000118 bool result = ResolveInternal();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000119
dan sinclair41e4d9a2022-05-01 14:40:55 +0000120 if (!result && !diagnostics_.contains_errors()) {
121 TINT_ICE(Resolver, diagnostics_) << "resolving failed, but no error was raised";
122 return false;
123 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000124
Ben Clayton7f2b8cd2022-05-18 22:41:48 +0000125 // Create the semantic module
126 builder_->Sem().SetModule(builder_->create<sem::Module>(
127 std::move(dependencies_.ordered_globals), std::move(enabled_extensions_)));
128
dan sinclair41e4d9a2022-05-01 14:40:55 +0000129 return result;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000130}
131
132bool Resolver::ResolveInternal() {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000133 Mark(&builder_->AST());
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000134
dan sinclair41e4d9a2022-05-01 14:40:55 +0000135 // Process all module-scope declarations in dependency order.
136 for (auto* decl : dependencies_.ordered_globals) {
137 Mark(decl);
Ben Clayton7f2b8cd2022-05-18 22:41:48 +0000138 if (!Switch<bool>(
dan sinclair41e4d9a2022-05-01 14:40:55 +0000139 decl, //
Ben Clayton7f2b8cd2022-05-18 22:41:48 +0000140 [&](const ast::Enable* e) { return Enable(e); },
dan sinclair41e4d9a2022-05-01 14:40:55 +0000141 [&](const ast::TypeDecl* td) { return TypeDecl(td); },
142 [&](const ast::Function* func) { return Function(func); },
143 [&](const ast::Variable* var) { return GlobalVariable(var); },
Ben Clayton02791e92022-08-02 23:28:28 +0000144 [&](const ast::StaticAssert* sa) { return StaticAssert(sa); },
dan sinclair41e4d9a2022-05-01 14:40:55 +0000145 [&](Default) {
146 TINT_UNREACHABLE(Resolver, diagnostics_)
147 << "unhandled global declaration: " << decl->TypeInfo().name;
Ben Clayton7f2b8cd2022-05-18 22:41:48 +0000148 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000149 })) {
150 return false;
151 }
Zhaoming Jiang7098d3d2022-04-27 02:27:52 +0000152 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000153
Ben Clayton9a6acc42022-07-27 20:50:40 +0000154 if (!AllocateOverridableConstantIds()) {
155 return false;
156 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000157
158 SetShadows();
159
160 if (!validator_.PipelineStages(entry_points_)) {
161 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000162 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000163
dan sinclair4abf28e2022-08-02 15:55:35 +0000164 if (!validator_.PushConstants(entry_points_)) {
165 return false;
166 }
167
Ben Claytondce63f52022-08-17 18:07:20 +0000168 if (!enabled_extensions_.Contains(ast::Extension::kChromiumDisableUniformityAnalysis)) {
Ben Clayton7f2b8cd2022-05-18 22:41:48 +0000169 if (!AnalyzeUniformity(builder_, dependencies_)) {
Ben Clayton4c70d7f2022-09-26 12:30:05 +0000170 return false;
Ben Clayton7f2b8cd2022-05-18 22:41:48 +0000171 }
James Pricebe656f72022-05-11 22:05:15 +0000172 }
173
dan sinclair41e4d9a2022-05-01 14:40:55 +0000174 bool result = true;
175 for (auto* node : builder_->ASTNodes().Objects()) {
Ben Claytonaf473882022-07-21 23:46:15 +0000176 if (!marked_[node->node_id.value]) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000177 TINT_ICE(Resolver, diagnostics_)
178 << "AST node '" << node->TypeInfo().name << "' was not reached by the resolver\n"
179 << "At: " << node->source << "\n"
180 << "Pointer: " << node;
181 result = false;
182 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000183 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000184
dan sinclair41e4d9a2022-05-01 14:40:55 +0000185 return result;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000186}
187
188sem::Type* Resolver::Type(const ast::Type* ty) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000189 Mark(ty);
190 auto* s = Switch(
191 ty, //
192 [&](const ast::Void*) { return builder_->create<sem::Void>(); },
193 [&](const ast::Bool*) { return builder_->create<sem::Bool>(); },
194 [&](const ast::I32*) { return builder_->create<sem::I32>(); },
195 [&](const ast::U32*) { return builder_->create<sem::U32>(); },
Zhaoming Jiang62bfd312022-05-13 12:01:11 +0000196 [&](const ast::F16* t) -> sem::F16* {
197 // Validate if f16 type is allowed.
Ben Claytondce63f52022-08-17 18:07:20 +0000198 if (!enabled_extensions_.Contains(ast::Extension::kF16)) {
Zhaoming Jiang62bfd312022-05-13 12:01:11 +0000199 AddError("f16 used without 'f16' extension enabled", t->source);
200 return nullptr;
201 }
202 return builder_->create<sem::F16>();
203 },
dan sinclair41e4d9a2022-05-01 14:40:55 +0000204 [&](const ast::F32*) { return builder_->create<sem::F32>(); },
205 [&](const ast::Vector* t) -> sem::Vector* {
206 if (!t->type) {
207 AddError("missing vector element type", t->source.End());
208 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000209 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000210 if (auto* el = Type(t->type)) {
211 if (auto* vector = builder_->create<sem::Vector>(el, t->width)) {
212 if (validator_.Vector(vector, t->source)) {
213 return vector;
214 }
Ben Claytone07c40a2022-03-22 14:04:41 +0000215 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000216 }
217 return nullptr;
218 },
219 [&](const ast::Matrix* t) -> sem::Matrix* {
220 if (!t->type) {
221 AddError("missing matrix element type", t->source.End());
222 return nullptr;
223 }
224 if (auto* el = Type(t->type)) {
225 if (auto* column_type = builder_->create<sem::Vector>(el, t->rows)) {
226 if (auto* matrix = builder_->create<sem::Matrix>(column_type, t->columns)) {
227 if (validator_.Matrix(matrix, t->source)) {
228 return matrix;
229 }
230 }
231 }
232 }
233 return nullptr;
234 },
235 [&](const ast::Array* t) { return Array(t); },
236 [&](const ast::Atomic* t) -> sem::Atomic* {
237 if (auto* el = Type(t->type)) {
238 auto* a = builder_->create<sem::Atomic>(el);
239 if (!validator_.Atomic(t, a)) {
240 return nullptr;
241 }
242 return a;
243 }
244 return nullptr;
245 },
246 [&](const ast::Pointer* t) -> sem::Pointer* {
247 if (auto* el = Type(t->type)) {
248 auto access = t->access;
249 if (access == ast::kUndefined) {
250 access = DefaultAccessForStorageClass(t->storage_class);
251 }
252 return builder_->create<sem::Pointer>(el, t->storage_class, access);
253 }
254 return nullptr;
255 },
256 [&](const ast::Sampler* t) { return builder_->create<sem::Sampler>(t->kind); },
257 [&](const ast::SampledTexture* t) -> sem::SampledTexture* {
258 if (auto* el = Type(t->type)) {
Ben Claytonecb46b32022-06-17 12:52:20 +0000259 auto* sem = builder_->create<sem::SampledTexture>(t->dim, el);
260 if (!validator_.SampledTexture(sem, t->source)) {
261 return nullptr;
262 }
263 return sem;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000264 }
265 return nullptr;
266 },
267 [&](const ast::MultisampledTexture* t) -> sem::MultisampledTexture* {
268 if (auto* el = Type(t->type)) {
Ben Claytonecb46b32022-06-17 12:52:20 +0000269 auto* sem = builder_->create<sem::MultisampledTexture>(t->dim, el);
270 if (!validator_.MultisampledTexture(sem, t->source)) {
271 return nullptr;
272 }
273 return sem;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000274 }
275 return nullptr;
276 },
277 [&](const ast::DepthTexture* t) { return builder_->create<sem::DepthTexture>(t->dim); },
278 [&](const ast::DepthMultisampledTexture* t) {
279 return builder_->create<sem::DepthMultisampledTexture>(t->dim);
280 },
281 [&](const ast::StorageTexture* t) -> sem::StorageTexture* {
282 if (auto* el = Type(t->type)) {
283 if (!validator_.StorageTexture(t)) {
284 return nullptr;
285 }
286 return builder_->create<sem::StorageTexture>(t->dim, t->format, t->access, el);
287 }
288 return nullptr;
289 },
290 [&](const ast::ExternalTexture*) { return builder_->create<sem::ExternalTexture>(); },
291 [&](Default) {
292 auto* resolved = sem_.ResolvedSymbol(ty);
293 return Switch(
294 resolved, //
295 [&](sem::Type* type) { return type; },
296 [&](sem::Variable* var) {
297 auto name = builder_->Symbols().NameFor(var->Declaration()->symbol);
298 AddError("cannot use variable '" + name + "' as type", ty->source);
299 AddNote("'" + name + "' declared here", var->Declaration()->source);
300 return nullptr;
301 },
302 [&](sem::Function* func) {
303 auto name = builder_->Symbols().NameFor(func->Declaration()->symbol);
304 AddError("cannot use function '" + name + "' as type", ty->source);
305 AddNote("'" + name + "' declared here", func->Declaration()->source);
306 return nullptr;
307 },
308 [&](Default) {
309 if (auto* tn = ty->As<ast::TypeName>()) {
310 if (IsBuiltin(tn->name)) {
311 auto name = builder_->Symbols().NameFor(tn->name);
312 AddError("cannot use builtin '" + name + "' as type", ty->source);
313 return nullptr;
314 }
315 }
316 TINT_UNREACHABLE(Resolver, diagnostics_)
317 << "Unhandled resolved type '"
318 << (resolved ? resolved->TypeInfo().name : "<null>")
319 << "' resolved from ast::Type '" << ty->TypeInfo().name << "'";
320 return nullptr;
321 });
322 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000323
dan sinclair41e4d9a2022-05-01 14:40:55 +0000324 if (s) {
325 builder_->Sem().Add(ty, s);
326 }
327 return s;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000328}
329
Ben Claytona05e31b2022-06-18 17:26:40 +0000330sem::Variable* Resolver::Variable(const ast::Variable* v, bool is_global) {
Ben Claytone4e48542022-06-20 19:17:52 +0000331 return Switch(
332 v, //
333 [&](const ast::Var* var) { return Var(var, is_global); },
334 [&](const ast::Let* let) { return Let(let, is_global); },
335 [&](const ast::Override* override) { return Override(override); },
Ben Claytone3834c42022-06-25 23:21:39 +0000336 [&](const ast::Const* const_) { return Const(const_, is_global); },
Ben Claytone4e48542022-06-20 19:17:52 +0000337 [&](Default) {
338 TINT_ICE(Resolver, diagnostics_)
339 << "Resolver::GlobalVariable() called with a unknown variable type: "
340 << v->TypeInfo().name;
341 return nullptr;
342 });
343}
344
345sem::Variable* Resolver::Let(const ast::Let* v, bool is_global) {
Ben Claytonee49b1e2022-06-20 15:30:41 +0000346 const sem::Type* ty = nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000347
dan sinclair41e4d9a2022-05-01 14:40:55 +0000348 // If the variable has a declared type, resolve it.
Ben Claytonee49b1e2022-06-20 15:30:41 +0000349 if (v->type) {
350 ty = Type(v->type);
351 if (!ty) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000352 return nullptr;
353 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000354 }
355
Ben Claytone4e48542022-06-20 19:17:52 +0000356 if (!v->constructor) {
357 AddError("'let' declaration must have an initializer", v->source);
358 return nullptr;
359 }
360
361 auto* rhs = Materialize(Expression(v->constructor), ty);
362 if (!rhs) {
363 return nullptr;
364 }
365
366 // If the variable has no declared type, infer it from the RHS
367 if (!ty) {
368 ty = rhs->Type()->UnwrapRef(); // Implicit load of RHS
369 }
370
Ben Claytonf9031332022-06-26 11:36:10 +0000371 if (rhs && !validator_.VariableInitializer(v, ast::StorageClass::kNone, ty, rhs)) {
Ben Claytone4e48542022-06-20 19:17:52 +0000372 return nullptr;
373 }
374
375 if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, const_cast<sem::Type*>(ty),
376 v->source)) {
377 AddNote("while instantiating 'let' " + builder_->Symbols().NameFor(v->symbol), v->source);
378 return nullptr;
379 }
380
381 sem::Variable* sem = nullptr;
382 if (is_global) {
Ben Claytonaa037ac2022-06-29 19:07:30 +0000383 sem = builder_->create<sem::GlobalVariable>(
Ben Clayton83bd7382022-07-15 23:46:31 +0000384 v, ty, sem::EvaluationStage::kRuntime, ast::StorageClass::kNone,
dan sinclair23cf74c2022-09-07 20:05:54 +0000385 ast::Access::kUndefined, /* constant_value */ nullptr, sem::BindingPoint{},
386 std::nullopt);
Ben Claytone4e48542022-06-20 19:17:52 +0000387 } else {
Ben Clayton83bd7382022-07-15 23:46:31 +0000388 sem = builder_->create<sem::LocalVariable>(v, ty, sem::EvaluationStage::kRuntime,
389 ast::StorageClass::kNone,
Ben Claytone4e48542022-06-20 19:17:52 +0000390 ast::Access::kUndefined, current_statement_,
Ben Claytonaa037ac2022-06-29 19:07:30 +0000391 /* constant_value */ nullptr);
Ben Claytone4e48542022-06-20 19:17:52 +0000392 }
393
394 sem->SetConstructor(rhs);
395 builder_->Sem().Add(v, sem);
396 return sem;
397}
398
399sem::Variable* Resolver::Override(const ast::Override* v) {
400 const sem::Type* ty = nullptr;
401
402 // If the variable has a declared type, resolve it.
403 if (v->type) {
404 ty = Type(v->type);
405 if (!ty) {
406 return nullptr;
407 }
408 }
Ben Claytondcdf66e2022-06-17 12:48:51 +0000409
dan sinclair41e4d9a2022-05-01 14:40:55 +0000410 const sem::Expression* rhs = nullptr;
411
412 // Does the variable have a constructor?
Ben Claytondcdf66e2022-06-17 12:48:51 +0000413 if (v->constructor) {
Ben Claytonee49b1e2022-06-20 15:30:41 +0000414 rhs = Materialize(Expression(v->constructor), ty);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000415 if (!rhs) {
416 return nullptr;
417 }
418
419 // If the variable has no declared type, infer it from the RHS
Ben Claytonee49b1e2022-06-20 15:30:41 +0000420 if (!ty) {
421 ty = rhs->Type()->UnwrapRef(); // Implicit load of RHS
dan sinclair41e4d9a2022-05-01 14:40:55 +0000422 }
Ben Claytonee49b1e2022-06-20 15:30:41 +0000423 } else if (!ty) {
Ben Claytonf9031332022-06-26 11:36:10 +0000424 AddError("override declaration requires a type or initializer", v->source);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000425 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000426 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000427
Ben Claytonf9031332022-06-26 11:36:10 +0000428 if (rhs && !validator_.VariableInitializer(v, ast::StorageClass::kNone, ty, rhs)) {
Ben Claytonee49b1e2022-06-20 15:30:41 +0000429 return nullptr;
430 }
431
432 if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, const_cast<sem::Type*>(ty),
433 v->source)) {
Ben Claytone4e48542022-06-20 19:17:52 +0000434 AddNote("while instantiating 'override' " + builder_->Symbols().NameFor(v->symbol),
Ben Claytonee49b1e2022-06-20 15:30:41 +0000435 v->source);
436 return nullptr;
437 }
438
Ben Claytonaa037ac2022-06-29 19:07:30 +0000439 auto* sem = builder_->create<sem::GlobalVariable>(
Ben Clayton83bd7382022-07-15 23:46:31 +0000440 v, ty, sem::EvaluationStage::kOverride, ast::StorageClass::kNone, ast::Access::kUndefined,
dan sinclair23cf74c2022-09-07 20:05:54 +0000441 /* constant_value */ nullptr, sem::BindingPoint{}, std::nullopt);
dan sinclair5361d9e2022-08-31 13:39:48 +0000442 sem->SetConstructor(rhs);
Ben Claytonee49b1e2022-06-20 15:30:41 +0000443
dan sinclair5361d9e2022-08-31 13:39:48 +0000444 if (auto* id_attr = ast::GetAttribute<ast::IdAttribute>(v->attributes)) {
dan sinclairc4e076f2022-09-08 01:04:34 +0000445 auto* materialize = Materialize(Expression(id_attr->expr));
dan sinclair5361d9e2022-08-31 13:39:48 +0000446 if (!materialize) {
447 return nullptr;
448 }
449 auto* c = materialize->ConstantValue();
450 if (!c) {
451 // TODO(crbug.com/tint/1633): Handle invalid materialization when expressions
452 // are supported.
453 return nullptr;
454 }
455
456 auto value = c->As<uint32_t>();
457 if (value > std::numeric_limits<decltype(OverrideId::value)>::max()) {
458 AddError("override IDs must be between 0 and " +
459 std::to_string(std::numeric_limits<decltype(OverrideId::value)>::max()),
460 id_attr->source);
461 return nullptr;
462 }
463
464 auto o = OverrideId{static_cast<decltype(OverrideId::value)>(value)};
465 sem->SetOverrideId(o);
466
467 // Track the constant IDs that are specified in the shader.
468 override_ids_.emplace(o, sem);
Ben Claytonee49b1e2022-06-20 15:30:41 +0000469 }
470
Ben Claytonee49b1e2022-06-20 15:30:41 +0000471 builder_->Sem().Add(v, sem);
472 return sem;
473}
474
Ben Claytone3834c42022-06-25 23:21:39 +0000475sem::Variable* Resolver::Const(const ast::Const* c, bool is_global) {
476 const sem::Type* ty = nullptr;
477
478 // If the variable has a declared type, resolve it.
479 if (c->type) {
480 ty = Type(c->type);
481 if (!ty) {
482 return nullptr;
483 }
484 }
485
486 if (!c->constructor) {
487 AddError("'const' declaration must have an initializer", c->source);
488 return nullptr;
489 }
490
491 const auto* rhs = Expression(c->constructor);
492 if (!rhs) {
493 return nullptr;
494 }
495
496 if (ty) {
497 // If an explicit type was specified, materialize to that type
498 rhs = Materialize(rhs, ty);
Ben Clayton89bdea02022-07-07 16:56:21 +0000499 if (!rhs) {
500 return nullptr;
501 }
Ben Claytone3834c42022-06-25 23:21:39 +0000502 } else {
503 // If no type was specified, infer it from the RHS
504 ty = rhs->Type();
505 }
506
507 const auto value = rhs->ConstantValue();
508 if (!value) {
509 AddError("'const' initializer must be constant expression", c->constructor->source);
510 return nullptr;
511 }
512
Ben Claytonf9031332022-06-26 11:36:10 +0000513 if (!validator_.VariableInitializer(c, ast::StorageClass::kNone, ty, rhs)) {
Ben Claytone3834c42022-06-25 23:21:39 +0000514 return nullptr;
515 }
516
517 if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, const_cast<sem::Type*>(ty),
518 c->source)) {
519 AddNote("while instantiating 'const' " + builder_->Symbols().NameFor(c->symbol), c->source);
520 return nullptr;
521 }
522
523 auto* sem = is_global ? static_cast<sem::Variable*>(builder_->create<sem::GlobalVariable>(
Ben Clayton83bd7382022-07-15 23:46:31 +0000524 c, ty, sem::EvaluationStage::kConstant, ast::StorageClass::kNone,
dan sinclair23cf74c2022-09-07 20:05:54 +0000525 ast::Access::kUndefined, value, sem::BindingPoint{}, std::nullopt))
Ben Claytone3834c42022-06-25 23:21:39 +0000526 : static_cast<sem::Variable*>(builder_->create<sem::LocalVariable>(
Ben Clayton83bd7382022-07-15 23:46:31 +0000527 c, ty, sem::EvaluationStage::kConstant, ast::StorageClass::kNone,
528 ast::Access::kUndefined, current_statement_, value));
Ben Claytone3834c42022-06-25 23:21:39 +0000529
530 sem->SetConstructor(rhs);
531 builder_->Sem().Add(c, sem);
532 return sem;
533}
534
Ben Claytonee49b1e2022-06-20 15:30:41 +0000535sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) {
536 const sem::Type* storage_ty = nullptr;
537
538 // If the variable has a declared type, resolve it.
539 if (auto* ty = var->type) {
540 storage_ty = Type(ty);
541 if (!storage_ty) {
542 return nullptr;
543 }
544 }
545
546 const sem::Expression* rhs = nullptr;
547
548 // Does the variable have a constructor?
549 if (var->constructor) {
550 rhs = Materialize(Expression(var->constructor), storage_ty);
551 if (!rhs) {
552 return nullptr;
553 }
554 // If the variable has no declared type, infer it from the RHS
555 if (!storage_ty) {
556 storage_ty = rhs->Type()->UnwrapRef(); // Implicit load of RHS
557 }
558 }
559
dan sinclair41e4d9a2022-05-01 14:40:55 +0000560 if (!storage_ty) {
Ben Claytonf9031332022-06-26 11:36:10 +0000561 AddError("var declaration requires a type or initializer", var->source);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000562 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000563 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000564
Ben Claytonee49b1e2022-06-20 15:30:41 +0000565 auto storage_class = var->declared_storage_class;
566 if (storage_class == ast::StorageClass::kNone) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000567 // No declared storage class. Infer from usage / type.
Ben Claytondcdf66e2022-06-17 12:48:51 +0000568 if (!is_global) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000569 storage_class = ast::StorageClass::kFunction;
570 } else if (storage_ty->UnwrapRef()->is_handle()) {
571 // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
572 // If the store type is a texture type or a sampler type, then the
573 // variable declaration must not have a storage class attribute. The
574 // storage class will always be handle.
575 storage_class = ast::StorageClass::kHandle;
576 }
577 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000578
Ben Claytonee49b1e2022-06-20 15:30:41 +0000579 if (!is_global && storage_class != ast::StorageClass::kFunction &&
580 validator_.IsValidationEnabled(var->attributes,
dan sinclair41e4d9a2022-05-01 14:40:55 +0000581 ast::DisabledValidation::kIgnoreStorageClass)) {
Ben Claytonee49b1e2022-06-20 15:30:41 +0000582 AddError("function-scope 'var' declaration must use 'function' storage class", var->source);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000583 return nullptr;
584 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000585
Ben Claytonee49b1e2022-06-20 15:30:41 +0000586 auto access = var->declared_access;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000587 if (access == ast::Access::kUndefined) {
588 access = DefaultAccessForStorageClass(storage_class);
589 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000590
Ben Claytonf9031332022-06-26 11:36:10 +0000591 if (rhs && !validator_.VariableInitializer(var, storage_class, storage_ty, rhs)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000592 return nullptr;
593 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000594
Ben Claytonee49b1e2022-06-20 15:30:41 +0000595 auto* var_ty = builder_->create<sem::Reference>(storage_ty, storage_class, access);
596
597 if (!ApplyStorageClassUsageToType(storage_class, var_ty, var->source)) {
598 AddNote("while instantiating 'var' " + builder_->Symbols().NameFor(var->symbol),
599 var->source);
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000600 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000601 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000602
Ben Claytonee49b1e2022-06-20 15:30:41 +0000603 sem::Variable* sem = nullptr;
Ben Claytondcdf66e2022-06-17 12:48:51 +0000604 if (is_global) {
605 sem::BindingPoint binding_point;
dan sinclairacdf6e12022-08-24 15:47:25 +0000606 if (var->HasBindingPoint()) {
607 uint32_t binding = 0;
608 {
609 auto* attr = ast::GetAttribute<ast::BindingAttribute>(var->attributes);
dan sinclairc4e076f2022-09-08 01:04:34 +0000610 auto* materialize = Materialize(Expression(attr->expr));
dan sinclairf9b831c2022-08-29 21:13:00 +0000611 if (!materialize) {
612 return nullptr;
613 }
614 auto* c = materialize->ConstantValue();
615 if (!c) {
616 // TODO(crbug.com/tint/1633): Add error message about invalid materialization
617 // when binding can be an expression.
618 return nullptr;
619 }
620 binding = c->As<uint32_t>();
dan sinclairacdf6e12022-08-24 15:47:25 +0000621 }
622
623 uint32_t group = 0;
624 {
625 auto* attr = ast::GetAttribute<ast::GroupAttribute>(var->attributes);
dan sinclairc4e076f2022-09-08 01:04:34 +0000626 auto* materialize = Materialize(Expression(attr->expr));
dan sinclairbe4c9f42022-08-29 21:22:31 +0000627 if (!materialize) {
628 return nullptr;
629 }
630 auto* c = materialize->ConstantValue();
631 if (!c) {
632 // TODO(crbug.com/tint/1633): Add error message about invalid materialization
633 // when binding can be an expression.
634 return nullptr;
635 }
636 group = c->As<uint32_t>();
dan sinclairacdf6e12022-08-24 15:47:25 +0000637 }
638 binding_point = {group, binding};
Ben Claytondcdf66e2022-06-17 12:48:51 +0000639 }
dan sinclair23cf74c2022-09-07 20:05:54 +0000640
641 std::optional<uint32_t> location;
642 if (auto* attr = ast::GetAttribute<ast::LocationAttribute>(var->attributes)) {
dan sinclairc4e076f2022-09-08 01:04:34 +0000643 auto* materialize = Materialize(Expression(attr->expr));
dan sinclairf9eeed62022-09-07 22:25:24 +0000644 if (!materialize) {
645 return nullptr;
646 }
647 auto* c = materialize->ConstantValue();
648 if (!c) {
649 // TODO(crbug.com/tint/1633): Add error message about invalid materialization
650 // when location can be an expression.
651 return nullptr;
652 }
653 location = c->As<uint32_t>();
dan sinclair23cf74c2022-09-07 20:05:54 +0000654 }
655
656 sem = builder_->create<sem::GlobalVariable>(
657 var, var_ty, sem::EvaluationStage::kRuntime, storage_class, access,
658 /* constant_value */ nullptr, binding_point, location);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000659
Ben Claytonee49b1e2022-06-20 15:30:41 +0000660 } else {
Ben Clayton83bd7382022-07-15 23:46:31 +0000661 sem = builder_->create<sem::LocalVariable>(var, var_ty, sem::EvaluationStage::kRuntime,
662 storage_class, access, current_statement_,
663 /* constant_value */ nullptr);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000664 }
665
Ben Claytonee49b1e2022-06-20 15:30:41 +0000666 sem->SetConstructor(rhs);
667 builder_->Sem().Add(var, sem);
668 return sem;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000669}
670
Ben Claytona05e31b2022-06-18 17:26:40 +0000671sem::Parameter* Resolver::Parameter(const ast::Parameter* param, uint32_t index) {
672 auto add_note = [&] {
673 AddNote("while instantiating parameter " + builder_->Symbols().NameFor(param->symbol),
674 param->source);
675 };
676
677 for (auto* attr : param->attributes) {
678 Mark(attr);
679 }
680 if (!validator_.NoDuplicateAttributes(param->attributes)) {
681 return nullptr;
682 }
683
684 sem::Type* ty = Type(param->type);
685 if (!ty) {
686 return nullptr;
687 }
688
689 if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, ty, param->source)) {
690 add_note();
691 return nullptr;
692 }
693
694 if (auto* ptr = ty->As<sem::Pointer>()) {
695 // For MSL, we push module-scope variables into the entry point as pointer
696 // parameters, so we also need to handle their store type.
697 if (!ApplyStorageClassUsageToType(
698 ptr->StorageClass(), const_cast<sem::Type*>(ptr->StoreType()), param->source)) {
699 add_note();
700 return nullptr;
701 }
702 }
703
dan sinclairacdf6e12022-08-24 15:47:25 +0000704 sem::BindingPoint binding_point;
705 if (param->HasBindingPoint()) {
706 {
707 auto* attr = ast::GetAttribute<ast::BindingAttribute>(param->attributes);
dan sinclairc4e076f2022-09-08 01:04:34 +0000708 auto* materialize = Materialize(Expression(attr->expr));
dan sinclairf9b831c2022-08-29 21:13:00 +0000709 if (!materialize) {
710 return nullptr;
711 }
712 auto* c = materialize->ConstantValue();
713 if (!c) {
714 // TODO(crbug.com/tint/1633): Add error message about invalid materialization when
715 // binding can be an expression.
716 return nullptr;
717 }
718 binding_point.binding = c->As<uint32_t>();
dan sinclairacdf6e12022-08-24 15:47:25 +0000719 }
720 {
721 auto* attr = ast::GetAttribute<ast::GroupAttribute>(param->attributes);
dan sinclairc4e076f2022-09-08 01:04:34 +0000722 auto* materialize = Materialize(Expression(attr->expr));
dan sinclairbe4c9f42022-08-29 21:22:31 +0000723 if (!materialize) {
724 return nullptr;
725 }
726 auto* c = materialize->ConstantValue();
727 if (!c) {
728 // TODO(crbug.com/tint/1633): Add error message about invalid materialization when
729 // binding can be an expression.
730 return nullptr;
731 }
732 binding_point.group = c->As<uint32_t>();
dan sinclairacdf6e12022-08-24 15:47:25 +0000733 }
734 }
735
dan sinclair72340d02022-09-06 16:05:54 +0000736 std::optional<uint32_t> location;
737 if (auto* l = ast::GetAttribute<ast::LocationAttribute>(param->attributes)) {
dan sinclairc4e076f2022-09-08 01:04:34 +0000738 auto* materialize = Materialize(Expression(l->expr));
dan sinclairf9eeed62022-09-07 22:25:24 +0000739 if (!materialize) {
740 return nullptr;
741 }
742 auto* c = materialize->ConstantValue();
743 if (!c) {
744 // TODO(crbug.com/tint/1633): Add error message about invalid materialization when
745 // location can be an expression.
746 return nullptr;
747 }
748 location = c->As<uint32_t>();
dan sinclair72340d02022-09-06 16:05:54 +0000749 }
750
751 auto* sem = builder_->create<sem::Parameter>(
752 param, index, ty, ast::StorageClass::kNone, ast::Access::kUndefined,
753 sem::ParameterUsage::kNone, binding_point, location);
Ben Claytona05e31b2022-06-18 17:26:40 +0000754 builder_->Sem().Add(param, sem);
755 return sem;
756}
757
dan sinclair41e4d9a2022-05-01 14:40:55 +0000758ast::Access Resolver::DefaultAccessForStorageClass(ast::StorageClass storage_class) {
759 // https://gpuweb.github.io/gpuweb/wgsl/#storage-class
760 switch (storage_class) {
761 case ast::StorageClass::kStorage:
762 case ast::StorageClass::kUniform:
763 case ast::StorageClass::kHandle:
764 return ast::Access::kRead;
765 default:
766 break;
767 }
768 return ast::Access::kReadWrite;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000769}
770
Ben Clayton9a6acc42022-07-27 20:50:40 +0000771bool Resolver::AllocateOverridableConstantIds() {
772 constexpr size_t kLimit = std::numeric_limits<decltype(OverrideId::value)>::max();
dan sinclair41e4d9a2022-05-01 14:40:55 +0000773 // The next pipeline constant ID to try to allocate.
Ben Clayton9a6acc42022-07-27 20:50:40 +0000774 OverrideId next_id;
775 bool ids_exhausted = false;
776
777 auto increment_next_id = [&] {
778 if (next_id.value == kLimit) {
779 ids_exhausted = true;
780 } else {
781 next_id.value = next_id.value + 1;
782 }
783 };
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000784
dan sinclair41e4d9a2022-05-01 14:40:55 +0000785 // Allocate constant IDs in global declaration order, so that they are
786 // deterministic.
787 // TODO(crbug.com/tint/1192): If a transform changes the order or removes an
788 // unused constant, the allocation may change on the next Resolver pass.
789 for (auto* decl : builder_->AST().GlobalDeclarations()) {
Ben Claytondcdf66e2022-06-17 12:48:51 +0000790 auto* override = decl->As<ast::Override>();
791 if (!override) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000792 continue;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000793 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000794
Ben Clayton9a6acc42022-07-27 20:50:40 +0000795 OverrideId id;
dan sinclair5361d9e2022-08-31 13:39:48 +0000796 if (ast::HasAttribute<ast::IdAttribute>(override->attributes)) {
797 id = builder_->Sem().Get<sem::GlobalVariable>(override)->OverrideId();
dan sinclair41e4d9a2022-05-01 14:40:55 +0000798 } else {
799 // No ID was specified, so allocate the next available ID.
Ben Clayton9a6acc42022-07-27 20:50:40 +0000800 while (!ids_exhausted && override_ids_.count(next_id)) {
801 increment_next_id();
dan sinclair41e4d9a2022-05-01 14:40:55 +0000802 }
Ben Clayton9a6acc42022-07-27 20:50:40 +0000803 if (ids_exhausted) {
804 AddError(
805 "number of 'override' variables exceeded limit of " + std::to_string(kLimit),
806 decl->source);
807 return false;
808 }
809 id = next_id;
810 increment_next_id();
dan sinclair41e4d9a2022-05-01 14:40:55 +0000811 }
812
Ben Claytondcdf66e2022-06-17 12:48:51 +0000813 auto* sem = sem_.Get<sem::GlobalVariable>(override);
Ben Clayton9a6acc42022-07-27 20:50:40 +0000814 const_cast<sem::GlobalVariable*>(sem)->SetOverrideId(id);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000815 }
Ben Clayton9a6acc42022-07-27 20:50:40 +0000816 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000817}
818
819void Resolver::SetShadows() {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000820 for (auto it : dependencies_.shadows) {
821 Switch(
822 sem_.Get(it.first), //
823 [&](sem::LocalVariable* local) { local->SetShadows(sem_.Get(it.second)); },
824 [&](sem::Parameter* param) { param->SetShadows(sem_.Get(it.second)); });
825 }
dan sinclaird2093792022-04-07 17:45:45 +0000826}
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000827
Ben Claytondcdf66e2022-06-17 12:48:51 +0000828sem::GlobalVariable* Resolver::GlobalVariable(const ast::Variable* v) {
Ben Claytone4e48542022-06-20 19:17:52 +0000829 auto* sem = As<sem::GlobalVariable>(Variable(v, /* is_global */ true));
dan sinclair41e4d9a2022-05-01 14:40:55 +0000830 if (!sem) {
831 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000832 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000833
Ben Claytondcdf66e2022-06-17 12:48:51 +0000834 for (auto* attr : v->attributes) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000835 Mark(attr);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000836 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000837
Ben Claytondcdf66e2022-06-17 12:48:51 +0000838 if (!validator_.NoDuplicateAttributes(v->attributes)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000839 return nullptr;
840 }
841
Ben Clayton9a6acc42022-07-27 20:50:40 +0000842 if (!validator_.GlobalVariable(sem, override_ids_, atomic_composite_info_)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000843 return nullptr;
844 }
845
846 // TODO(bclayton): Call this at the end of resolve on all uniform and storage
847 // referenced structs
dan sinclair4abf28e2022-08-02 15:55:35 +0000848 if (!validator_.StorageClassLayout(sem, enabled_extensions_, valid_type_storage_layouts_)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000849 return nullptr;
850 }
851
Ben Claytonee49b1e2022-06-20 15:30:41 +0000852 return sem;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000853}
854
Ben Clayton02791e92022-08-02 23:28:28 +0000855sem::Statement* Resolver::StaticAssert(const ast::StaticAssert* assertion) {
856 auto* expr = Expression(assertion->condition);
857 if (!expr) {
858 return nullptr;
859 }
860 auto* cond = expr->ConstantValue();
861 if (!cond) {
862 AddError("static assertion condition must be a constant expression",
863 assertion->condition->source);
864 return nullptr;
865 }
866 if (auto* ty = cond->Type(); !ty->Is<sem::Bool>()) {
867 AddError(
868 "static assertion condition must be a bool, got '" + builder_->FriendlyName(ty) + "'",
869 assertion->condition->source);
870 return nullptr;
871 }
872 if (!cond->As<bool>()) {
873 AddError("static assertion failed", assertion->source);
874 return nullptr;
875 }
876 auto* sem =
877 builder_->create<sem::Statement>(assertion, current_compound_statement_, current_function_);
878 builder_->Sem().Add(assertion, sem);
879 return sem;
880}
881
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000882sem::Function* Resolver::Function(const ast::Function* decl) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000883 uint32_t parameter_index = 0;
884 std::unordered_map<Symbol, Source> parameter_names;
Ben Clayton958a4642022-07-26 07:55:24 +0000885 utils::Vector<sem::Parameter*, 8> parameters;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000886
dan sinclair41e4d9a2022-05-01 14:40:55 +0000887 // Resolve all the parameters
888 for (auto* param : decl->params) {
889 Mark(param);
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000890
dan sinclair41e4d9a2022-05-01 14:40:55 +0000891 { // Check the parameter name is unique for the function
892 auto emplaced = parameter_names.emplace(param->symbol, param->source);
893 if (!emplaced.second) {
894 auto name = builder_->Symbols().NameFor(param->symbol);
895 AddError("redefinition of parameter '" + name + "'", param->source);
896 AddNote("previous definition is here", emplaced.first->second);
897 return nullptr;
898 }
899 }
900
Ben Claytona05e31b2022-06-18 17:26:40 +0000901 auto* p = Parameter(param, parameter_index++);
Ben Claytondcdf66e2022-06-17 12:48:51 +0000902 if (!p) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000903 return nullptr;
904 }
905
Ben Claytona05e31b2022-06-18 17:26:40 +0000906 if (!validator_.Parameter(decl, p)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000907 return nullptr;
908 }
909
Ben Clayton958a4642022-07-26 07:55:24 +0000910 parameters.Push(p);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000911
Ben Claytondcdf66e2022-06-17 12:48:51 +0000912 auto* p_ty = const_cast<sem::Type*>(p->Type());
913 if (auto* str = p_ty->As<sem::Struct>()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000914 switch (decl->PipelineStage()) {
915 case ast::PipelineStage::kVertex:
916 str->AddUsage(sem::PipelineStageUsage::kVertexInput);
917 break;
918 case ast::PipelineStage::kFragment:
919 str->AddUsage(sem::PipelineStageUsage::kFragmentInput);
920 break;
921 case ast::PipelineStage::kCompute:
922 str->AddUsage(sem::PipelineStageUsage::kComputeInput);
923 break;
924 case ast::PipelineStage::kNone:
925 break;
926 }
927 }
928 }
929
930 // Resolve the return type
931 sem::Type* return_type = nullptr;
932 if (auto* ty = decl->return_type) {
933 return_type = Type(ty);
934 if (!return_type) {
935 return nullptr;
936 }
937 } else {
938 return_type = builder_->create<sem::Void>();
939 }
940
dan sinclair766a4582022-09-06 16:00:13 +0000941 // Determine if the return type has a location
942 std::optional<uint32_t> return_location;
943 for (auto* attr : decl->return_type_attributes) {
944 Mark(attr);
945
946 if (auto* a = attr->As<ast::LocationAttribute>()) {
dan sinclairc4e076f2022-09-08 01:04:34 +0000947 auto* materialize = Materialize(Expression(a->expr));
dan sinclairf9eeed62022-09-07 22:25:24 +0000948 if (!materialize) {
949 return nullptr;
950 }
951 auto* c = materialize->ConstantValue();
952 if (!c) {
953 // TODO(crbug.com/tint/1633): Add error message about invalid materialization when
954 // location can be an expression.
955 return nullptr;
956 }
957 return_location = c->As<uint32_t>();
dan sinclair766a4582022-09-06 16:00:13 +0000958 }
959 }
960 if (!validator_.NoDuplicateAttributes(decl->attributes)) {
961 return nullptr;
962 }
963
dan sinclair41e4d9a2022-05-01 14:40:55 +0000964 if (auto* str = return_type->As<sem::Struct>()) {
965 if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, str, decl->source)) {
966 AddNote(
967 "while instantiating return type for " + builder_->Symbols().NameFor(decl->symbol),
968 decl->source);
969 return nullptr;
970 }
971
972 switch (decl->PipelineStage()) {
973 case ast::PipelineStage::kVertex:
974 str->AddUsage(sem::PipelineStageUsage::kVertexOutput);
975 break;
976 case ast::PipelineStage::kFragment:
977 str->AddUsage(sem::PipelineStageUsage::kFragmentOutput);
978 break;
979 case ast::PipelineStage::kCompute:
980 str->AddUsage(sem::PipelineStageUsage::kComputeOutput);
981 break;
982 case ast::PipelineStage::kNone:
983 break;
984 }
985 }
986
dan sinclair766a4582022-09-06 16:00:13 +0000987 auto* func =
988 builder_->create<sem::Function>(decl, return_type, return_location, std::move(parameters));
dan sinclair41e4d9a2022-05-01 14:40:55 +0000989 builder_->Sem().Add(decl, func);
990
991 TINT_SCOPED_ASSIGNMENT(current_function_, func);
992
993 if (!WorkgroupSize(decl)) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000994 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000995 }
996
dan sinclair41e4d9a2022-05-01 14:40:55 +0000997 if (decl->IsEntryPoint()) {
998 entry_points_.emplace_back(func);
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000999 }
1000
dan sinclair41e4d9a2022-05-01 14:40:55 +00001001 if (decl->body) {
1002 Mark(decl->body);
1003 if (current_compound_statement_) {
1004 TINT_ICE(Resolver, diagnostics_)
1005 << "Resolver::Function() called with a current compound statement";
1006 return nullptr;
1007 }
1008 auto* body = StatementScope(decl->body, builder_->create<sem::FunctionBlockStatement>(func),
1009 [&] { return Statements(decl->body->statements); });
1010 if (!body) {
1011 return nullptr;
1012 }
1013 func->Behaviors() = body->Behaviors();
1014 if (func->Behaviors().Contains(sem::Behavior::kReturn)) {
1015 // https://www.w3.org/TR/WGSL/#behaviors-rules
1016 // We assign a behavior to each function: it is its body’s behavior
1017 // (treating the body as a regular statement), with any "Return" replaced
1018 // by "Next".
1019 func->Behaviors().Remove(sem::Behavior::kReturn);
1020 func->Behaviors().Add(sem::Behavior::kNext);
1021 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001022 }
1023
dan sinclair41e4d9a2022-05-01 14:40:55 +00001024 for (auto* attr : decl->attributes) {
1025 Mark(attr);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001026 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001027
dan sinclair41e4d9a2022-05-01 14:40:55 +00001028 if (!validator_.NoDuplicateAttributes(decl->return_type_attributes)) {
1029 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001030 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001031
1032 auto stage = current_function_ ? current_function_->Declaration()->PipelineStage()
1033 : ast::PipelineStage::kNone;
1034 if (!validator_.Function(func, stage)) {
1035 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001036 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001037
1038 // If this is an entry point, mark all transitively called functions as being
1039 // used by this entry point.
1040 if (decl->IsEntryPoint()) {
1041 for (auto* f : func->TransitivelyCalledFunctions()) {
1042 const_cast<sem::Function*>(f)->AddAncestorEntryPoint(func);
1043 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001044 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001045
dan sinclair41e4d9a2022-05-01 14:40:55 +00001046 return func;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001047}
1048
1049bool Resolver::WorkgroupSize(const ast::Function* func) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001050 // Set work-group size defaults.
1051 sem::WorkgroupSize ws;
dan sinclair3a2a2792022-06-29 14:38:15 +00001052 for (size_t i = 0; i < 3; i++) {
Ben Clayton490d9882022-09-21 21:05:45 +00001053 ws[i] = 1;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001054 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001055
dan sinclair41e4d9a2022-05-01 14:40:55 +00001056 auto* attr = ast::GetAttribute<ast::WorkgroupAttribute>(func->attributes);
1057 if (!attr) {
1058 return true;
1059 }
1060
1061 auto values = attr->Values();
Ben Clayton8e0368e2022-07-29 14:15:01 +00001062 utils::Vector<const sem::Expression*, 3> args;
1063 utils::Vector<const sem::Type*, 3> arg_tys;
Ben Claytonb8ac93392022-05-28 10:34:06 +00001064
Ben Clayton511529f2022-06-26 09:05:00 +00001065 constexpr const char* kErrBadExpr =
Ben Clayton490d9882022-09-21 21:05:45 +00001066 "workgroup_size argument must be a constant or override expression of type "
Ben Clayton511529f2022-06-26 09:05:00 +00001067 "abstract-integer, i32 or u32";
Ben Claytonb8ac93392022-05-28 10:34:06 +00001068
dan sinclair3a2a2792022-06-29 14:38:15 +00001069 for (size_t i = 0; i < 3; i++) {
Ben Claytonb8ac93392022-05-28 10:34:06 +00001070 // Each argument to this attribute can either be a literal, an identifier for a module-scope
dan sinclair7517e212022-08-24 21:31:45 +00001071 // constants, a constant expression, or nullptr if not specified.
Ben Claytonb8ac93392022-05-28 10:34:06 +00001072 auto* value = values[i];
1073 if (!value) {
1074 break;
1075 }
1076 const auto* expr = Expression(value);
dan sinclair41e4d9a2022-05-01 14:40:55 +00001077 if (!expr) {
Ben Claytonb8ac93392022-05-28 10:34:06 +00001078 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001079 }
Ben Claytonb8ac93392022-05-28 10:34:06 +00001080 auto* ty = expr->Type();
1081 if (!ty->IsAnyOf<sem::I32, sem::U32, sem::AbstractInt>()) {
Ben Clayton511529f2022-06-26 09:05:00 +00001082 AddError(kErrBadExpr, value->source);
dan sinclair41e4d9a2022-05-01 14:40:55 +00001083 return false;
1084 }
1085
Ben Clayton490d9882022-09-21 21:05:45 +00001086 if (expr->Stage() != sem::EvaluationStage::kConstant &&
1087 expr->Stage() != sem::EvaluationStage::kOverride) {
1088 AddError(kErrBadExpr, value->source);
1089 return false;
1090 }
1091
Ben Clayton8e0368e2022-07-29 14:15:01 +00001092 args.Push(expr);
1093 arg_tys.Push(ty);
Ben Claytonb8ac93392022-05-28 10:34:06 +00001094 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001095
Ben Clayton8e0368e2022-07-29 14:15:01 +00001096 auto* common_ty = sem::Type::Common(arg_tys);
Ben Claytonb8ac93392022-05-28 10:34:06 +00001097 if (!common_ty) {
1098 AddError("workgroup_size arguments must be of the same type, either i32 or u32",
1099 attr->source);
1100 return false;
1101 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001102
Ben Claytonb8ac93392022-05-28 10:34:06 +00001103 // If all arguments are abstract-integers, then materialize to i32.
1104 if (common_ty->Is<sem::AbstractInt>()) {
1105 common_ty = builder_->create<sem::I32>();
1106 }
1107
Ben Clayton8e0368e2022-07-29 14:15:01 +00001108 for (size_t i = 0; i < args.Length(); i++) {
Ben Claytonb8ac93392022-05-28 10:34:06 +00001109 auto* materialized = Materialize(args[i], common_ty);
1110 if (!materialized) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001111 return false;
1112 }
Ben Clayton490d9882022-09-21 21:05:45 +00001113 if (auto* value = materialized->ConstantValue()) {
1114 if (value->As<AInt>() < 1) {
1115 AddError("workgroup_size argument must be at least 1", values[i]->source);
dan sinclair41e4d9a2022-05-01 14:40:55 +00001116 return false;
1117 }
Ben Clayton490d9882022-09-21 21:05:45 +00001118 ws[i] = value->As<uint32_t>();
dan sinclair41e4d9a2022-05-01 14:40:55 +00001119 } else {
Ben Clayton490d9882022-09-21 21:05:45 +00001120 ws[i] = std::nullopt;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001121 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001122 }
1123
1124 current_function_->SetWorkgroupSize(std::move(ws));
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001125 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001126}
1127
Ben Clayton783b1692022-08-02 17:03:35 +00001128bool Resolver::Statements(utils::VectorRef<const ast::Statement*> stmts) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001129 sem::Behaviors behaviors{sem::Behavior::kNext};
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001130
dan sinclair41e4d9a2022-05-01 14:40:55 +00001131 bool reachable = true;
1132 for (auto* stmt : stmts) {
1133 Mark(stmt);
1134 auto* sem = Statement(stmt);
1135 if (!sem) {
1136 return false;
1137 }
1138 // s1 s2:(B1∖{Next}) ∪ B2
1139 sem->SetIsReachable(reachable);
1140 if (reachable) {
1141 behaviors = (behaviors - sem::Behavior::kNext) + sem->Behaviors();
1142 }
1143 reachable = reachable && sem->Behaviors().Contains(sem::Behavior::kNext);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001144 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001145
1146 current_statement_->Behaviors() = behaviors;
1147
1148 if (!validator_.Statements(stmts)) {
1149 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001150 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001151
dan sinclair41e4d9a2022-05-01 14:40:55 +00001152 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001153}
1154
1155sem::Statement* Resolver::Statement(const ast::Statement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001156 return Switch(
1157 stmt,
1158 // Compound statements. These create their own sem::CompoundStatement
1159 // bindings.
1160 [&](const ast::BlockStatement* b) { return BlockStatement(b); },
1161 [&](const ast::ForLoopStatement* l) { return ForLoopStatement(l); },
1162 [&](const ast::LoopStatement* l) { return LoopStatement(l); },
dan sinclair49d1a2d2022-06-16 12:01:27 +00001163 [&](const ast::WhileStatement* w) { return WhileStatement(w); },
dan sinclair41e4d9a2022-05-01 14:40:55 +00001164 [&](const ast::IfStatement* i) { return IfStatement(i); },
1165 [&](const ast::SwitchStatement* s) { return SwitchStatement(s); },
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001166
dan sinclair41e4d9a2022-05-01 14:40:55 +00001167 // Non-Compound statements
1168 [&](const ast::AssignmentStatement* a) { return AssignmentStatement(a); },
1169 [&](const ast::BreakStatement* b) { return BreakStatement(b); },
1170 [&](const ast::CallStatement* c) { return CallStatement(c); },
1171 [&](const ast::CompoundAssignmentStatement* c) { return CompoundAssignmentStatement(c); },
1172 [&](const ast::ContinueStatement* c) { return ContinueStatement(c); },
1173 [&](const ast::DiscardStatement* d) { return DiscardStatement(d); },
1174 [&](const ast::FallthroughStatement* f) { return FallthroughStatement(f); },
1175 [&](const ast::IncrementDecrementStatement* i) { return IncrementDecrementStatement(i); },
1176 [&](const ast::ReturnStatement* r) { return ReturnStatement(r); },
1177 [&](const ast::VariableDeclStatement* v) { return VariableDeclStatement(v); },
Ben Clayton02791e92022-08-02 23:28:28 +00001178 [&](const ast::StaticAssert* sa) { return StaticAssert(sa); },
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001179
dan sinclair41e4d9a2022-05-01 14:40:55 +00001180 // Error cases
1181 [&](const ast::CaseStatement*) {
1182 AddError("case statement can only be used inside a switch statement", stmt->source);
1183 return nullptr;
1184 },
1185 [&](Default) {
1186 AddError("unknown statement type: " + std::string(stmt->TypeInfo().name), stmt->source);
1187 return nullptr;
1188 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001189}
1190
1191sem::CaseStatement* Resolver::CaseStatement(const ast::CaseStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001192 auto* sem =
1193 builder_->create<sem::CaseStatement>(stmt, current_compound_statement_, current_function_);
1194 return StatementScope(stmt, sem, [&] {
Ben Clayton783b1692022-08-02 17:03:35 +00001195 sem->Selectors().reserve(stmt->selectors.Length());
dan sinclair41e4d9a2022-05-01 14:40:55 +00001196 for (auto* sel : stmt->selectors) {
Ben Clayton43581f12022-05-20 12:28:00 +00001197 auto* expr = Expression(sel);
1198 if (!expr) {
Ben Clayton8822e292022-05-04 22:18:49 +00001199 return false;
1200 }
Ben Clayton43581f12022-05-20 12:28:00 +00001201 sem->Selectors().emplace_back(expr);
dan sinclair41e4d9a2022-05-01 14:40:55 +00001202 }
1203 Mark(stmt->body);
1204 auto* body = BlockStatement(stmt->body);
1205 if (!body) {
1206 return false;
1207 }
1208 sem->SetBlock(body);
1209 sem->Behaviors() = body->Behaviors();
1210 return true;
1211 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001212}
1213
1214sem::IfStatement* Resolver::IfStatement(const ast::IfStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001215 auto* sem =
1216 builder_->create<sem::IfStatement>(stmt, current_compound_statement_, current_function_);
1217 return StatementScope(stmt, sem, [&] {
1218 auto* cond = Expression(stmt->condition);
1219 if (!cond) {
1220 return false;
1221 }
1222 sem->SetCondition(cond);
1223 sem->Behaviors() = cond->Behaviors();
1224 sem->Behaviors().Remove(sem::Behavior::kNext);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001225
dan sinclair41e4d9a2022-05-01 14:40:55 +00001226 Mark(stmt->body);
1227 auto* body = builder_->create<sem::BlockStatement>(stmt->body, current_compound_statement_,
1228 current_function_);
1229 if (!StatementScope(stmt->body, body, [&] { return Statements(stmt->body->statements); })) {
1230 return false;
1231 }
1232 sem->Behaviors().Add(body->Behaviors());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001233
dan sinclair41e4d9a2022-05-01 14:40:55 +00001234 if (stmt->else_statement) {
1235 Mark(stmt->else_statement);
1236 auto* else_sem = Statement(stmt->else_statement);
1237 if (!else_sem) {
1238 return false;
1239 }
1240 sem->Behaviors().Add(else_sem->Behaviors());
1241 } else {
1242 // https://www.w3.org/TR/WGSL/#behaviors-rules
1243 // if statements without an else branch are treated as if they had an
1244 // empty else branch (which adds Next to their behavior)
1245 sem->Behaviors().Add(sem::Behavior::kNext);
1246 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001247
dan sinclair41e4d9a2022-05-01 14:40:55 +00001248 return validator_.IfStatement(sem);
1249 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001250}
1251
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001252sem::BlockStatement* Resolver::BlockStatement(const ast::BlockStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001253 auto* sem = builder_->create<sem::BlockStatement>(
1254 stmt->As<ast::BlockStatement>(), current_compound_statement_, current_function_);
1255 return StatementScope(stmt, sem, [&] { return Statements(stmt->statements); });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001256}
1257
1258sem::LoopStatement* Resolver::LoopStatement(const ast::LoopStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001259 auto* sem =
1260 builder_->create<sem::LoopStatement>(stmt, current_compound_statement_, current_function_);
1261 return StatementScope(stmt, sem, [&] {
1262 Mark(stmt->body);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001263
dan sinclair41e4d9a2022-05-01 14:40:55 +00001264 auto* body = builder_->create<sem::LoopBlockStatement>(
1265 stmt->body, current_compound_statement_, current_function_);
1266 return StatementScope(stmt->body, body, [&] {
1267 if (!Statements(stmt->body->statements)) {
1268 return false;
1269 }
1270 auto& behaviors = sem->Behaviors();
1271 behaviors = body->Behaviors();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001272
dan sinclair41e4d9a2022-05-01 14:40:55 +00001273 if (stmt->continuing) {
1274 Mark(stmt->continuing);
James Price8e68f0a2022-05-11 13:50:33 +00001275 auto* continuing = StatementScope(
1276 stmt->continuing,
1277 builder_->create<sem::LoopContinuingBlockStatement>(
1278 stmt->continuing, current_compound_statement_, current_function_),
1279 [&] { return Statements(stmt->continuing->statements); });
1280 if (!continuing) {
1281 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001282 }
James Price8e68f0a2022-05-11 13:50:33 +00001283 behaviors.Add(continuing->Behaviors());
dan sinclair41e4d9a2022-05-01 14:40:55 +00001284 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001285
dan sinclair41e4d9a2022-05-01 14:40:55 +00001286 if (behaviors.Contains(sem::Behavior::kBreak)) { // Does the loop exit?
1287 behaviors.Add(sem::Behavior::kNext);
1288 } else {
1289 behaviors.Remove(sem::Behavior::kNext);
1290 }
1291 behaviors.Remove(sem::Behavior::kBreak, sem::Behavior::kContinue);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001292
dan sinclair41e4d9a2022-05-01 14:40:55 +00001293 return validator_.LoopStatement(sem);
1294 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001295 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001296}
1297
dan sinclair41e4d9a2022-05-01 14:40:55 +00001298sem::ForLoopStatement* Resolver::ForLoopStatement(const ast::ForLoopStatement* stmt) {
1299 auto* sem = builder_->create<sem::ForLoopStatement>(stmt, current_compound_statement_,
1300 current_function_);
1301 return StatementScope(stmt, sem, [&] {
1302 auto& behaviors = sem->Behaviors();
1303 if (auto* initializer = stmt->initializer) {
1304 Mark(initializer);
1305 auto* init = Statement(initializer);
1306 if (!init) {
1307 return false;
1308 }
1309 behaviors.Add(init->Behaviors());
1310 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001311
dan sinclair41e4d9a2022-05-01 14:40:55 +00001312 if (auto* cond_expr = stmt->condition) {
1313 auto* cond = Expression(cond_expr);
1314 if (!cond) {
1315 return false;
1316 }
1317 sem->SetCondition(cond);
1318 behaviors.Add(cond->Behaviors());
1319 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001320
dan sinclair41e4d9a2022-05-01 14:40:55 +00001321 if (auto* continuing = stmt->continuing) {
1322 Mark(continuing);
1323 auto* cont = Statement(continuing);
1324 if (!cont) {
1325 return false;
1326 }
1327 behaviors.Add(cont->Behaviors());
1328 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001329
dan sinclair41e4d9a2022-05-01 14:40:55 +00001330 Mark(stmt->body);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001331
dan sinclair41e4d9a2022-05-01 14:40:55 +00001332 auto* body = builder_->create<sem::LoopBlockStatement>(
1333 stmt->body, current_compound_statement_, current_function_);
1334 if (!StatementScope(stmt->body, body, [&] { return Statements(stmt->body->statements); })) {
1335 return false;
1336 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001337
dan sinclair41e4d9a2022-05-01 14:40:55 +00001338 behaviors.Add(body->Behaviors());
1339 if (stmt->condition || behaviors.Contains(sem::Behavior::kBreak)) { // Does the loop exit?
1340 behaviors.Add(sem::Behavior::kNext);
1341 } else {
1342 behaviors.Remove(sem::Behavior::kNext);
1343 }
1344 behaviors.Remove(sem::Behavior::kBreak, sem::Behavior::kContinue);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001345
dan sinclair41e4d9a2022-05-01 14:40:55 +00001346 return validator_.ForLoopStatement(sem);
1347 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001348}
1349
dan sinclair49d1a2d2022-06-16 12:01:27 +00001350sem::WhileStatement* Resolver::WhileStatement(const ast::WhileStatement* stmt) {
1351 auto* sem =
1352 builder_->create<sem::WhileStatement>(stmt, current_compound_statement_, current_function_);
1353 return StatementScope(stmt, sem, [&] {
1354 auto& behaviors = sem->Behaviors();
1355
1356 auto* cond = Expression(stmt->condition);
1357 if (!cond) {
1358 return false;
1359 }
1360 sem->SetCondition(cond);
1361 behaviors.Add(cond->Behaviors());
1362
1363 Mark(stmt->body);
1364
1365 auto* body = builder_->create<sem::LoopBlockStatement>(
1366 stmt->body, current_compound_statement_, current_function_);
1367 if (!StatementScope(stmt->body, body, [&] { return Statements(stmt->body->statements); })) {
1368 return false;
1369 }
1370
1371 behaviors.Add(body->Behaviors());
1372 // Always consider the while as having a 'next' behaviour because it has
1373 // a condition. We don't check if the condition will terminate but it isn't
1374 // valid to have an infinite loop in a WGSL program, so a non-terminating
1375 // condition is already an invalid program.
1376 behaviors.Add(sem::Behavior::kNext);
1377 behaviors.Remove(sem::Behavior::kBreak, sem::Behavior::kContinue);
1378
1379 return validator_.WhileStatement(sem);
1380 });
1381}
1382
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001383sem::Expression* Resolver::Expression(const ast::Expression* root) {
Ben Clayton6d669212022-08-03 08:00:15 +00001384 utils::Vector<const ast::Expression*, 64> sorted;
Antonio Maiorano8ba6e1d2022-05-17 15:01:42 +00001385 constexpr size_t kMaxExpressionDepth = 512U;
1386 bool failed = false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001387 if (!ast::TraverseExpressions<ast::TraverseOrder::RightToLeft>(
Antonio Maiorano8ba6e1d2022-05-17 15:01:42 +00001388 root, diagnostics_, [&](const ast::Expression* expr, size_t depth) {
1389 if (depth > kMaxExpressionDepth) {
1390 AddError(
1391 "reached max expression depth of " + std::to_string(kMaxExpressionDepth),
1392 expr->source);
1393 failed = true;
1394 return ast::TraverseAction::Stop;
1395 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001396 if (!Mark(expr)) {
Antonio Maiorano8ba6e1d2022-05-17 15:01:42 +00001397 failed = true;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001398 return ast::TraverseAction::Stop;
1399 }
Ben Clayton958a4642022-07-26 07:55:24 +00001400 sorted.Push(expr);
dan sinclair41e4d9a2022-05-01 14:40:55 +00001401 return ast::TraverseAction::Descend;
1402 })) {
1403 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001404 }
1405
Antonio Maiorano8ba6e1d2022-05-17 15:01:42 +00001406 if (failed) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001407 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001408 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001409
dan sinclair41e4d9a2022-05-01 14:40:55 +00001410 for (auto* expr : utils::Reverse(sorted)) {
1411 auto* sem_expr = Switch(
1412 expr,
1413 [&](const ast::IndexAccessorExpression* array) -> sem::Expression* {
1414 return IndexAccessor(array);
1415 },
1416 [&](const ast::BinaryExpression* bin_op) -> sem::Expression* { return Binary(bin_op); },
1417 [&](const ast::BitcastExpression* bitcast) -> sem::Expression* {
1418 return Bitcast(bitcast);
1419 },
1420 [&](const ast::CallExpression* call) -> sem::Expression* { return Call(call); },
1421 [&](const ast::IdentifierExpression* ident) -> sem::Expression* {
1422 return Identifier(ident);
1423 },
1424 [&](const ast::LiteralExpression* literal) -> sem::Expression* {
1425 return Literal(literal);
1426 },
1427 [&](const ast::MemberAccessorExpression* member) -> sem::Expression* {
1428 return MemberAccessor(member);
1429 },
1430 [&](const ast::UnaryOpExpression* unary) -> sem::Expression* { return UnaryOp(unary); },
1431 [&](const ast::PhonyExpression*) -> sem::Expression* {
1432 return builder_->create<sem::Expression>(expr, builder_->create<sem::Void>(),
Ben Clayton83bd7382022-07-15 23:46:31 +00001433 sem::EvaluationStage::kRuntime,
Ben Claytonaa037ac2022-06-29 19:07:30 +00001434 current_statement_,
1435 /* constant_value */ nullptr,
dan sinclair41e4d9a2022-05-01 14:40:55 +00001436 /* has_side_effects */ false);
1437 },
1438 [&](Default) {
1439 TINT_ICE(Resolver, diagnostics_)
1440 << "unhandled expression type: " << expr->TypeInfo().name;
1441 return nullptr;
1442 });
1443 if (!sem_expr) {
1444 return nullptr;
1445 }
1446
1447 builder_->Sem().Add(expr, sem_expr);
1448 if (expr == root) {
1449 return sem_expr;
1450 }
1451 }
1452
1453 TINT_ICE(Resolver, diagnostics_) << "Expression() did not find root node";
1454 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001455}
1456
Ben Clayton4a656202022-08-01 17:36:54 +00001457const sem::Type* Resolver::ConcreteType(const sem::Type* ty,
1458 const sem::Type* target_ty,
1459 const Source& source) {
Ben Clayton932418e2022-05-24 21:42:03 +00001460 auto i32 = [&] { return builder_->create<sem::I32>(); };
1461 auto f32 = [&] { return builder_->create<sem::F32>(); };
1462 auto i32v = [&](uint32_t width) { return builder_->create<sem::Vector>(i32(), width); };
1463 auto f32v = [&](uint32_t width) { return builder_->create<sem::Vector>(f32(), width); };
1464 auto f32m = [&](uint32_t columns, uint32_t rows) {
Ben Clayton3a236c62022-05-28 05:28:11 +00001465 return builder_->create<sem::Matrix>(f32v(rows), columns);
Ben Clayton932418e2022-05-24 21:42:03 +00001466 };
1467
Ben Clayton05a76182022-07-29 13:19:12 +00001468 return Switch(
1469 ty, //
1470 [&](const sem::AbstractInt*) { return target_ty ? target_ty : i32(); },
1471 [&](const sem::AbstractFloat*) { return target_ty ? target_ty : f32(); },
Ben Clayton932418e2022-05-24 21:42:03 +00001472 [&](const sem::Vector* v) {
1473 return Switch(
1474 v->type(), //
Ben Clayton05a76182022-07-29 13:19:12 +00001475 [&](const sem::AbstractInt*) { return target_ty ? target_ty : i32v(v->Width()); },
Ben Clayton932418e2022-05-24 21:42:03 +00001476 [&](const sem::AbstractFloat*) {
Ben Clayton05a76182022-07-29 13:19:12 +00001477 return target_ty ? target_ty : f32v(v->Width());
1478 });
Ben Clayton932418e2022-05-24 21:42:03 +00001479 },
1480 [&](const sem::Matrix* m) {
Ben Clayton05a76182022-07-29 13:19:12 +00001481 return Switch(m->type(), //
1482 [&](const sem::AbstractFloat*) {
1483 return target_ty ? target_ty : f32m(m->columns(), m->rows());
1484 });
Ben Clayton4a656202022-08-01 17:36:54 +00001485 },
1486 [&](const sem::Array* a) -> const sem::Type* {
1487 const sem::Type* target_el_ty = nullptr;
1488 if (auto* target_arr_ty = As<sem::Array>(target_ty)) {
1489 target_el_ty = target_arr_ty->ElemType();
1490 }
1491 if (auto* el_ty = ConcreteType(a->ElemType(), target_el_ty, source)) {
dan sinclair78f80672022-09-22 22:28:21 +00001492 return Array(source, source, el_ty, a->Count(), /* explicit_stride */ 0);
Ben Clayton4a656202022-08-01 17:36:54 +00001493 }
1494 return nullptr;
Ben Clayton05a76182022-07-29 13:19:12 +00001495 });
1496}
1497
1498const sem::Expression* Resolver::Materialize(const sem::Expression* expr,
1499 const sem::Type* target_type /* = nullptr */) {
1500 if (!expr) {
1501 // Allow for Materialize(Expression(blah)), where failures pass through Materialize()
1502 return nullptr;
1503 }
1504
1505 auto* decl = expr->Declaration();
1506
Ben Clayton4a656202022-08-01 17:36:54 +00001507 auto* concrete_ty = ConcreteType(expr->Type(), target_type, decl->source);
Ben Clayton05a76182022-07-29 13:19:12 +00001508 if (!concrete_ty) {
1509 return expr; // Does not require materialization
1510 }
1511
1512 auto* src_ty = expr->Type();
1513 if (!validator_.Materialize(concrete_ty, src_ty, decl->source)) {
1514 return nullptr;
1515 }
1516
1517 auto expr_val = expr->ConstantValue();
1518 if (!expr_val) {
1519 TINT_ICE(Resolver, builder_->Diagnostics())
1520 << decl->source << "Materialize(" << decl->TypeInfo().name
1521 << ") called on expression with no constant value";
1522 return nullptr;
1523 }
1524
1525 auto materialized_val = const_eval_.Convert(concrete_ty, expr_val, decl->source);
1526 if (!materialized_val) {
1527 // ConvertValue() has already failed and raised an diagnostic error.
1528 return nullptr;
1529 }
1530
1531 if (!materialized_val.Get()) {
1532 TINT_ICE(Resolver, builder_->Diagnostics())
1533 << decl->source << "ConvertValue(" << builder_->FriendlyName(expr_val->Type()) << " -> "
1534 << builder_->FriendlyName(concrete_ty) << ") returned invalid value";
1535 return nullptr;
1536 }
1537 auto* m = builder_->create<sem::Materialize>(expr, current_statement_, materialized_val.Get());
1538 m->Behaviors() = expr->Behaviors();
1539 builder_->Sem().Replace(decl, m);
1540 return m;
Ben Clayton932418e2022-05-24 21:42:03 +00001541}
1542
Ben Clayton34d46732022-08-02 16:52:25 +00001543template <size_t N>
1544bool Resolver::MaterializeArguments(utils::Vector<const sem::Expression*, N>& args,
Ben Clayton932418e2022-05-24 21:42:03 +00001545 const sem::CallTarget* target) {
Ben Clayton958a4642022-07-26 07:55:24 +00001546 for (size_t i = 0, n = std::min(args.Length(), target->Parameters().Length()); i < n; i++) {
Ben Clayton932418e2022-05-24 21:42:03 +00001547 const auto* param_ty = target->Parameters()[i]->Type();
1548 if (ShouldMaterializeArgument(param_ty)) {
1549 auto* materialized = Materialize(args[i], param_ty);
1550 if (!materialized) {
1551 return false;
1552 }
1553 args[i] = materialized;
1554 }
1555 }
1556 return true;
1557}
1558
1559bool Resolver::ShouldMaterializeArgument(const sem::Type* parameter_ty) const {
Antonio Maioranoa8e9a6e2022-08-10 14:32:19 +00001560 const auto* param_el_ty = sem::Type::DeepestElementOf(parameter_ty);
Ben Clayton932418e2022-05-24 21:42:03 +00001561 return param_el_ty && !param_el_ty->Is<sem::AbstractNumeric>();
1562}
1563
Antonio Maioranoa58d8c92022-08-10 20:01:17 +00001564bool Resolver::Convert(const sem::Constant*& c, const sem::Type* target_ty, const Source& source) {
1565 auto r = const_eval_.Convert(target_ty, c, source);
1566 if (!r) {
1567 return false;
1568 }
1569 c = r.Get();
1570 return true;
1571}
1572
1573template <size_t N>
1574utils::Result<utils::Vector<const sem::Constant*, N>> Resolver::ConvertArguments(
1575 const utils::Vector<const sem::Expression*, N>& args,
1576 const sem::CallTarget* target) {
1577 auto const_args = utils::Transform(args, [](auto* arg) { return arg->ConstantValue(); });
1578 for (size_t i = 0, n = std::min(args.Length(), target->Parameters().Length()); i < n; i++) {
1579 if (!Convert(const_args[i], target->Parameters()[i]->Type(),
1580 args[i]->Declaration()->source)) {
1581 return utils::Failure;
1582 }
1583 }
1584 return const_args;
1585}
1586
dan sinclair41e4d9a2022-05-01 14:40:55 +00001587sem::Expression* Resolver::IndexAccessor(const ast::IndexAccessorExpression* expr) {
Ben Clayton49a09142022-05-31 15:22:21 +00001588 auto* idx = Materialize(sem_.Get(expr->index));
1589 if (!idx) {
1590 return nullptr;
1591 }
Ben Clayton57166112022-07-21 15:25:35 +00001592 const auto* obj = sem_.Get(expr->object);
1593 if (idx->Stage() != sem::EvaluationStage::kConstant) {
1594 // If the index is non-constant, then the resulting expression is non-constant, so we'll
1595 // have to materialize the object. For example, consider:
1596 // vec2(1, 2)[runtime-index]
1597 obj = Materialize(obj);
1598 }
Ben Clayton4aa360d2022-08-01 18:46:54 +00001599 if (!obj) {
1600 return nullptr;
1601 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001602 auto* obj_raw_ty = obj->Type();
1603 auto* obj_ty = obj_raw_ty->UnwrapRef();
1604 auto* ty = Switch(
1605 obj_ty, //
1606 [&](const sem::Array* arr) { return arr->ElemType(); },
1607 [&](const sem::Vector* vec) { return vec->type(); },
1608 [&](const sem::Matrix* mat) {
1609 return builder_->create<sem::Vector>(mat->type(), mat->rows());
1610 },
1611 [&](Default) {
1612 AddError("cannot index type '" + sem_.TypeNameOf(obj_ty) + "'", expr->source);
1613 return nullptr;
1614 });
1615 if (ty == nullptr) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001616 return nullptr;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001617 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001618
dan sinclair41e4d9a2022-05-01 14:40:55 +00001619 auto* idx_ty = idx->Type()->UnwrapRef();
1620 if (!idx_ty->IsAnyOf<sem::I32, sem::U32>()) {
1621 AddError("index must be of type 'i32' or 'u32', found: '" + sem_.TypeNameOf(idx_ty) + "'",
1622 idx->Declaration()->source);
1623 return nullptr;
1624 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001625
dan sinclair41e4d9a2022-05-01 14:40:55 +00001626 // If we're extracting from a reference, we return a reference.
1627 if (auto* ref = obj_raw_ty->As<sem::Reference>()) {
1628 ty = builder_->create<sem::Reference>(ty, ref->StorageClass(), ref->Access());
1629 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001630
Ben Clayton83bd7382022-07-15 23:46:31 +00001631 auto stage = sem::EarliestStage(obj->Stage(), idx->Stage());
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00001632 const sem::Constant* val = nullptr;
1633 if (auto r = const_eval_.Index(obj, idx)) {
1634 val = r.Get();
1635 } else {
1636 return nullptr;
1637 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001638 bool has_side_effects = idx->HasSideEffects() || obj->HasSideEffects();
Antonio Maioranodfeaf2902022-06-24 20:34:00 +00001639 auto* sem = builder_->create<sem::IndexAccessorExpression>(
Ben Clayton83bd7382022-07-15 23:46:31 +00001640 expr, ty, stage, obj, idx, current_statement_, std::move(val), has_side_effects,
Antonio Maioranodfeaf2902022-06-24 20:34:00 +00001641 obj->SourceVariable());
dan sinclair41e4d9a2022-05-01 14:40:55 +00001642 sem->Behaviors() = idx->Behaviors() + obj->Behaviors();
1643 return sem;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001644}
1645
1646sem::Expression* Resolver::Bitcast(const ast::BitcastExpression* expr) {
Ben Clayton8c7ffbd2022-05-28 07:09:17 +00001647 auto* inner = Materialize(sem_.Get(expr->expr));
1648 if (!inner) {
1649 return nullptr;
1650 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001651 auto* ty = Type(expr->type);
1652 if (!ty) {
1653 return nullptr;
1654 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001655
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00001656 const sem::Constant* val = nullptr;
1657 if (auto r = const_eval_.Bitcast(ty, inner)) {
1658 val = r.Get();
1659 } else {
1660 return nullptr;
1661 }
Ben Clayton83bd7382022-07-15 23:46:31 +00001662 auto stage = sem::EvaluationStage::kRuntime; // TODO(crbug.com/tint/1581)
1663 auto* sem = builder_->create<sem::Expression>(expr, ty, stage, current_statement_,
1664 std::move(val), inner->HasSideEffects());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001665
dan sinclair41e4d9a2022-05-01 14:40:55 +00001666 sem->Behaviors() = inner->Behaviors();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001667
dan sinclair41e4d9a2022-05-01 14:40:55 +00001668 if (!validator_.Bitcast(expr, ty)) {
1669 return nullptr;
1670 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001671
dan sinclair41e4d9a2022-05-01 14:40:55 +00001672 return sem;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001673}
1674
1675sem::Call* Resolver::Call(const ast::CallExpression* expr) {
Ben Clayton6ae608c2022-05-16 20:54:42 +00001676 // A CallExpression can resolve to one of:
1677 // * A function call.
1678 // * A builtin call.
1679 // * A type constructor.
1680 // * A type conversion.
1681
1682 // Resolve all of the arguments, their types and the set of behaviors.
Ben Clayton958a4642022-07-26 07:55:24 +00001683 utils::Vector<const sem::Expression*, 8> args;
Ben Clayton783b1692022-08-02 17:03:35 +00001684 args.Reserve(expr->args.Length());
Ben Clayton83bd7382022-07-15 23:46:31 +00001685 auto args_stage = sem::EvaluationStage::kConstant;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001686 sem::Behaviors arg_behaviors;
Ben Clayton783b1692022-08-02 17:03:35 +00001687 for (size_t i = 0; i < expr->args.Length(); i++) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001688 auto* arg = sem_.Get(expr->args[i]);
1689 if (!arg) {
1690 return nullptr;
1691 }
Ben Clayton958a4642022-07-26 07:55:24 +00001692 args.Push(arg);
Ben Clayton83bd7382022-07-15 23:46:31 +00001693 args_stage = sem::EarliestStage(args_stage, arg->Stage());
dan sinclair41e4d9a2022-05-01 14:40:55 +00001694 arg_behaviors.Add(arg->Behaviors());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001695 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001696 arg_behaviors.Remove(sem::Behavior::kNext);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001697
Ben Clayton6ae608c2022-05-16 20:54:42 +00001698 // Did any arguments have side effects?
1699 bool has_side_effects =
1700 std::any_of(args.begin(), args.end(), [](auto* e) { return e->HasSideEffects(); });
1701
Ben Clayton6ae608c2022-05-16 20:54:42 +00001702 // ct_ctor_or_conv is a helper for building either a sem::TypeConstructor or sem::TypeConversion
1703 // call for a CtorConvIntrinsic with an optional template argument type.
1704 auto ct_ctor_or_conv = [&](CtorConvIntrinsic ty, const sem::Type* template_arg) -> sem::Call* {
Ben Clayton7b921fb2022-05-19 20:19:49 +00001705 auto arg_tys = utils::Transform(args, [](auto* arg) { return arg->Type(); });
Ben Clayton6a80ce62022-07-15 22:27:50 +00001706 auto ctor_or_conv = intrinsic_table_->Lookup(ty, template_arg, arg_tys, expr->source);
1707 if (!ctor_or_conv.target) {
Ben Clayton6ae608c2022-05-16 20:54:42 +00001708 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001709 }
Ben Clayton6a80ce62022-07-15 22:27:50 +00001710 if (!MaterializeArguments(args, ctor_or_conv.target)) {
Ben Clayton932418e2022-05-24 21:42:03 +00001711 return nullptr;
1712 }
Ben Clayton6a80ce62022-07-15 22:27:50 +00001713 const sem::Constant* value = nullptr;
Ben Clayton83bd7382022-07-15 23:46:31 +00001714 auto stage = sem::EarliestStage(ctor_or_conv.target->Stage(), args_stage);
1715 if (stage == sem::EvaluationStage::kConstant) {
Antonio Maioranoff0295e2022-08-31 13:48:41 +00001716 auto const_args = ConvertArguments(args, ctor_or_conv.target);
1717 if (!const_args) {
1718 return nullptr;
1719 }
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00001720 if (auto r = (const_eval_.*ctor_or_conv.const_eval_fn)(
Antonio Maioranoff0295e2022-08-31 13:48:41 +00001721 ctor_or_conv.target->ReturnType(), const_args.Get(), expr->source)) {
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00001722 value = r.Get();
1723 } else {
1724 return nullptr;
1725 }
Ben Clayton6a80ce62022-07-15 22:27:50 +00001726 }
Ben Clayton83bd7382022-07-15 23:46:31 +00001727 return builder_->create<sem::Call>(expr, ctor_or_conv.target, stage, std::move(args),
Ben Clayton6a80ce62022-07-15 22:27:50 +00001728 current_statement_, value, has_side_effects);
Ben Clayton6ae608c2022-05-16 20:54:42 +00001729 };
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001730
Ben Claytonac660c22022-07-15 23:54:10 +00001731 // arr_or_str_ctor is a helper for building a sem::TypeConstructor for an array or structure
Ben Clayton83bd7382022-07-15 23:46:31 +00001732 // constructor call target.
1733 auto arr_or_str_ctor = [&](const sem::Type* ty,
1734 const sem::CallTarget* call_target) -> sem::Call* {
1735 if (!MaterializeArguments(args, call_target)) {
1736 return nullptr;
1737 }
1738
1739 auto stage = args_stage; // The evaluation stage of the call
1740 const sem::Constant* value = nullptr; // The constant value for the call
1741 if (stage == sem::EvaluationStage::kConstant) {
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00001742 if (auto r = const_eval_.ArrayOrStructCtor(ty, args)) {
1743 value = r.Get();
1744 } else {
1745 return nullptr;
1746 }
Ben Clayton83bd7382022-07-15 23:46:31 +00001747 if (!value) {
1748 // Constant evaluation failed.
1749 // Can happen for expressions that will fail validation (later).
Ben Claytonac660c22022-07-15 23:54:10 +00001750 // Use the kRuntime EvaluationStage, as kConstant will trigger an assertion in the
1751 // sem::Expression constructor, which checks that kConstant is paired with a
1752 // constant value.
Ben Clayton83bd7382022-07-15 23:46:31 +00001753 stage = sem::EvaluationStage::kRuntime;
1754 }
1755 }
1756
1757 return builder_->create<sem::Call>(expr, call_target, stage, std::move(args),
Ben Claytonac660c22022-07-15 23:54:10 +00001758 current_statement_, value, has_side_effects);
Ben Clayton83bd7382022-07-15 23:46:31 +00001759 };
1760
Ben Claytonac660c22022-07-15 23:54:10 +00001761 // ty_ctor_or_conv is a helper for building either a sem::TypeConstructor or sem::TypeConversion
Ben Clayton6ae608c2022-05-16 20:54:42 +00001762 // call for the given semantic type.
1763 auto ty_ctor_or_conv = [&](const sem::Type* ty) {
1764 return Switch(
1765 ty, //
1766 [&](const sem::Vector* v) {
1767 return ct_ctor_or_conv(VectorCtorConvIntrinsic(v->Width()), v->type());
1768 },
1769 [&](const sem::Matrix* m) {
1770 return ct_ctor_or_conv(MatrixCtorConvIntrinsic(m->columns(), m->rows()), m->type());
1771 },
1772 [&](const sem::I32*) { return ct_ctor_or_conv(CtorConvIntrinsic::kI32, nullptr); },
1773 [&](const sem::U32*) { return ct_ctor_or_conv(CtorConvIntrinsic::kU32, nullptr); },
Zhaoming Jiang60588822022-06-28 14:03:36 +00001774 [&](const sem::F16*) { return ct_ctor_or_conv(CtorConvIntrinsic::kF16, nullptr); },
Ben Clayton6ae608c2022-05-16 20:54:42 +00001775 [&](const sem::F32*) { return ct_ctor_or_conv(CtorConvIntrinsic::kF32, nullptr); },
1776 [&](const sem::Bool*) { return ct_ctor_or_conv(CtorConvIntrinsic::kBool, nullptr); },
Ben Clayton7b921fb2022-05-19 20:19:49 +00001777 [&](const sem::Array* arr) -> sem::Call* {
1778 auto* call_target = utils::GetOrCreate(
Ben Clayton958a4642022-07-26 07:55:24 +00001779 array_ctors_, ArrayConstructorSig{{arr, args.Length(), args_stage}},
Ben Clayton7b921fb2022-05-19 20:19:49 +00001780 [&]() -> sem::TypeConstructor* {
Ben Clayton4a656202022-08-01 17:36:54 +00001781 auto params = utils::Transform(args, [&](auto, size_t i) {
1782 return builder_->create<sem::Parameter>(
1783 nullptr, // declaration
1784 static_cast<uint32_t>(i), // index
1785 arr->ElemType(), // type
1786 ast::StorageClass::kNone, // storage_class
1787 ast::Access::kUndefined);
1788 });
Ben Clayton83bd7382022-07-15 23:46:31 +00001789 return builder_->create<sem::TypeConstructor>(arr, std::move(params),
1790 args_stage);
Ben Clayton7b921fb2022-05-19 20:19:49 +00001791 });
Ben Clayton8c9cc862022-08-01 15:19:54 +00001792
1793 auto* call = arr_or_str_ctor(arr, call_target);
1794 if (!call) {
1795 return nullptr;
1796 }
1797
1798 // Validation must occur after argument materialization in arr_or_str_ctor().
1799 if (!validator_.ArrayConstructor(expr, arr)) {
1800 return nullptr;
1801 }
1802 return call;
Ben Clayton7b921fb2022-05-19 20:19:49 +00001803 },
1804 [&](const sem::Struct* str) -> sem::Call* {
1805 auto* call_target = utils::GetOrCreate(
Ben Clayton958a4642022-07-26 07:55:24 +00001806 struct_ctors_, StructConstructorSig{{str, args.Length(), args_stage}},
Ben Clayton7b921fb2022-05-19 20:19:49 +00001807 [&]() -> sem::TypeConstructor* {
Ben Clayton958a4642022-07-26 07:55:24 +00001808 utils::Vector<const sem::Parameter*, 8> params;
1809 params.Resize(std::min(args.Length(), str->Members().size()));
1810 for (size_t i = 0, n = params.Length(); i < n; i++) {
Ben Clayton7b921fb2022-05-19 20:19:49 +00001811 params[i] = builder_->create<sem::Parameter>(
1812 nullptr, // declaration
1813 static_cast<uint32_t>(i), // index
1814 str->Members()[i]->Type(), // type
1815 ast::StorageClass::kNone, // storage_class
1816 ast::Access::kUndefined); // access
1817 }
Ben Clayton83bd7382022-07-15 23:46:31 +00001818 return builder_->create<sem::TypeConstructor>(str, std::move(params),
1819 args_stage);
Ben Clayton7b921fb2022-05-19 20:19:49 +00001820 });
Ben Clayton8c9cc862022-08-01 15:19:54 +00001821
1822 auto* call = arr_or_str_ctor(str, call_target);
1823 if (!call) {
1824 return nullptr;
1825 }
1826
1827 // Validation must occur after argument materialization in arr_or_str_ctor().
1828 if (!validator_.StructureConstructor(expr, str)) {
1829 return nullptr;
1830 }
1831 return call;
Ben Clayton7b921fb2022-05-19 20:19:49 +00001832 },
Ben Clayton6ae608c2022-05-16 20:54:42 +00001833 [&](Default) {
1834 AddError("type is not constructible", expr->source);
dan sinclair41e4d9a2022-05-01 14:40:55 +00001835 return nullptr;
Ben Clayton6ae608c2022-05-16 20:54:42 +00001836 });
1837 };
dan sinclair41e4d9a2022-05-01 14:40:55 +00001838
Ben Clayton6ae608c2022-05-16 20:54:42 +00001839 // ast::CallExpression has a target which is either an ast::Type or an ast::IdentifierExpression
1840 sem::Call* call = nullptr;
1841 if (expr->target.type) {
1842 // ast::CallExpression has an ast::Type as the target.
1843 // This call is either a type constructor or type conversion.
1844 call = Switch(
1845 expr->target.type,
1846 [&](const ast::Vector* v) -> sem::Call* {
1847 Mark(v);
1848 // vector element type must be inferred if it was not specified.
1849 sem::Type* template_arg = nullptr;
1850 if (v->type) {
1851 template_arg = Type(v->type);
1852 if (!template_arg) {
1853 return nullptr;
1854 }
1855 }
1856 if (auto* c = ct_ctor_or_conv(VectorCtorConvIntrinsic(v->width), template_arg)) {
1857 builder_->Sem().Add(expr->target.type, c->Target()->ReturnType());
1858 return c;
1859 }
1860 return nullptr;
1861 },
1862 [&](const ast::Matrix* m) -> sem::Call* {
1863 Mark(m);
1864 // matrix element type must be inferred if it was not specified.
1865 sem::Type* template_arg = nullptr;
1866 if (m->type) {
1867 template_arg = Type(m->type);
1868 if (!template_arg) {
1869 return nullptr;
1870 }
1871 }
1872 if (auto* c = ct_ctor_or_conv(MatrixCtorConvIntrinsic(m->columns, m->rows),
1873 template_arg)) {
1874 builder_->Sem().Add(expr->target.type, c->Target()->ReturnType());
1875 return c;
1876 }
1877 return nullptr;
1878 },
Ben Clayton4a656202022-08-01 17:36:54 +00001879 [&](const ast::Array* a) -> sem::Call* {
1880 Mark(a);
1881 // array element type must be inferred if it was not specified.
dan sinclair78f80672022-09-22 22:28:21 +00001882 sem::ArrayCount el_count =
1883 sem::ConstantArrayCount{static_cast<uint32_t>(args.Length())};
Ben Clayton4a656202022-08-01 17:36:54 +00001884 const sem::Type* el_ty = nullptr;
1885 if (a->type) {
1886 el_ty = Type(a->type);
1887 if (!el_ty) {
1888 return nullptr;
1889 }
1890 if (!a->count) {
1891 AddError("cannot construct a runtime-sized array", expr->source);
1892 return nullptr;
1893 }
1894 if (auto count = ArrayCount(a->count)) {
1895 el_count = count.Get();
1896 } else {
1897 return nullptr;
1898 }
1899 // Note: validation later will detect any mismatches between explicit array
1900 // size and number of constructor expressions.
1901 } else {
1902 auto arg_tys =
1903 utils::Transform(args, [](auto* arg) { return arg->Type()->UnwrapRef(); });
1904 el_ty = sem::Type::Common(arg_tys);
1905 if (!el_ty) {
1906 AddError(
1907 "cannot infer common array element type from constructor arguments",
1908 expr->source);
1909 std::unordered_set<const sem::Type*> types;
1910 for (size_t i = 0; i < args.Length(); i++) {
1911 if (types.emplace(args[i]->Type()).second) {
1912 AddNote("argument " + std::to_string(i) + " is of type '" +
1913 sem_.TypeNameOf(args[i]->Type()) + "'",
1914 args[i]->Declaration()->source);
1915 }
1916 }
1917 return nullptr;
1918 }
1919 }
1920 uint32_t explicit_stride = 0;
1921 if (!ArrayAttributes(a->attributes, el_ty, explicit_stride)) {
1922 return nullptr;
1923 }
1924
dan sinclair78f80672022-09-22 22:28:21 +00001925 auto* arr = Array(a->type ? a->type->source : a->source,
1926 a->count ? a->count->source : a->source, //
1927 el_ty, el_count, explicit_stride);
Ben Clayton4a656202022-08-01 17:36:54 +00001928 if (!arr) {
1929 return nullptr;
1930 }
1931 builder_->Sem().Add(a, arr);
1932
1933 return ty_ctor_or_conv(arr);
1934 },
Ben Clayton6ae608c2022-05-16 20:54:42 +00001935 [&](const ast::Type* ast) -> sem::Call* {
1936 // Handler for AST types that do not have an optional element type.
1937 if (auto* ty = Type(ast)) {
1938 return ty_ctor_or_conv(ty);
1939 }
1940 return nullptr;
1941 },
1942 [&](Default) {
1943 TINT_ICE(Resolver, diagnostics_)
1944 << expr->source << " unhandled CallExpression target:\n"
1945 << "type: "
1946 << (expr->target.type ? expr->target.type->TypeInfo().name : "<null>");
1947 return nullptr;
1948 });
1949 } else {
1950 // ast::CallExpression has an ast::IdentifierExpression as the target.
1951 // This call is either a function call, builtin call, type constructor or type conversion.
1952 auto* ident = expr->target.name;
1953 Mark(ident);
1954 auto* resolved = sem_.ResolvedSymbol(ident);
1955 call = Switch<sem::Call*>(
1956 resolved, //
1957 [&](sem::Type* ty) {
1958 // A type constructor or conversions.
Ben Clayton7b921fb2022-05-19 20:19:49 +00001959 // Note: Unlike the code path where we're resolving the call target from an
Ben Clayton6ae608c2022-05-16 20:54:42 +00001960 // ast::Type, all types must already have the element type explicitly specified, so
1961 // there's no need to infer element types.
1962 return ty_ctor_or_conv(ty);
1963 },
Ben Clayton34d46732022-08-02 16:52:25 +00001964 [&](sem::Function* func) { return FunctionCall(expr, func, args, arg_behaviors); },
Ben Clayton6ae608c2022-05-16 20:54:42 +00001965 [&](sem::Variable* var) {
1966 auto name = builder_->Symbols().NameFor(var->Declaration()->symbol);
1967 AddError("cannot call variable '" + name + "'", ident->source);
1968 AddNote("'" + name + "' declared here", var->Declaration()->source);
1969 return nullptr;
1970 },
1971 [&](Default) -> sem::Call* {
1972 auto name = builder_->Symbols().NameFor(ident->symbol);
1973 auto builtin_type = sem::ParseBuiltinType(name);
1974 if (builtin_type != sem::BuiltinType::kNone) {
Ben Clayton34d46732022-08-02 16:52:25 +00001975 return BuiltinCall(expr, builtin_type, args);
Ben Clayton6ae608c2022-05-16 20:54:42 +00001976 }
1977
1978 TINT_ICE(Resolver, diagnostics_)
1979 << expr->source << " unhandled CallExpression target:\n"
1980 << "resolved: " << (resolved ? resolved->TypeInfo().name : "<null>") << "\n"
1981 << "name: " << builder_->Symbols().NameFor(ident->symbol);
1982 return nullptr;
1983 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001984 }
1985
Ben Clayton6ae608c2022-05-16 20:54:42 +00001986 if (!call) {
1987 return nullptr;
1988 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001989
Ben Clayton6ae608c2022-05-16 20:54:42 +00001990 return validator_.Call(call, current_statement_) ? call : nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001991}
1992
Ben Clayton34d46732022-08-02 16:52:25 +00001993template <size_t N>
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001994sem::Call* Resolver::BuiltinCall(const ast::CallExpression* expr,
1995 sem::BuiltinType builtin_type,
Ben Clayton34d46732022-08-02 16:52:25 +00001996 utils::Vector<const sem::Expression*, N>& args) {
Ben Clayton451eee02022-06-01 23:57:20 +00001997 IntrinsicTable::Builtin builtin;
Ben Clayton7b921fb2022-05-19 20:19:49 +00001998 {
Ben Clayton34d46732022-08-02 16:52:25 +00001999 auto arg_tys = utils::Transform(args, [](auto* arg) { return arg->Type(); });
Ben Clayton7b921fb2022-05-19 20:19:49 +00002000 builtin = intrinsic_table_->Lookup(builtin_type, arg_tys, expr->source);
Ben Clayton451eee02022-06-01 23:57:20 +00002001 if (!builtin.sem) {
Ben Clayton7b921fb2022-05-19 20:19:49 +00002002 return nullptr;
2003 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002004 }
2005
Ben Clayton451eee02022-06-01 23:57:20 +00002006 if (!MaterializeArguments(args, builtin.sem)) {
Ben Clayton932418e2022-05-24 21:42:03 +00002007 return nullptr;
2008 }
2009
Ben Clayton451eee02022-06-01 23:57:20 +00002010 if (builtin.sem->IsDeprecated()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002011 AddWarning("use of deprecated builtin", expr->source);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002012 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002013
Ben Clayton83bd7382022-07-15 23:46:31 +00002014 auto stage = builtin.sem->Stage();
Ben Claytonac660c22022-07-15 23:54:10 +00002015 if (stage == sem::EvaluationStage::kConstant) { // <-- Optimization
2016 // If the builtin is not annotated with @const, then it can only be evaluated
2017 // at runtime, in which case there's no point checking the evaluation stage of the
2018 // arguments.
2019
2020 // The builtin is @const annotated. Check all arguments are also constant.
Ben Clayton83bd7382022-07-15 23:46:31 +00002021 for (auto* arg : args) {
2022 stage = sem::EarliestStage(stage, arg->Stage());
2023 }
2024 }
2025
Ben Clayton451eee02022-06-01 23:57:20 +00002026 // If the builtin is @const, and all arguments have constant values, evaluate the builtin now.
Ben Clayton65c5c9d2022-07-15 14:14:09 +00002027 const sem::Constant* value = nullptr;
Ben Claytonac660c22022-07-15 23:54:10 +00002028 if (stage == sem::EvaluationStage::kConstant) {
Antonio Maioranoa58d8c92022-08-10 20:01:17 +00002029 auto const_args = ConvertArguments(args, builtin.sem);
2030 if (!const_args) {
2031 return nullptr;
2032 }
2033 if (auto r = (const_eval_.*builtin.const_eval_fn)(builtin.sem->ReturnType(),
2034 const_args.Get(), expr->source)) {
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00002035 value = r.Get();
2036 } else {
2037 return nullptr;
2038 }
Ben Clayton451eee02022-06-01 23:57:20 +00002039 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002040
Ben Clayton451eee02022-06-01 23:57:20 +00002041 bool has_side_effects =
2042 builtin.sem->HasSideEffects() ||
2043 std::any_of(args.begin(), args.end(), [](auto* e) { return e->HasSideEffects(); });
Ben Clayton83bd7382022-07-15 23:46:31 +00002044 auto* call = builder_->create<sem::Call>(expr, builtin.sem, stage, std::move(args),
2045 current_statement_, value, has_side_effects);
Ben Clayton451eee02022-06-01 23:57:20 +00002046
Ben Claytonb6c8ea92022-07-07 17:04:21 +00002047 if (current_function_) {
2048 current_function_->AddDirectlyCalledBuiltin(builtin.sem);
2049 current_function_->AddDirectCall(call);
2050 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002051
Ben Clayton7f2b8cd2022-05-18 22:41:48 +00002052 if (!validator_.RequiredExtensionForBuiltinFunction(call, enabled_extensions_)) {
Jiawei Shao6c9e4752022-05-10 09:05:54 +00002053 return nullptr;
2054 }
2055
dan sinclair41e4d9a2022-05-01 14:40:55 +00002056 if (IsTextureBuiltin(builtin_type)) {
2057 if (!validator_.TextureBuiltinFunction(call)) {
2058 return nullptr;
2059 }
Ben Clayton451eee02022-06-01 23:57:20 +00002060 CollectTextureSamplerPairs(builtin.sem, call->Arguments());
dan sinclair41e4d9a2022-05-01 14:40:55 +00002061 }
2062
2063 if (!validator_.BuiltinCall(call)) {
2064 return nullptr;
2065 }
2066
dan sinclair41e4d9a2022-05-01 14:40:55 +00002067 return call;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002068}
2069
Ben Clayton34d46732022-08-02 16:52:25 +00002070void Resolver::CollectTextureSamplerPairs(const sem::Builtin* builtin,
2071 utils::VectorRef<const sem::Expression*> args) const {
Ben Clayton7b921fb2022-05-19 20:19:49 +00002072 // Collect a texture/sampler pair for this builtin.
2073 const auto& signature = builtin->Signature();
2074 int texture_index = signature.IndexOf(sem::ParameterUsage::kTexture);
2075 if (texture_index == -1) {
2076 TINT_ICE(Resolver, diagnostics_) << "texture builtin without texture parameter";
2077 }
dan sinclair3a2a2792022-06-29 14:38:15 +00002078 auto* texture = args[static_cast<size_t>(texture_index)]->As<sem::VariableUser>()->Variable();
Ben Clayton7b921fb2022-05-19 20:19:49 +00002079 if (!texture->Type()->UnwrapRef()->Is<sem::StorageTexture>()) {
2080 int sampler_index = signature.IndexOf(sem::ParameterUsage::kSampler);
2081 const sem::Variable* sampler =
dan sinclair3a2a2792022-06-29 14:38:15 +00002082 sampler_index != -1
2083 ? args[static_cast<size_t>(sampler_index)]->As<sem::VariableUser>()->Variable()
2084 : nullptr;
Ben Clayton7b921fb2022-05-19 20:19:49 +00002085 current_function_->AddTextureSamplerPair(texture, sampler);
2086 }
2087}
2088
Ben Clayton34d46732022-08-02 16:52:25 +00002089template <size_t N>
dan sinclair41e4d9a2022-05-01 14:40:55 +00002090sem::Call* Resolver::FunctionCall(const ast::CallExpression* expr,
2091 sem::Function* target,
Ben Clayton34d46732022-08-02 16:52:25 +00002092 utils::Vector<const sem::Expression*, N>& args,
dan sinclair41e4d9a2022-05-01 14:40:55 +00002093 sem::Behaviors arg_behaviors) {
2094 auto sym = expr->target.name->symbol;
2095 auto name = builder_->Symbols().NameFor(sym);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002096
Ben Clayton932418e2022-05-24 21:42:03 +00002097 if (!MaterializeArguments(args, target)) {
2098 return nullptr;
2099 }
2100
dan sinclair41e4d9a2022-05-01 14:40:55 +00002101 // TODO(crbug.com/tint/1420): For now, assume all function calls have side
2102 // effects.
2103 bool has_side_effects = true;
Ben Clayton83bd7382022-07-15 23:46:31 +00002104 auto* call = builder_->create<sem::Call>(expr, target, sem::EvaluationStage::kRuntime,
2105 std::move(args), current_statement_,
Ben Claytonaa037ac2022-06-29 19:07:30 +00002106 /* constant_value */ nullptr, has_side_effects);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002107
Ben Clayton4b047212022-05-27 20:16:26 +00002108 target->AddCallSite(call);
2109
2110 call->Behaviors() = arg_behaviors + target->Behaviors();
2111
2112 if (!validator_.FunctionCall(call, current_statement_)) {
2113 return nullptr;
2114 }
2115
dan sinclair41e4d9a2022-05-01 14:40:55 +00002116 if (current_function_) {
2117 // Note: Requires called functions to be resolved first.
2118 // This is currently guaranteed as functions must be declared before
2119 // use.
2120 current_function_->AddTransitivelyCalledFunction(target);
2121 current_function_->AddDirectCall(call);
2122 for (auto* transitive_call : target->TransitivelyCalledFunctions()) {
2123 current_function_->AddTransitivelyCalledFunction(transitive_call);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002124 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002125
2126 // We inherit any referenced variables from the callee.
2127 for (auto* var : target->TransitivelyReferencedGlobals()) {
2128 current_function_->AddTransitivelyReferencedGlobal(var);
2129 }
2130
Ben Clayton4b047212022-05-27 20:16:26 +00002131 // Note: Validation *must* be performed before calling this method.
Ben Clayton7b921fb2022-05-19 20:19:49 +00002132 CollectTextureSamplerPairs(target, call->Arguments());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002133 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002134
dan sinclair41e4d9a2022-05-01 14:40:55 +00002135 return call;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002136}
2137
Ben Clayton34d46732022-08-02 16:52:25 +00002138void Resolver::CollectTextureSamplerPairs(sem::Function* func,
2139 utils::VectorRef<const sem::Expression*> args) const {
Ben Clayton7b921fb2022-05-19 20:19:49 +00002140 // Map all texture/sampler pairs from the target function to the
2141 // current function. These can only be global or parameter
2142 // variables. Resolve any parameter variables to the corresponding
2143 // argument passed to the current function. Leave global variables
2144 // as-is. Then add the mapped pair to the current function's list of
2145 // texture/sampler pairs.
2146 for (sem::VariablePair pair : func->TextureSamplerPairs()) {
2147 const sem::Variable* texture = pair.first;
2148 const sem::Variable* sampler = pair.second;
2149 if (auto* param = texture->As<sem::Parameter>()) {
2150 texture = args[param->Index()]->As<sem::VariableUser>()->Variable();
2151 }
2152 if (sampler) {
2153 if (auto* param = sampler->As<sem::Parameter>()) {
2154 sampler = args[param->Index()]->As<sem::VariableUser>()->Variable();
2155 }
2156 }
2157 current_function_->AddTextureSamplerPair(texture, sampler);
2158 }
2159}
2160
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002161sem::Expression* Resolver::Literal(const ast::LiteralExpression* literal) {
Ben Clayton8822e292022-05-04 22:18:49 +00002162 auto* ty = Switch(
2163 literal,
2164 [&](const ast::IntLiteralExpression* i) -> sem::Type* {
2165 switch (i->suffix) {
2166 case ast::IntLiteralExpression::Suffix::kNone:
Ben Clayton0d757d22022-06-01 20:34:40 +00002167 return builder_->create<sem::AbstractInt>();
Ben Clayton8822e292022-05-04 22:18:49 +00002168 case ast::IntLiteralExpression::Suffix::kI:
2169 return builder_->create<sem::I32>();
2170 case ast::IntLiteralExpression::Suffix::kU:
2171 return builder_->create<sem::U32>();
2172 }
2173 return nullptr;
2174 },
Ben Clayton65228372022-05-20 17:18:50 +00002175 [&](const ast::FloatLiteralExpression* f) -> sem::Type* {
Zhaoming Jiang0fb4e2c2022-06-10 18:18:35 +00002176 switch (f->suffix) {
2177 case ast::FloatLiteralExpression::Suffix::kNone:
2178 return builder_->create<sem::AbstractFloat>();
2179 case ast::FloatLiteralExpression::Suffix::kF:
2180 return builder_->create<sem::F32>();
2181 case ast::FloatLiteralExpression::Suffix::kH:
2182 return builder_->create<sem::F16>();
Ben Clayton65228372022-05-20 17:18:50 +00002183 }
Zhaoming Jiang0fb4e2c2022-06-10 18:18:35 +00002184 return nullptr;
Ben Clayton65228372022-05-20 17:18:50 +00002185 },
Ben Clayton8822e292022-05-04 22:18:49 +00002186 [&](const ast::BoolLiteralExpression*) { return builder_->create<sem::Bool>(); },
2187 [&](Default) { return nullptr; });
2188
2189 if (ty == nullptr) {
2190 TINT_UNREACHABLE(Resolver, builder_->Diagnostics())
2191 << "Unhandled literal type: " << literal->TypeInfo().name;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002192 return nullptr;
2193 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002194
Ben Claytondce63f52022-08-17 18:07:20 +00002195 if ((ty->Is<sem::F16>()) && (!enabled_extensions_.Contains(tint::ast::Extension::kF16))) {
Zhaoming Jiang0fb4e2c2022-06-10 18:18:35 +00002196 AddError("f16 literal used without 'f16' extension enabled", literal->source);
2197 return nullptr;
2198 }
2199
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00002200 const sem::Constant* val = nullptr;
2201 if (auto r = const_eval_.Literal(ty, literal)) {
2202 val = r.Get();
2203 } else {
2204 return nullptr;
2205 }
Ben Clayton83bd7382022-07-15 23:46:31 +00002206 return builder_->create<sem::Expression>(literal, ty, sem::EvaluationStage::kConstant,
2207 current_statement_, std::move(val),
dan sinclair41e4d9a2022-05-01 14:40:55 +00002208 /* has_side_effects */ false);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002209}
2210
2211sem::Expression* Resolver::Identifier(const ast::IdentifierExpression* expr) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002212 auto symbol = expr->symbol;
2213 auto* resolved = sem_.ResolvedSymbol(expr);
Ben Claytonf9031332022-06-26 11:36:10 +00002214 if (auto* variable = As<sem::Variable>(resolved)) {
2215 auto* user = builder_->create<sem::VariableUser>(expr, current_statement_, variable);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002216
dan sinclair41e4d9a2022-05-01 14:40:55 +00002217 if (current_statement_) {
2218 // If identifier is part of a loop continuing block, make sure it
2219 // doesn't refer to a variable that is bypassed by a continue statement
2220 // in the loop's body block.
2221 if (auto* continuing_block =
2222 current_statement_->FindFirstParent<sem::LoopContinuingBlockStatement>()) {
2223 auto* loop_block = continuing_block->FindFirstParent<sem::LoopBlockStatement>();
2224 if (loop_block->FirstContinue()) {
2225 auto& decls = loop_block->Decls();
2226 // If our identifier is in loop_block->decls, make sure its index is
2227 // less than first_continue
2228 auto iter = std::find_if(decls.begin(), decls.end(),
2229 [&symbol](auto* v) { return v->symbol == symbol; });
2230 if (iter != decls.end()) {
2231 auto var_decl_index =
2232 static_cast<size_t>(std::distance(decls.begin(), iter));
2233 if (var_decl_index >= loop_block->NumDeclsAtFirstContinue()) {
2234 AddError("continue statement bypasses declaration of '" +
2235 builder_->Symbols().NameFor(symbol) + "'",
2236 loop_block->FirstContinue()->source);
2237 AddNote("identifier '" + builder_->Symbols().NameFor(symbol) +
2238 "' declared here",
2239 (*iter)->source);
2240 AddNote("identifier '" + builder_->Symbols().NameFor(symbol) +
2241 "' referenced in continuing block here",
2242 expr->source);
2243 return nullptr;
2244 }
2245 }
2246 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002247 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002248 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002249
2250 if (current_function_) {
Ben Claytonf9031332022-06-26 11:36:10 +00002251 if (auto* global = variable->As<sem::GlobalVariable>()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002252 current_function_->AddDirectlyReferencedGlobal(global);
2253 }
Ben Claytonf9031332022-06-26 11:36:10 +00002254 } else if (variable->Declaration()->Is<ast::Var>()) {
2255 // Use of a module-scope 'var' outside of a function.
2256 // Note: The spec is currently vague around the rules here. See
2257 // https://github.com/gpuweb/gpuweb/issues/3081. Remove this comment when resolved.
2258 std::string desc = "var '" + builder_->Symbols().NameFor(symbol) + "' ";
dan sinclair1662f552022-09-21 14:44:33 +00002259 AddError(desc + "cannot be referenced at module-scope", expr->source);
Ben Claytonf9031332022-06-26 11:36:10 +00002260 AddNote(desc + "declared here", variable->Declaration()->source);
2261 return nullptr;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002262 }
2263
Ben Claytonf9031332022-06-26 11:36:10 +00002264 variable->AddUser(user);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002265 return user;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002266 }
2267
dan sinclair41e4d9a2022-05-01 14:40:55 +00002268 if (Is<sem::Function>(resolved)) {
2269 AddError("missing '(' for function call", expr->source.End());
2270 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002271 }
2272
dan sinclair41e4d9a2022-05-01 14:40:55 +00002273 if (IsBuiltin(symbol)) {
2274 AddError("missing '(' for builtin call", expr->source.End());
2275 return nullptr;
2276 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002277
dan sinclair41e4d9a2022-05-01 14:40:55 +00002278 if (resolved->Is<sem::Type>()) {
2279 AddError("missing '(' for type constructor or cast", expr->source.End());
2280 return nullptr;
2281 }
2282
2283 TINT_ICE(Resolver, diagnostics_)
2284 << expr->source << " unresolved identifier:\n"
2285 << "resolved: " << (resolved ? resolved->TypeInfo().name : "<null>") << "\n"
2286 << "name: " << builder_->Symbols().NameFor(symbol);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002287 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002288}
2289
dan sinclair41e4d9a2022-05-01 14:40:55 +00002290sem::Expression* Resolver::MemberAccessor(const ast::MemberAccessorExpression* expr) {
2291 auto* structure = sem_.TypeOf(expr->structure);
2292 auto* storage_ty = structure->UnwrapRef();
Ben Clayton6c098ba2022-07-14 20:46:39 +00002293 auto* object = sem_.Get(expr->structure);
2294 auto* source_var = object->SourceVariable();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002295
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002296 const sem::Type* ty = nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002297
Antonio Maioranodfeaf2902022-06-24 20:34:00 +00002298 // Object may be a side-effecting expression (e.g. function call).
Antonio Maioranodfeaf2902022-06-24 20:34:00 +00002299 bool has_side_effects = object && object->HasSideEffects();
Antonio Maiorano93e600d2022-03-15 13:51:17 +00002300
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002301 return Switch(
2302 storage_ty, //
2303 [&](const sem::Struct* str) -> sem::Expression* {
2304 Mark(expr->member);
2305 auto symbol = expr->member->symbol;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002306
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002307 const sem::StructMember* member = nullptr;
2308 for (auto* m : str->Members()) {
2309 if (m->Name() == symbol) {
2310 ty = m->Type();
2311 member = m;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002312 break;
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002313 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002314 }
2315
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002316 if (ty == nullptr) {
2317 AddError("struct member " + builder_->Symbols().NameFor(symbol) + " not found",
2318 expr->source);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002319 return nullptr;
2320 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002321
dan sinclair41e4d9a2022-05-01 14:40:55 +00002322 // If we're extracting from a reference, we return a reference.
2323 if (auto* ref = structure->As<sem::Reference>()) {
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002324 ty = builder_->create<sem::Reference>(ty, ref->StorageClass(), ref->Access());
dan sinclair41e4d9a2022-05-01 14:40:55 +00002325 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002326
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002327 auto val = const_eval_.MemberAccess(object, member);
2328 if (!val) {
2329 return nullptr;
2330 }
2331 return builder_->create<sem::StructMemberAccess>(expr, ty, current_statement_,
2332 val.Get(), object, member,
2333 has_side_effects, source_var);
2334 },
2335
2336 [&](const sem::Vector* vec) -> sem::Expression* {
2337 Mark(expr->member);
2338 std::string s = builder_->Symbols().NameFor(expr->member->symbol);
2339 auto size = s.size();
2340 utils::Vector<uint32_t, 4> swizzle;
2341 swizzle.Reserve(s.size());
2342
2343 for (auto c : s) {
2344 switch (c) {
2345 case 'x':
2346 case 'r':
2347 swizzle.Push(0u);
2348 break;
2349 case 'y':
2350 case 'g':
2351 swizzle.Push(1u);
2352 break;
2353 case 'z':
2354 case 'b':
2355 swizzle.Push(2u);
2356 break;
2357 case 'w':
2358 case 'a':
2359 swizzle.Push(3u);
2360 break;
2361 default:
2362 AddError("invalid vector swizzle character",
2363 expr->member->source.Begin() + swizzle.Length());
2364 return nullptr;
2365 }
2366
2367 if (swizzle.Back() >= vec->Width()) {
2368 AddError("invalid vector swizzle member", expr->member->source);
2369 return nullptr;
2370 }
2371 }
2372
2373 if (size < 1 || size > 4) {
2374 AddError("invalid vector swizzle size", expr->member->source);
2375 return nullptr;
2376 }
2377
2378 // All characters are valid, check if they're being mixed
2379 auto is_rgba = [](char c) { return c == 'r' || c == 'g' || c == 'b' || c == 'a'; };
2380 auto is_xyzw = [](char c) { return c == 'x' || c == 'y' || c == 'z' || c == 'w'; };
2381 if (!std::all_of(s.begin(), s.end(), is_rgba) &&
2382 !std::all_of(s.begin(), s.end(), is_xyzw)) {
2383 AddError("invalid mixing of vector swizzle characters rgba with xyzw",
2384 expr->member->source);
2385 return nullptr;
2386 }
2387
2388 if (size == 1) {
2389 // A single element swizzle is just the type of the vector.
2390 ty = vec->type();
2391 // If we're extracting from a reference, we return a reference.
2392 if (auto* ref = structure->As<sem::Reference>()) {
2393 ty = builder_->create<sem::Reference>(ty, ref->StorageClass(), ref->Access());
2394 }
2395 } else {
2396 // The vector will have a number of components equal to the length of
2397 // the swizzle.
2398 ty = builder_->create<sem::Vector>(vec->type(), static_cast<uint32_t>(size));
2399 }
2400 auto val = const_eval_.Swizzle(ty, object, swizzle);
2401 if (!val) {
2402 return nullptr;
2403 }
2404 return builder_->create<sem::Swizzle>(expr, ty, current_statement_, val.Get(), object,
2405 std::move(swizzle), has_side_effects, source_var);
2406 },
2407
2408 [&](Default) {
2409 AddError("invalid member accessor expression. Expected vector or struct, got '" +
2410 sem_.TypeNameOf(storage_ty) + "'",
2411 expr->structure->source);
2412 return nullptr;
2413 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002414}
2415
2416sem::Expression* Resolver::Binary(const ast::BinaryExpression* expr) {
Ben Clayton7b921fb2022-05-19 20:19:49 +00002417 const auto* lhs = sem_.Get(expr->lhs);
2418 const auto* rhs = sem_.Get(expr->rhs);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002419 auto* lhs_ty = lhs->Type()->UnwrapRef();
2420 auto* rhs_ty = rhs->Type()->UnwrapRef();
James Pricedaea0342022-03-31 22:30:10 +00002421
Ben Clayton7b921fb2022-05-19 20:19:49 +00002422 auto op = intrinsic_table_->Lookup(expr->op, lhs_ty, rhs_ty, expr->source, false);
2423 if (!op.result) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002424 return nullptr;
2425 }
Ben Clayton932418e2022-05-24 21:42:03 +00002426 if (ShouldMaterializeArgument(op.lhs)) {
2427 lhs = Materialize(lhs, op.lhs);
2428 if (!lhs) {
2429 return nullptr;
2430 }
2431 }
2432 if (ShouldMaterializeArgument(op.rhs)) {
2433 rhs = Materialize(rhs, op.rhs);
2434 if (!rhs) {
2435 return nullptr;
2436 }
2437 }
James Pricedaea0342022-03-31 22:30:10 +00002438
Ben Clayton65c5c9d2022-07-15 14:14:09 +00002439 const sem::Constant* value = nullptr;
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00002440 auto stage = sem::EarliestStage(lhs->Stage(), rhs->Stage());
2441 if (stage == sem::EvaluationStage::kConstant) {
2442 if (op.const_eval_fn) {
Antonio Maiorano7058a172022-08-10 18:06:43 +00002443 auto const_args = utils::Vector{lhs->ConstantValue(), rhs->ConstantValue()};
Antonio Maioranoa58d8c92022-08-10 20:01:17 +00002444 // Implicit conversion (e.g. AInt -> AFloat)
Antonio Maiorano2de4dee2022-08-16 16:00:02 +00002445 if (!Convert(const_args[0], op.lhs, lhs->Declaration()->source)) {
Antonio Maioranoa58d8c92022-08-10 20:01:17 +00002446 return nullptr;
2447 }
Antonio Maiorano2de4dee2022-08-16 16:00:02 +00002448 if (!Convert(const_args[1], op.rhs, rhs->Declaration()->source)) {
Antonio Maioranoa58d8c92022-08-10 20:01:17 +00002449 return nullptr;
2450 }
2451
Antonio Maiorano7058a172022-08-10 18:06:43 +00002452 if (auto r = (const_eval_.*op.const_eval_fn)(op.result, const_args, expr->source)) {
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00002453 value = r.Get();
2454 } else {
2455 return nullptr;
2456 }
2457 } else {
2458 stage = sem::EvaluationStage::kRuntime;
2459 }
Ben Clayton65c5c9d2022-07-15 14:14:09 +00002460 }
2461
dan sinclair41e4d9a2022-05-01 14:40:55 +00002462 bool has_side_effects = lhs->HasSideEffects() || rhs->HasSideEffects();
Ben Clayton83bd7382022-07-15 23:46:31 +00002463 auto* sem = builder_->create<sem::Expression>(expr, op.result, stage, current_statement_, value,
Ben Claytond5f53ab2022-07-07 17:30:11 +00002464 has_side_effects);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002465 sem->Behaviors() = lhs->Behaviors() + rhs->Behaviors();
James Pricedaea0342022-03-31 22:30:10 +00002466
dan sinclair41e4d9a2022-05-01 14:40:55 +00002467 return sem;
James Pricedaea0342022-03-31 22:30:10 +00002468}
2469
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002470sem::Expression* Resolver::UnaryOp(const ast::UnaryOpExpression* unary) {
Ben Clayton7b921fb2022-05-19 20:19:49 +00002471 const auto* expr = sem_.Get(unary->expr);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002472 auto* expr_ty = expr->Type();
2473 if (!expr_ty) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002474 return nullptr;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002475 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002476
dan sinclair41e4d9a2022-05-01 14:40:55 +00002477 const sem::Type* ty = nullptr;
2478 const sem::Variable* source_var = nullptr;
Ben Clayton65c5c9d2022-07-15 14:14:09 +00002479 const sem::Constant* value = nullptr;
Ben Claytonac660c22022-07-15 23:54:10 +00002480 auto stage = sem::EvaluationStage::kRuntime;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002481
dan sinclair41e4d9a2022-05-01 14:40:55 +00002482 switch (unary->op) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002483 case ast::UnaryOp::kAddressOf:
2484 if (auto* ref = expr_ty->As<sem::Reference>()) {
2485 if (ref->StoreType()->UnwrapRef()->is_handle()) {
2486 AddError("cannot take the address of expression in handle storage class",
2487 unary->expr->source);
2488 return nullptr;
2489 }
James Price98eb1692022-04-22 17:40:33 +00002490
dan sinclair41e4d9a2022-05-01 14:40:55 +00002491 auto* array = unary->expr->As<ast::IndexAccessorExpression>();
2492 auto* member = unary->expr->As<ast::MemberAccessorExpression>();
2493 if ((array && sem_.TypeOf(array->object)->UnwrapRef()->Is<sem::Vector>()) ||
2494 (member && sem_.TypeOf(member->structure)->UnwrapRef()->Is<sem::Vector>())) {
2495 AddError("cannot take the address of a vector component", unary->expr->source);
2496 return nullptr;
2497 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002498
dan sinclair41e4d9a2022-05-01 14:40:55 +00002499 ty = builder_->create<sem::Pointer>(ref->StoreType(), ref->StorageClass(),
2500 ref->Access());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002501
dan sinclair41e4d9a2022-05-01 14:40:55 +00002502 source_var = expr->SourceVariable();
2503 } else {
2504 AddError("cannot take the address of expression", unary->expr->source);
2505 return nullptr;
2506 }
2507 break;
2508
2509 case ast::UnaryOp::kIndirection:
2510 if (auto* ptr = expr_ty->As<sem::Pointer>()) {
2511 ty = builder_->create<sem::Reference>(ptr->StoreType(), ptr->StorageClass(),
2512 ptr->Access());
2513 source_var = expr->SourceVariable();
2514 } else {
2515 AddError("cannot dereference expression of type '" + sem_.TypeNameOf(expr_ty) + "'",
2516 unary->expr->source);
2517 return nullptr;
2518 }
2519 break;
Ben Claytonb61e0452022-05-09 21:22:24 +00002520
2521 default: {
Ben Clayton932418e2022-05-24 21:42:03 +00002522 auto op = intrinsic_table_->Lookup(unary->op, expr_ty, unary->source);
2523 if (!op.result) {
Ben Claytonb61e0452022-05-09 21:22:24 +00002524 return nullptr;
2525 }
Antonio Maiorano29aa6132022-09-07 19:34:44 +00002526 ty = op.result;
Ben Clayton932418e2022-05-24 21:42:03 +00002527 if (ShouldMaterializeArgument(op.parameter)) {
2528 expr = Materialize(expr, op.parameter);
2529 if (!expr) {
2530 return nullptr;
2531 }
2532 }
Ben Claytonac660c22022-07-15 23:54:10 +00002533 stage = expr->Stage();
2534 if (stage == sem::EvaluationStage::kConstant) {
2535 if (op.const_eval_fn) {
Antonio Maiorano7058a172022-08-10 18:06:43 +00002536 if (auto r = (const_eval_.*op.const_eval_fn)(
2537 ty, utils::Vector{expr->ConstantValue()},
2538 expr->Declaration()->source)) {
Antonio Maioranoc2a052e2022-08-04 13:59:36 +00002539 value = r.Get();
2540 } else {
2541 return nullptr;
2542 }
Ben Claytonac660c22022-07-15 23:54:10 +00002543 } else {
2544 stage = sem::EvaluationStage::kRuntime;
2545 }
Ben Clayton65c5c9d2022-07-15 14:14:09 +00002546 }
Ben Clayton7b921fb2022-05-19 20:19:49 +00002547 break;
Ben Claytonb61e0452022-05-09 21:22:24 +00002548 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002549 }
2550
Ben Clayton83bd7382022-07-15 23:46:31 +00002551 auto* sem = builder_->create<sem::Expression>(unary, ty, stage, current_statement_, value,
dan sinclair41e4d9a2022-05-01 14:40:55 +00002552 expr->HasSideEffects(), source_var);
2553 sem->Behaviors() = expr->Behaviors();
2554 return sem;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002555}
2556
Ben Clayton7f2b8cd2022-05-18 22:41:48 +00002557bool Resolver::Enable(const ast::Enable* enable) {
Ben Claytondce63f52022-08-17 18:07:20 +00002558 enabled_extensions_.Add(enable->extension);
Ben Clayton7f2b8cd2022-05-18 22:41:48 +00002559 return true;
2560}
2561
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002562sem::Type* Resolver::TypeDecl(const ast::TypeDecl* named_type) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002563 sem::Type* result = nullptr;
2564 if (auto* alias = named_type->As<ast::Alias>()) {
2565 result = Alias(alias);
2566 } else if (auto* str = named_type->As<ast::Struct>()) {
2567 result = Structure(str);
2568 } else {
2569 TINT_UNREACHABLE(Resolver, diagnostics_) << "Unhandled TypeDecl";
2570 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002571
dan sinclair41e4d9a2022-05-01 14:40:55 +00002572 if (!result) {
2573 return nullptr;
2574 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002575
dan sinclair41e4d9a2022-05-01 14:40:55 +00002576 builder_->Sem().Add(named_type, result);
2577 return result;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002578}
2579
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002580sem::Array* Resolver::Array(const ast::Array* arr) {
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002581 if (!arr->type) {
2582 AddError("missing array element type", arr->source.End());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002583 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002584 }
2585
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002586 auto* el_ty = Type(arr->type);
2587 if (!el_ty) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002588 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002589 }
2590
dan sinclair41e4d9a2022-05-01 14:40:55 +00002591 // Look for explicit stride via @stride(n) attribute
2592 uint32_t explicit_stride = 0;
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002593 if (!ArrayAttributes(arr->attributes, el_ty, explicit_stride)) {
2594 return nullptr;
2595 }
2596
dan sinclair78f80672022-09-22 22:28:21 +00002597 sem::ArrayCount el_count = sem::RuntimeArrayCount{};
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002598
2599 // Evaluate the constant array size expression.
2600 if (auto* count_expr = arr->count) {
2601 if (auto count = ArrayCount(count_expr)) {
2602 el_count = count.Get();
2603 } else {
2604 return nullptr;
2605 }
2606 }
2607
dan sinclair78f80672022-09-22 22:28:21 +00002608 auto* out = Array(arr->type->source, //
2609 arr->count ? arr->count->source : arr->source, //
2610 el_ty, el_count, explicit_stride);
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002611 if (out == nullptr) {
2612 return nullptr;
2613 }
2614
2615 if (el_ty->Is<sem::Atomic>()) {
2616 atomic_composite_info_.emplace(out, arr->type->source);
2617 } else {
2618 auto found = atomic_composite_info_.find(el_ty);
2619 if (found != atomic_composite_info_.end()) {
2620 atomic_composite_info_.emplace(out, found->second);
2621 }
2622 }
2623
2624 return out;
2625}
2626
dan sinclair78f80672022-09-22 22:28:21 +00002627utils::Result<sem::ArrayCount> Resolver::ArrayCount(const ast::Expression* count_expr) {
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002628 // Evaluate the constant array size expression.
2629 const auto* count_sem = Materialize(Expression(count_expr));
2630 if (!count_sem) {
2631 return utils::Failure;
2632 }
2633
dan sinclair78f80672022-09-22 22:28:21 +00002634 // Note: If the array count is an 'override', but not a identifier expression, we do not return
2635 // here, but instead continue to the ConstantValue() check below.
2636 if (auto* user = count_sem->UnwrapMaterialize()->As<sem::VariableUser>()) {
2637 if (auto* global = user->Variable()->As<sem::GlobalVariable>()) {
2638 if (global->Declaration()->Is<ast::Override>()) {
2639 return sem::ArrayCount{sem::OverrideArrayCount{global}};
2640 }
2641 }
2642 }
2643
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002644 auto* count_val = count_sem->ConstantValue();
2645 if (!count_val) {
dan sinclair78f80672022-09-22 22:28:21 +00002646 AddError("array size must evaluate to a constant integer expression or override variable",
2647 count_expr->source);
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002648 return utils::Failure;
2649 }
2650
2651 if (auto* ty = count_val->Type(); !ty->is_integer_scalar()) {
2652 AddError("array size must evaluate to a constant integer expression, but is type '" +
2653 builder_->FriendlyName(ty) + "'",
2654 count_expr->source);
2655 return utils::Failure;
2656 }
2657
2658 int64_t count = count_val->As<AInt>();
2659 if (count < 1) {
2660 AddError("array size (" + std::to_string(count) + ") must be greater than 0",
2661 count_expr->source);
2662 return utils::Failure;
2663 }
2664
dan sinclair78f80672022-09-22 22:28:21 +00002665 return sem::ArrayCount{sem::ConstantArrayCount{static_cast<uint32_t>(count)}};
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002666}
2667
Ben Clayton783b1692022-08-02 17:03:35 +00002668bool Resolver::ArrayAttributes(utils::VectorRef<const ast::Attribute*> attributes,
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002669 const sem::Type* el_ty,
2670 uint32_t& explicit_stride) {
2671 if (!validator_.NoDuplicateAttributes(attributes)) {
2672 return false;
2673 }
2674
2675 for (auto* attr : attributes) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002676 Mark(attr);
2677 if (auto* sd = attr->As<ast::StrideAttribute>()) {
2678 explicit_stride = sd->stride;
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002679 if (!validator_.ArrayStrideAttribute(sd, el_ty->Size(), el_ty->Align())) {
2680 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002681 }
2682 continue;
2683 }
2684
2685 AddError("attribute is not valid for array types", attr->source);
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002686 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002687 }
2688
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002689 return true;
2690}
dan sinclair41e4d9a2022-05-01 14:40:55 +00002691
dan sinclair78f80672022-09-22 22:28:21 +00002692sem::Array* Resolver::Array(const Source& el_source,
2693 const Source& count_source,
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002694 const sem::Type* el_ty,
dan sinclair78f80672022-09-22 22:28:21 +00002695 sem::ArrayCount el_count,
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002696 uint32_t explicit_stride) {
2697 uint32_t el_align = el_ty->Align();
2698 uint32_t el_size = el_ty->Size();
2699 uint64_t implicit_stride = el_size ? utils::RoundUp<uint64_t>(el_align, el_size) : 0;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002700 uint64_t stride = explicit_stride ? explicit_stride : implicit_stride;
dan sinclair78f80672022-09-22 22:28:21 +00002701 uint64_t size = 0;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002702
dan sinclair78f80672022-09-22 22:28:21 +00002703 if (auto const_count = std::get_if<sem::ConstantArrayCount>(&el_count)) {
2704 size = const_count->value * stride;
2705 if (size > std::numeric_limits<uint32_t>::max()) {
2706 std::stringstream msg;
2707 msg << "array size (0x" << std::hex << size << ") must not exceed 0xffffffff bytes";
2708 AddError(msg.str(), count_source);
2709 return nullptr;
2710 }
2711 } else if (std::holds_alternative<sem::RuntimeArrayCount>(el_count)) {
2712 size = stride;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002713 }
Ben Clayton9c4d0c92022-07-31 19:33:04 +00002714 auto* out = builder_->create<sem::Array>(el_ty, el_count, el_align, static_cast<uint32_t>(size),
2715 static_cast<uint32_t>(stride),
2716 static_cast<uint32_t>(implicit_stride));
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002717
dan sinclair78f80672022-09-22 22:28:21 +00002718 if (!validator_.Array(out, el_source)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002719 return nullptr;
2720 }
2721
dan sinclair41e4d9a2022-05-01 14:40:55 +00002722 return out;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002723}
2724
2725sem::Type* Resolver::Alias(const ast::Alias* alias) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002726 auto* ty = Type(alias->type);
2727 if (!ty) {
2728 return nullptr;
2729 }
2730 if (!validator_.Alias(alias)) {
2731 return nullptr;
2732 }
2733 return ty;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002734}
2735
2736sem::Struct* Resolver::Structure(const ast::Struct* str) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002737 if (!validator_.NoDuplicateAttributes(str->attributes)) {
2738 return nullptr;
2739 }
2740 for (auto* attr : str->attributes) {
2741 Mark(attr);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002742 }
2743
dan sinclair41e4d9a2022-05-01 14:40:55 +00002744 sem::StructMemberList sem_members;
Ben Clayton783b1692022-08-02 17:03:35 +00002745 sem_members.reserve(str->members.Length());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002746
dan sinclair41e4d9a2022-05-01 14:40:55 +00002747 // Calculate the effective size and alignment of each field, and the overall
2748 // size of the structure.
2749 // For size, use the size attribute if provided, otherwise use the default
2750 // size for the type.
2751 // For alignment, use the alignment attribute if provided, otherwise use the
2752 // default alignment for the member type.
2753 // Diagnostic errors are raised if a basic rule is violated.
2754 // Validation of storage-class rules requires analysing the actual variable
2755 // usage of the structure, and so is performed as part of the variable
2756 // validation.
2757 uint64_t struct_size = 0;
2758 uint64_t struct_align = 1;
2759 std::unordered_map<Symbol, const ast::StructMember*> member_map;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002760
dan sinclair41e4d9a2022-05-01 14:40:55 +00002761 for (auto* member : str->members) {
2762 Mark(member);
2763 auto result = member_map.emplace(member->symbol, member);
2764 if (!result.second) {
2765 AddError("redefinition of '" + builder_->Symbols().NameFor(member->symbol) + "'",
2766 member->source);
2767 AddNote("previous definition is here", result.first->second->source);
2768 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002769 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002770
2771 // Resolve member type
2772 auto* type = Type(member->type);
2773 if (!type) {
2774 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002775 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002776
2777 // validator_.Validate member type
2778 if (!validator_.IsPlain(type)) {
2779 AddError(sem_.TypeNameOf(type) + " cannot be used as the type of a structure member",
2780 member->source);
2781 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002782 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002783
2784 uint64_t offset = struct_size;
2785 uint64_t align = type->Align();
2786 uint64_t size = type->Size();
2787
2788 if (!validator_.NoDuplicateAttributes(member->attributes)) {
2789 return nullptr;
2790 }
2791
2792 bool has_offset_attr = false;
2793 bool has_align_attr = false;
2794 bool has_size_attr = false;
dan sinclair97c9e0b2022-09-06 14:41:16 +00002795 std::optional<uint32_t> location;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002796 for (auto* attr : member->attributes) {
2797 Mark(attr);
2798 if (auto* o = attr->As<ast::StructMemberOffsetAttribute>()) {
2799 // Offset attributes are not part of the WGSL spec, but are emitted
2800 // by the SPIR-V reader.
dan sinclair93df9672022-09-09 14:49:09 +00002801
2802 auto* materialized = Materialize(Expression(o->expr));
2803 if (!materialized) {
2804 return nullptr;
2805 }
2806 auto const_value = materialized->ConstantValue();
2807 if (!const_value) {
2808 AddError("'offset' must be constant expression", o->expr->source);
2809 return nullptr;
2810 }
2811 offset = const_value->As<uint64_t>();
2812
2813 if (offset < struct_size) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002814 AddError("offsets must be in ascending order", o->source);
2815 return nullptr;
2816 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002817 align = 1;
2818 has_offset_attr = true;
2819 } else if (auto* a = attr->As<ast::StructMemberAlignAttribute>()) {
dan sinclairc4e076f2022-09-08 01:04:34 +00002820 auto* materialized = Materialize(Expression(a->expr));
dan sinclair4964d9b2022-08-23 13:28:44 +00002821 if (!materialized) {
2822 return nullptr;
2823 }
2824 auto const_value = materialized->ConstantValue();
2825 if (!const_value) {
dan sinclairc4e076f2022-09-08 01:04:34 +00002826 AddError("'align' must be constant expression", a->expr->source);
dan sinclair4964d9b2022-08-23 13:28:44 +00002827 return nullptr;
2828 }
2829 auto value = const_value->As<AInt>();
2830
2831 if (value <= 0 || !utils::IsPowerOfTwo(value)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002832 AddError("align value must be a positive, power-of-two integer", a->source);
2833 return nullptr;
2834 }
dan sinclair4964d9b2022-08-23 13:28:44 +00002835 align = const_value->As<u32>();
dan sinclair41e4d9a2022-05-01 14:40:55 +00002836 has_align_attr = true;
2837 } else if (auto* s = attr->As<ast::StructMemberSizeAttribute>()) {
dan sinclair93df9672022-09-09 14:49:09 +00002838 auto* materialized = Materialize(Expression(s->expr));
2839 if (!materialized) {
2840 return nullptr;
2841 }
2842 auto const_value = materialized->ConstantValue();
2843 if (!const_value) {
2844 AddError("'size' must be constant expression", s->expr->source);
2845 return nullptr;
2846 }
2847 auto value = const_value->As<uint64_t>();
2848
2849 if (value < size) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002850 AddError("size must be at least as big as the type's size (" +
2851 std::to_string(size) + ")",
2852 s->source);
2853 return nullptr;
2854 }
dan sinclair93df9672022-09-09 14:49:09 +00002855 size = const_value->As<u32>();
dan sinclair41e4d9a2022-05-01 14:40:55 +00002856 has_size_attr = true;
dan sinclair97c9e0b2022-09-06 14:41:16 +00002857 } else if (auto* l = attr->As<ast::LocationAttribute>()) {
dan sinclairc4e076f2022-09-08 01:04:34 +00002858 auto* materialize = Materialize(Expression(l->expr));
dan sinclairf9eeed62022-09-07 22:25:24 +00002859 if (!materialize) {
2860 return nullptr;
2861 }
2862 auto* c = materialize->ConstantValue();
2863 if (!c) {
2864 // TODO(crbug.com/tint/1633): Add error message about invalid materialization
2865 // when location can be an expression.
2866 return nullptr;
2867 }
2868 location = c->As<uint32_t>();
dan sinclair41e4d9a2022-05-01 14:40:55 +00002869 }
2870 }
2871
2872 if (has_offset_attr && (has_align_attr || has_size_attr)) {
2873 AddError("offset attributes cannot be used with align or size attributes",
2874 member->source);
2875 return nullptr;
2876 }
2877
2878 offset = utils::RoundUp(align, offset);
2879 if (offset > std::numeric_limits<uint32_t>::max()) {
2880 std::stringstream msg;
Ben Claytonb4ff8c82022-06-24 23:30:59 +00002881 msg << "struct member offset (0x" << std::hex << offset << ") must not exceed 0x"
2882 << std::hex << std::numeric_limits<uint32_t>::max() << " bytes";
dan sinclair41e4d9a2022-05-01 14:40:55 +00002883 AddError(msg.str(), member->source);
2884 return nullptr;
2885 }
2886
2887 auto* sem_member = builder_->create<sem::StructMember>(
2888 member, member->symbol, type, static_cast<uint32_t>(sem_members.size()),
2889 static_cast<uint32_t>(offset), static_cast<uint32_t>(align),
dan sinclair97c9e0b2022-09-06 14:41:16 +00002890 static_cast<uint32_t>(size), location);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002891 builder_->Sem().Add(member, sem_member);
2892 sem_members.emplace_back(sem_member);
2893
2894 struct_size = offset + size;
2895 struct_align = std::max(struct_align, align);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002896 }
2897
dan sinclair41e4d9a2022-05-01 14:40:55 +00002898 uint64_t size_no_padding = struct_size;
2899 struct_size = utils::RoundUp(struct_align, struct_size);
2900
2901 if (struct_size > std::numeric_limits<uint32_t>::max()) {
2902 std::stringstream msg;
Ben Claytonb4ff8c82022-06-24 23:30:59 +00002903 msg << "struct size (0x" << std::hex << struct_size << ") must not exceed 0xffffffff bytes";
dan sinclair41e4d9a2022-05-01 14:40:55 +00002904 AddError(msg.str(), str->source);
2905 return nullptr;
2906 }
2907 if (struct_align > std::numeric_limits<uint32_t>::max()) {
2908 TINT_ICE(Resolver, diagnostics_) << "calculated struct stride exceeds uint32";
2909 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002910 }
2911
dan sinclair41e4d9a2022-05-01 14:40:55 +00002912 auto* out = builder_->create<sem::Struct>(
2913 str, str->name, sem_members, static_cast<uint32_t>(struct_align),
2914 static_cast<uint32_t>(struct_size), static_cast<uint32_t>(size_no_padding));
2915
2916 for (size_t i = 0; i < sem_members.size(); i++) {
2917 auto* mem_type = sem_members[i]->Type();
2918 if (mem_type->Is<sem::Atomic>()) {
2919 atomic_composite_info_.emplace(out, sem_members[i]->Declaration()->source);
2920 break;
2921 } else {
2922 auto found = atomic_composite_info_.find(mem_type);
2923 if (found != atomic_composite_info_.end()) {
2924 atomic_composite_info_.emplace(out, found->second);
2925 break;
2926 }
2927 }
Antonio Maioranodfeaf2902022-06-24 20:34:00 +00002928
2929 const_cast<sem::StructMember*>(sem_members[i])->SetStruct(out);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002930 }
2931
dan sinclair41e4d9a2022-05-01 14:40:55 +00002932 auto stage = current_function_ ? current_function_->Declaration()->PipelineStage()
2933 : ast::PipelineStage::kNone;
2934 if (!validator_.Structure(out, stage)) {
2935 return nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002936 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002937
dan sinclair41e4d9a2022-05-01 14:40:55 +00002938 return out;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002939}
2940
2941sem::Statement* Resolver::ReturnStatement(const ast::ReturnStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002942 auto* sem =
2943 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
2944 return StatementScope(stmt, sem, [&] {
2945 auto& behaviors = current_statement_->Behaviors();
2946 behaviors = sem::Behavior::kReturn;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002947
Ben Clayton7b921fb2022-05-19 20:19:49 +00002948 const sem::Type* value_ty = nullptr;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002949 if (auto* value = stmt->value) {
Ben Clayton33fe68e2022-06-06 15:11:04 +00002950 const auto* expr = Expression(value);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002951 if (!expr) {
2952 return false;
2953 }
Ben Clayton33fe68e2022-06-06 15:11:04 +00002954 if (auto* ret_ty = current_function_->ReturnType(); !ret_ty->Is<sem::Void>()) {
2955 expr = Materialize(expr, ret_ty);
2956 if (!expr) {
2957 return false;
2958 }
2959 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002960 behaviors.Add(expr->Behaviors() - sem::Behavior::kNext);
Ben Clayton7b921fb2022-05-19 20:19:49 +00002961 value_ty = expr->Type()->UnwrapRef();
2962 } else {
2963 value_ty = builder_->create<sem::Void>();
dan sinclair41e4d9a2022-05-01 14:40:55 +00002964 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002965
dan sinclair41e4d9a2022-05-01 14:40:55 +00002966 // Validate after processing the return value expression so that its type
2967 // is available for validation.
Ben Clayton7b921fb2022-05-19 20:19:49 +00002968 return validator_.Return(stmt, current_function_->ReturnType(), value_ty,
dan sinclair41e4d9a2022-05-01 14:40:55 +00002969 current_statement_);
2970 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002971}
2972
dan sinclair41e4d9a2022-05-01 14:40:55 +00002973sem::SwitchStatement* Resolver::SwitchStatement(const ast::SwitchStatement* stmt) {
2974 auto* sem = builder_->create<sem::SwitchStatement>(stmt, current_compound_statement_,
2975 current_function_);
2976 return StatementScope(stmt, sem, [&] {
2977 auto& behaviors = sem->Behaviors();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002978
Ben Clayton932418e2022-05-24 21:42:03 +00002979 const auto* cond = Expression(stmt->condition);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002980 if (!cond) {
2981 return false;
2982 }
2983 behaviors = cond->Behaviors() - sem::Behavior::kNext;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002984
Ben Clayton932418e2022-05-24 21:42:03 +00002985 auto* cond_ty = cond->Type()->UnwrapRef();
2986
Ben Clayton8e0368e2022-07-29 14:15:01 +00002987 utils::Vector<const sem::Type*, 8> types;
2988 types.Push(cond_ty);
Ben Clayton932418e2022-05-24 21:42:03 +00002989
Ben Clayton783b1692022-08-02 17:03:35 +00002990 utils::Vector<sem::CaseStatement*, 4> cases;
2991 cases.Reserve(stmt->body.Length());
dan sinclair41e4d9a2022-05-01 14:40:55 +00002992 for (auto* case_stmt : stmt->body) {
2993 Mark(case_stmt);
2994 auto* c = CaseStatement(case_stmt);
2995 if (!c) {
2996 return false;
2997 }
Ben Clayton932418e2022-05-24 21:42:03 +00002998 for (auto* expr : c->Selectors()) {
Ben Clayton8e0368e2022-07-29 14:15:01 +00002999 types.Push(expr->Type()->UnwrapRef());
Ben Clayton932418e2022-05-24 21:42:03 +00003000 }
Ben Clayton783b1692022-08-02 17:03:35 +00003001 cases.Push(c);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003002 behaviors.Add(c->Behaviors());
Ben Clayton43581f12022-05-20 12:28:00 +00003003 sem->Cases().emplace_back(c);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003004 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003005
Ben Clayton932418e2022-05-24 21:42:03 +00003006 // Determine the common type across all selectors and the switch expression
3007 // This must materialize to an integer scalar (non-abstract).
Ben Clayton8e0368e2022-07-29 14:15:01 +00003008 auto* common_ty = sem::Type::Common(types);
Ben Clayton932418e2022-05-24 21:42:03 +00003009 if (!common_ty || !common_ty->is_integer_scalar()) {
3010 // No common type found or the common type was abstract.
3011 // Pick i32 and let validation deal with any mismatches.
3012 common_ty = builder_->create<sem::I32>();
3013 }
3014 cond = Materialize(cond, common_ty);
3015 if (!cond) {
3016 return false;
3017 }
3018 for (auto* c : cases) {
3019 for (auto*& sel : c->Selectors()) { // Note: pointer reference
3020 sel = Materialize(sel, common_ty);
3021 if (!sel) {
3022 return false;
3023 }
3024 }
3025 }
3026
dan sinclair41e4d9a2022-05-01 14:40:55 +00003027 if (behaviors.Contains(sem::Behavior::kBreak)) {
3028 behaviors.Add(sem::Behavior::kNext);
3029 }
3030 behaviors.Remove(sem::Behavior::kBreak, sem::Behavior::kFallthrough);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003031
dan sinclair41e4d9a2022-05-01 14:40:55 +00003032 return validator_.SwitchStatement(stmt);
3033 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003034}
3035
dan sinclair41e4d9a2022-05-01 14:40:55 +00003036sem::Statement* Resolver::VariableDeclStatement(const ast::VariableDeclStatement* stmt) {
3037 auto* sem =
3038 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3039 return StatementScope(stmt, sem, [&] {
3040 Mark(stmt->variable);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003041
Ben Claytone4e48542022-06-20 19:17:52 +00003042 auto* variable = Variable(stmt->variable, /* is_global */ false);
Ben Claytonee49b1e2022-06-20 15:30:41 +00003043 if (!variable) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003044 return false;
3045 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003046
dan sinclair41e4d9a2022-05-01 14:40:55 +00003047 for (auto* attr : stmt->variable->attributes) {
3048 Mark(attr);
3049 if (!attr->Is<ast::InternalAttribute>()) {
3050 AddError("attributes are not valid on local variables", attr->source);
3051 return false;
3052 }
3053 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003054
dan sinclair41e4d9a2022-05-01 14:40:55 +00003055 if (current_block_) { // Not all statements are inside a block
3056 current_block_->AddDecl(stmt->variable);
3057 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003058
Ben Claytonee49b1e2022-06-20 15:30:41 +00003059 if (auto* ctor = variable->Constructor()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003060 sem->Behaviors() = ctor->Behaviors();
3061 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003062
Ben Claytoncfe07a12022-07-15 13:01:49 +00003063 return validator_.LocalVariable(variable);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003064 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003065}
3066
dan sinclair41e4d9a2022-05-01 14:40:55 +00003067sem::Statement* Resolver::AssignmentStatement(const ast::AssignmentStatement* stmt) {
3068 auto* sem =
3069 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3070 return StatementScope(stmt, sem, [&] {
3071 auto* lhs = Expression(stmt->lhs);
3072 if (!lhs) {
3073 return false;
3074 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003075
Ben Clayton22bd0042022-05-31 20:40:59 +00003076 const bool is_phony_assignment = stmt->lhs->Is<ast::PhonyExpression>();
3077
3078 const auto* rhs = Expression(stmt->rhs);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003079 if (!rhs) {
3080 return false;
3081 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003082
Ben Clayton22bd0042022-05-31 20:40:59 +00003083 if (!is_phony_assignment) {
3084 rhs = Materialize(rhs, lhs->Type()->UnwrapRef());
3085 if (!rhs) {
3086 return false;
3087 }
3088 }
3089
dan sinclair41e4d9a2022-05-01 14:40:55 +00003090 auto& behaviors = sem->Behaviors();
3091 behaviors = rhs->Behaviors();
Ben Clayton22bd0042022-05-31 20:40:59 +00003092 if (!is_phony_assignment) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003093 behaviors.Add(lhs->Behaviors());
3094 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003095
dan sinclair41e4d9a2022-05-01 14:40:55 +00003096 return validator_.Assignment(stmt, sem_.TypeOf(stmt->rhs));
3097 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003098}
3099
3100sem::Statement* Resolver::BreakStatement(const ast::BreakStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003101 auto* sem =
3102 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3103 return StatementScope(stmt, sem, [&] {
3104 sem->Behaviors() = sem::Behavior::kBreak;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003105
dan sinclair41e4d9a2022-05-01 14:40:55 +00003106 return validator_.BreakStatement(sem, current_statement_);
3107 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003108}
3109
3110sem::Statement* Resolver::CallStatement(const ast::CallStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003111 auto* sem =
3112 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3113 return StatementScope(stmt, sem, [&] {
3114 if (auto* expr = Expression(stmt->expr)) {
3115 sem->Behaviors() = expr->Behaviors();
3116 return true;
3117 }
3118 return false;
3119 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003120}
3121
James Pricec022ff52022-03-31 22:30:10 +00003122sem::Statement* Resolver::CompoundAssignmentStatement(
3123 const ast::CompoundAssignmentStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003124 auto* sem =
3125 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3126 return StatementScope(stmt, sem, [&] {
3127 auto* lhs = Expression(stmt->lhs);
3128 if (!lhs) {
3129 return false;
3130 }
James Pricec022ff52022-03-31 22:30:10 +00003131
dan sinclair41e4d9a2022-05-01 14:40:55 +00003132 auto* rhs = Expression(stmt->rhs);
3133 if (!rhs) {
3134 return false;
3135 }
James Pricec022ff52022-03-31 22:30:10 +00003136
dan sinclair41e4d9a2022-05-01 14:40:55 +00003137 sem->Behaviors() = rhs->Behaviors() + lhs->Behaviors();
James Pricec022ff52022-03-31 22:30:10 +00003138
dan sinclair41e4d9a2022-05-01 14:40:55 +00003139 auto* lhs_ty = lhs->Type()->UnwrapRef();
3140 auto* rhs_ty = rhs->Type()->UnwrapRef();
Ben Clayton9fb29a32022-05-09 20:00:13 +00003141 auto* ty = intrinsic_table_->Lookup(stmt->op, lhs_ty, rhs_ty, stmt->source, true).result;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003142 if (!ty) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003143 return false;
3144 }
3145 return validator_.Assignment(stmt, ty);
3146 });
James Pricec022ff52022-03-31 22:30:10 +00003147}
3148
dan sinclair41e4d9a2022-05-01 14:40:55 +00003149sem::Statement* Resolver::ContinueStatement(const ast::ContinueStatement* stmt) {
3150 auto* sem =
3151 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3152 return StatementScope(stmt, sem, [&] {
3153 sem->Behaviors() = sem::Behavior::kContinue;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003154
dan sinclair41e4d9a2022-05-01 14:40:55 +00003155 // Set if we've hit the first continue statement in our parent loop
3156 if (auto* block = sem->FindFirstParent<sem::LoopBlockStatement>()) {
3157 if (!block->FirstContinue()) {
3158 const_cast<sem::LoopBlockStatement*>(block)->SetFirstContinue(
3159 stmt, block->Decls().size());
3160 }
3161 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003162
dan sinclair41e4d9a2022-05-01 14:40:55 +00003163 return validator_.ContinueStatement(sem, current_statement_);
3164 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003165}
3166
3167sem::Statement* Resolver::DiscardStatement(const ast::DiscardStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003168 auto* sem =
3169 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3170 return StatementScope(stmt, sem, [&] {
3171 sem->Behaviors() = sem::Behavior::kDiscard;
3172 current_function_->SetHasDiscard();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003173
dan sinclair41e4d9a2022-05-01 14:40:55 +00003174 return validator_.DiscardStatement(sem, current_statement_);
3175 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003176}
3177
dan sinclair41e4d9a2022-05-01 14:40:55 +00003178sem::Statement* Resolver::FallthroughStatement(const ast::FallthroughStatement* stmt) {
3179 auto* sem =
3180 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3181 return StatementScope(stmt, sem, [&] {
3182 sem->Behaviors() = sem::Behavior::kFallthrough;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003183
dan sinclair41e4d9a2022-05-01 14:40:55 +00003184 return validator_.FallthroughStatement(sem);
3185 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003186}
3187
James Price2f9e31c2022-04-07 13:42:45 +00003188sem::Statement* Resolver::IncrementDecrementStatement(
3189 const ast::IncrementDecrementStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003190 auto* sem =
3191 builder_->create<sem::Statement>(stmt, current_compound_statement_, current_function_);
3192 return StatementScope(stmt, sem, [&] {
3193 auto* lhs = Expression(stmt->lhs);
3194 if (!lhs) {
3195 return false;
3196 }
3197 sem->Behaviors() = lhs->Behaviors();
James Price2f9e31c2022-04-07 13:42:45 +00003198
dan sinclair41e4d9a2022-05-01 14:40:55 +00003199 return validator_.IncrementDecrementStatement(stmt);
3200 });
James Price2f9e31c2022-04-07 13:42:45 +00003201}
3202
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003203bool Resolver::ApplyStorageClassUsageToType(ast::StorageClass sc,
3204 sem::Type* ty,
3205 const Source& usage) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003206 ty = const_cast<sem::Type*>(ty->UnwrapRef());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003207
dan sinclair41e4d9a2022-05-01 14:40:55 +00003208 if (auto* str = ty->As<sem::Struct>()) {
3209 if (str->StorageClassUsage().count(sc)) {
3210 return true; // Already applied
3211 }
3212
3213 str->AddUsage(sc);
3214
3215 for (auto* member : str->Members()) {
Ben Clayton6fb5d372022-07-29 14:32:51 +00003216 if (!ApplyStorageClassUsageToType(sc, const_cast<sem::Type*>(member->Type()), usage)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003217 std::stringstream err;
3218 err << "while analysing structure member " << sem_.TypeNameOf(str) << "."
3219 << builder_->Symbols().NameFor(member->Declaration()->symbol);
3220 AddNote(err.str(), member->Declaration()->source);
3221 return false;
3222 }
3223 }
3224 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003225 }
3226
dan sinclair41e4d9a2022-05-01 14:40:55 +00003227 if (auto* arr = ty->As<sem::Array>()) {
3228 if (arr->IsRuntimeSized() && sc != ast::StorageClass::kStorage) {
3229 AddError(
3230 "runtime-sized arrays can only be used in the <storage> storage "
3231 "class",
3232 usage);
3233 return false;
3234 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003235
dan sinclair41e4d9a2022-05-01 14:40:55 +00003236 return ApplyStorageClassUsageToType(sc, const_cast<sem::Type*>(arr->ElemType()), usage);
3237 }
3238
3239 if (ast::IsHostShareable(sc) && !validator_.IsHostShareable(ty)) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003240 std::stringstream err;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003241 err << "Type '" << sem_.TypeNameOf(ty) << "' cannot be used in storage class '" << sc
3242 << "' as it is non-host-shareable";
3243 AddError(err.str(), usage);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003244 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003245 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00003246
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003247 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003248}
3249
3250template <typename SEM, typename F>
dan sinclair41e4d9a2022-05-01 14:40:55 +00003251SEM* Resolver::StatementScope(const ast::Statement* ast, SEM* sem, F&& callback) {
3252 builder_->Sem().Add(ast, sem);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003253
dan sinclair41e4d9a2022-05-01 14:40:55 +00003254 auto* as_compound = As<sem::CompoundStatement, CastFlags::kDontErrorOnImpossibleCast>(sem);
3255 auto* as_block = As<sem::BlockStatement, CastFlags::kDontErrorOnImpossibleCast>(sem);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003256
dan sinclair41e4d9a2022-05-01 14:40:55 +00003257 TINT_SCOPED_ASSIGNMENT(current_statement_, sem);
3258 TINT_SCOPED_ASSIGNMENT(current_compound_statement_,
3259 as_compound ? as_compound : current_compound_statement_);
3260 TINT_SCOPED_ASSIGNMENT(current_block_, as_block ? as_block : current_block_);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003261
dan sinclair41e4d9a2022-05-01 14:40:55 +00003262 if (!callback()) {
3263 return nullptr;
3264 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003265
dan sinclair41e4d9a2022-05-01 14:40:55 +00003266 return sem;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003267}
3268
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003269bool Resolver::Mark(const ast::Node* node) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003270 if (node == nullptr) {
3271 TINT_ICE(Resolver, diagnostics_) << "Resolver::Mark() called with nullptr";
3272 return false;
3273 }
Ben Claytonaf473882022-07-21 23:46:15 +00003274 auto marked_bit_ref = marked_[node->node_id.value];
3275 if (!marked_bit_ref) {
3276 marked_bit_ref = true;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003277 return true;
3278 }
3279 TINT_ICE(Resolver, diagnostics_) << "AST node '" << node->TypeInfo().name
3280 << "' was encountered twice in the same AST of a Program\n"
3281 << "At: " << node->source << "\n"
3282 << "Pointer: " << node;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003283 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003284}
3285
3286void Resolver::AddError(const std::string& msg, const Source& source) const {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003287 diagnostics_.add_error(diag::System::Resolver, msg, source);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003288}
3289
3290void Resolver::AddWarning(const std::string& msg, const Source& source) const {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003291 diagnostics_.add_warning(diag::System::Resolver, msg, source);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003292}
3293
3294void Resolver::AddNote(const std::string& msg, const Source& source) const {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003295 diagnostics_.add_note(diag::System::Resolver, msg, source);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003296}
3297
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003298bool Resolver::IsBuiltin(Symbol symbol) const {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003299 std::string name = builder_->Symbols().NameFor(symbol);
3300 return sem::ParseBuiltinType(name) != sem::BuiltinType::kNone;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003301}
3302
dan sinclaird2093792022-04-07 17:45:45 +00003303} // namespace tint::resolver