blob: d9c34791519fb0cc4785e42a59b7efe74dff4df4 [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/writer/hlsl/generator_impl.h"
16
17#include <algorithm>
18#include <cmath>
19#include <functional>
20#include <iomanip>
21#include <set>
22#include <utility>
23#include <vector>
24
25#include "src/tint/ast/call_statement.h"
26#include "src/tint/ast/fallthrough_statement.h"
27#include "src/tint/ast/id_attribute.h"
28#include "src/tint/ast/internal_attribute.h"
29#include "src/tint/ast/interpolate_attribute.h"
30#include "src/tint/ast/variable_decl_statement.h"
31#include "src/tint/debug.h"
32#include "src/tint/sem/array.h"
Ben Clayton01004b72022-04-28 18:49:04 +000033#include "src/tint/sem/atomic.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000034#include "src/tint/sem/block_statement.h"
35#include "src/tint/sem/call.h"
Ben Claytone9f8b092022-06-01 13:14:39 +000036#include "src/tint/sem/constant.h"
Ben Clayton01004b72022-04-28 18:49:04 +000037#include "src/tint/sem/depth_multisampled_texture.h"
38#include "src/tint/sem/depth_texture.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000039#include "src/tint/sem/function.h"
40#include "src/tint/sem/member_accessor_expression.h"
41#include "src/tint/sem/module.h"
Ben Clayton01004b72022-04-28 18:49:04 +000042#include "src/tint/sem/multisampled_texture.h"
43#include "src/tint/sem/sampled_texture.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000044#include "src/tint/sem/statement.h"
Ben Clayton01004b72022-04-28 18:49:04 +000045#include "src/tint/sem/storage_texture.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000046#include "src/tint/sem/struct.h"
dan sinclaird32fbe02022-10-19 00:43:41 +000047#include "src/tint/sem/switch_statement.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000048#include "src/tint/sem/type_conversion.h"
dan sinclair6e77b472022-10-20 13:38:28 +000049#include "src/tint/sem/type_initializer.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000050#include "src/tint/sem/variable.h"
51#include "src/tint/transform/add_empty_entry_point.h"
52#include "src/tint/transform/array_length_from_uniform.h"
Ben Clayton27aa57c2022-02-22 23:13:39 +000053#include "src/tint/transform/builtin_polyfill.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000054#include "src/tint/transform/calculate_array_length.h"
55#include "src/tint/transform/canonicalize_entry_point_io.h"
56#include "src/tint/transform/decompose_memory_access.h"
James Price791b4352022-05-11 13:50:33 +000057#include "src/tint/transform/disable_uniformity_analysis.h"
James Price508a9662022-03-31 22:30:10 +000058#include "src/tint/transform/expand_compound_assignment.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000059#include "src/tint/transform/localize_struct_array_assignment.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000060#include "src/tint/transform/manager.h"
61#include "src/tint/transform/num_workgroups_from_uniform.h"
Ben Clayton7ebcfc72022-06-27 20:20:25 +000062#include "src/tint/transform/promote_initializers_to_let.h"
Antonio Maiorano93baaae2022-03-15 15:35:13 +000063#include "src/tint/transform/promote_side_effects_to_decl.h"
Antonio Maioranob3497102022-03-31 15:02:25 +000064#include "src/tint/transform/remove_continue_in_switch.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000065#include "src/tint/transform/remove_phonies.h"
66#include "src/tint/transform/simplify_pointers.h"
67#include "src/tint/transform/unshadow.h"
Antonio Maiorano66d66682022-03-28 20:51:32 +000068#include "src/tint/transform/unwind_discard_functions.h"
dan sinclair6e77b472022-10-20 13:38:28 +000069#include "src/tint/transform/vectorize_scalar_matrix_initializers.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000070#include "src/tint/transform/zero_init_workgroup_memory.h"
71#include "src/tint/utils/defer.h"
72#include "src/tint/utils/map.h"
73#include "src/tint/utils/scoped_assignment.h"
dan sinclair4abf28e2022-08-02 15:55:35 +000074#include "src/tint/utils/string.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000075#include "src/tint/writer/append_vector.h"
Ben Clayton1a567782022-10-14 13:38:27 +000076#include "src/tint/writer/check_supported_extensions.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000077#include "src/tint/writer/float_to_string.h"
Antonio Maioranoa730eb72022-04-06 13:57:54 +000078#include "src/tint/writer/generate_external_texture_bindings.h"
Ryan Harrisondbc13af2022-02-21 15:19:07 +000079
Ben Clayton0ce9ab02022-05-05 20:23:40 +000080using namespace tint::number_suffixes; // NOLINT
81
dan sinclair6a5bef12022-04-07 14:30:24 +000082namespace tint::writer::hlsl {
Ryan Harrisondbc13af2022-02-21 15:19:07 +000083namespace {
84
85const char kTempNamePrefix[] = "tint_tmp";
Ryan Harrisondbc13af2022-02-21 15:19:07 +000086
87const char* image_format_to_rwtexture_type(ast::TexelFormat image_format) {
dan sinclair41e4d9a2022-05-01 14:40:55 +000088 switch (image_format) {
89 case ast::TexelFormat::kRgba8Unorm:
90 case ast::TexelFormat::kRgba8Snorm:
91 case ast::TexelFormat::kRgba16Float:
92 case ast::TexelFormat::kR32Float:
93 case ast::TexelFormat::kRg32Float:
94 case ast::TexelFormat::kRgba32Float:
95 return "float4";
96 case ast::TexelFormat::kRgba8Uint:
97 case ast::TexelFormat::kRgba16Uint:
98 case ast::TexelFormat::kR32Uint:
99 case ast::TexelFormat::kRg32Uint:
100 case ast::TexelFormat::kRgba32Uint:
101 return "uint4";
102 case ast::TexelFormat::kRgba8Sint:
103 case ast::TexelFormat::kRgba16Sint:
104 case ast::TexelFormat::kR32Sint:
105 case ast::TexelFormat::kRg32Sint:
106 case ast::TexelFormat::kRgba32Sint:
107 return "int4";
108 default:
109 return nullptr;
110 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000111}
112
Ben Claytone9f8b092022-06-01 13:14:39 +0000113void PrintF32(std::ostream& out, float value) {
Ben Claytone9f8b092022-06-01 13:14:39 +0000114 if (std::isinf(value)) {
Antonio Maiorano679cf4f2022-09-03 21:43:01 +0000115 out << "0.0f " << (value >= 0 ? "/* inf */" : "/* -inf */");
Ben Claytone9f8b092022-06-01 13:14:39 +0000116 } else if (std::isnan(value)) {
Antonio Maiorano679cf4f2022-09-03 21:43:01 +0000117 out << "0.0f /* nan */";
Ben Claytone9f8b092022-06-01 13:14:39 +0000118 } else {
119 out << FloatToString(value) << "f";
120 }
121}
122
Antonio Maiorano679cf4f2022-09-03 21:43:01 +0000123void PrintF16(std::ostream& out, float value) {
124 if (std::isinf(value)) {
125 out << "0.0h " << (value >= 0 ? "/* inf */" : "/* -inf */");
126 } else if (std::isnan(value)) {
127 out << "0.0h /* nan */";
Zhaoming Jianga5988a32022-07-11 15:43:38 +0000128 } else {
129 out << FloatToString(value) << "h";
Zhaoming Jianga5988a32022-07-11 15:43:38 +0000130 }
131}
132
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000133// Helper for writing " : register(RX, spaceY)", where R is the register, X is
134// the binding point binding value, and Y is the binding point group value.
135struct RegisterAndSpace {
dan sinclairacdf6e12022-08-24 15:47:25 +0000136 RegisterAndSpace(char r, sem::BindingPoint bp) : reg(r), binding_point(bp) {}
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000137
dan sinclair41e4d9a2022-05-01 14:40:55 +0000138 const char reg;
dan sinclairacdf6e12022-08-24 15:47:25 +0000139 sem::BindingPoint const binding_point;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000140};
141
142std::ostream& operator<<(std::ostream& s, const RegisterAndSpace& rs) {
dan sinclairacdf6e12022-08-24 15:47:25 +0000143 s << " : register(" << rs.reg << rs.binding_point.binding << ", space" << rs.binding_point.group
144 << ")";
dan sinclair41e4d9a2022-05-01 14:40:55 +0000145 return s;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000146}
147
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000148} // namespace
149
150SanitizedResult::SanitizedResult() = default;
151SanitizedResult::~SanitizedResult() = default;
152SanitizedResult::SanitizedResult(SanitizedResult&&) = default;
153
Antonio Maiorano7eaab382022-04-11 16:33:30 +0000154SanitizedResult Sanitize(const Program* in, const Options& options) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000155 transform::Manager manager;
156 transform::DataMap data;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000157
James Price791b4352022-05-11 13:50:33 +0000158 manager.Add<transform::DisableUniformityAnalysis>();
159
dan sinclair41e4d9a2022-05-01 14:40:55 +0000160 { // Builtin polyfills
161 transform::BuiltinPolyfill::Builtins polyfills;
dan sinclaird23f2962022-06-28 15:27:44 +0000162 polyfills.acosh = transform::BuiltinPolyfill::Level::kFull;
163 polyfills.asinh = true;
164 polyfills.atanh = transform::BuiltinPolyfill::Level::kFull;
Ben Clayton6dbb4632022-10-31 17:54:49 +0000165 polyfills.clamp_int = true;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000166 // TODO(crbug.com/tint/1449): Some of these can map to HLSL's `firstbitlow`
167 // and `firstbithigh`.
168 polyfills.count_leading_zeros = true;
169 polyfills.count_trailing_zeros = true;
170 polyfills.extract_bits = transform::BuiltinPolyfill::Level::kFull;
171 polyfills.first_leading_bit = true;
172 polyfills.first_trailing_bit = true;
173 polyfills.insert_bits = transform::BuiltinPolyfill::Level::kFull;
Ben Claytonc4ebf2c2022-09-22 22:59:16 +0000174 polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000175 data.Add<transform::BuiltinPolyfill::Config>(polyfills);
176 manager.Add<transform::BuiltinPolyfill>();
177 }
Ben Clayton27aa57c2022-02-22 23:13:39 +0000178
dan sinclair41e4d9a2022-05-01 14:40:55 +0000179 // Build the config for the internal ArrayLengthFromUniform transform.
180 auto& array_length_from_uniform = options.array_length_from_uniform;
181 transform::ArrayLengthFromUniform::Config array_length_from_uniform_cfg(
182 array_length_from_uniform.ubo_binding);
183 array_length_from_uniform_cfg.bindpoint_to_size_index =
184 array_length_from_uniform.bindpoint_to_size_index;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000185
dan sinclair41e4d9a2022-05-01 14:40:55 +0000186 if (options.generate_external_texture_bindings) {
187 auto new_bindings_map = GenerateExternalTextureBindings(in);
188 data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
189 }
190 manager.Add<transform::MultiplanarExternalTexture>();
Antonio Maioranoa730eb72022-04-06 13:57:54 +0000191
dan sinclair41e4d9a2022-05-01 14:40:55 +0000192 manager.Add<transform::Unshadow>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000193
dan sinclair41e4d9a2022-05-01 14:40:55 +0000194 // LocalizeStructArrayAssignment must come after:
195 // * SimplifyPointers, because it assumes assignment to arrays in structs are
196 // done directly, not indirectly.
197 // TODO(crbug.com/tint/1340): See if we can get rid of the duplicate
198 // SimplifyPointers transform. Can't do it right now because
199 // LocalizeStructArrayAssignment introduces pointers.
200 manager.Add<transform::SimplifyPointers>();
201 manager.Add<transform::LocalizeStructArrayAssignment>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000202
dan sinclair41e4d9a2022-05-01 14:40:55 +0000203 if (!options.disable_workgroup_init) {
204 // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
205 // ZeroInitWorkgroupMemory may inject new builtin parameters.
206 manager.Add<transform::ZeroInitWorkgroupMemory>();
207 }
208 manager.Add<transform::CanonicalizeEntryPointIO>();
209 // NumWorkgroupsFromUniform must come after CanonicalizeEntryPointIO, as it
210 // assumes that num_workgroups builtins only appear as struct members and are
211 // only accessed directly via member accessors.
212 manager.Add<transform::NumWorkgroupsFromUniform>();
213 manager.Add<transform::ExpandCompoundAssignment>();
214 manager.Add<transform::PromoteSideEffectsToDecl>();
215 manager.Add<transform::UnwindDiscardFunctions>();
dan sinclair6e77b472022-10-20 13:38:28 +0000216 manager.Add<transform::VectorizeScalarMatrixInitializers>();
dan sinclair41e4d9a2022-05-01 14:40:55 +0000217 manager.Add<transform::SimplifyPointers>();
218 manager.Add<transform::RemovePhonies>();
219 // ArrayLengthFromUniform must come after InlinePointerLets and Simplify, as
220 // it assumes that the form of the array length argument is &var.array.
221 manager.Add<transform::ArrayLengthFromUniform>();
222 data.Add<transform::ArrayLengthFromUniform::Config>(std::move(array_length_from_uniform_cfg));
223 // DecomposeMemoryAccess must come after:
224 // * InlinePointerLets, as we cannot take the address of calls to
225 // DecomposeMemoryAccess::Intrinsic.
226 // * Simplify, as we need to fold away the address-of and dereferences of
227 // `*(&(intrinsic_load()))` expressions.
228 // * RemovePhonies, as phonies can be assigned a pointer to a
229 // non-constructible buffer, or dynamic array, which DMA cannot cope with.
230 manager.Add<transform::DecomposeMemoryAccess>();
231 // CalculateArrayLength must come after DecomposeMemoryAccess, as
232 // DecomposeMemoryAccess special-cases the arrayLength() intrinsic, which
233 // will be transformed by CalculateArrayLength
234 manager.Add<transform::CalculateArrayLength>();
Ben Clayton7ebcfc72022-06-27 20:20:25 +0000235 manager.Add<transform::PromoteInitializersToLet>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000236
dan sinclair41e4d9a2022-05-01 14:40:55 +0000237 manager.Add<transform::RemoveContinueInSwitch>();
Antonio Maioranob3497102022-03-31 15:02:25 +0000238
dan sinclair41e4d9a2022-05-01 14:40:55 +0000239 manager.Add<transform::AddEmptyEntryPoint>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000240
dan sinclair41e4d9a2022-05-01 14:40:55 +0000241 data.Add<transform::CanonicalizeEntryPointIO::Config>(
242 transform::CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
243 data.Add<transform::NumWorkgroupsFromUniform::Config>(options.root_constant_binding_point);
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000244
dan sinclair41e4d9a2022-05-01 14:40:55 +0000245 auto out = manager.Run(in, data);
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000246
dan sinclair41e4d9a2022-05-01 14:40:55 +0000247 SanitizedResult result;
248 result.program = std::move(out.program);
249 if (auto* res = out.data.Get<transform::ArrayLengthFromUniform::Result>()) {
250 result.used_array_length_from_uniform_indices = std::move(res->used_size_indices);
251 }
252 return result;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000253}
254
255GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
256
257GeneratorImpl::~GeneratorImpl() = default;
258
259bool GeneratorImpl::Generate() {
Ben Clayton1a567782022-10-14 13:38:27 +0000260 if (!CheckSupportedExtensions("HLSL", program_->AST(), diagnostics_,
261 utils::Vector{
262 ast::Extension::kChromiumDisableUniformityAnalysis,
263 ast::Extension::kChromiumExperimentalDp4A,
264 ast::Extension::kChromiumExperimentalPushConstant,
265 ast::Extension::kF16,
266 })) {
267 return false;
268 }
269
dan sinclair41e4d9a2022-05-01 14:40:55 +0000270 const TypeInfo* last_kind = nullptr;
271 size_t last_padding_line = 0;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000272
dan sinclair41e4d9a2022-05-01 14:40:55 +0000273 auto* mod = builder_.Sem().Module();
274 for (auto* decl : mod->DependencyOrderedDeclarations()) {
Ben Claytonb4744ac2022-08-03 07:01:08 +0000275 if (decl->IsAnyOf<ast::Alias, ast::Enable, ast::StaticAssert>()) {
276 continue; // These are not emitted.
James Price791b4352022-05-11 13:50:33 +0000277 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000278
dan sinclair41e4d9a2022-05-01 14:40:55 +0000279 // Emit a new line between declarations if the type of declaration has
280 // changed, or we're about to emit a function
281 auto* kind = &decl->TypeInfo();
282 if (current_buffer_->lines.size() != last_padding_line) {
283 if (last_kind && (last_kind != kind || decl->Is<ast::Function>())) {
284 line();
285 last_padding_line = current_buffer_->lines.size();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000286 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000287 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000288 last_kind = kind;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000289
dan sinclair41e4d9a2022-05-01 14:40:55 +0000290 bool ok = Switch(
291 decl,
292 [&](const ast::Variable* global) { //
293 return EmitGlobalVariable(global);
294 },
295 [&](const ast::Struct* str) {
296 auto* ty = builder_.Sem().Get(str);
dan sinclairff7cf212022-10-03 14:05:23 +0000297 auto address_space_uses = ty->AddressSpaceUsage();
298 if (address_space_uses.size() !=
299 (address_space_uses.count(ast::AddressSpace::kStorage) +
300 address_space_uses.count(ast::AddressSpace::kUniform))) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000301 // The structure is used as something other than a storage buffer or
302 // uniform buffer, so it needs to be emitted.
303 // Storage buffer are read and written to via a ByteAddressBuffer
304 // instead of true structure.
305 // Structures used as uniform buffer are read from an array of
306 // vectors instead of true structure.
307 return EmitStructType(current_buffer_, ty);
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000308 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000309 return true;
310 },
311 [&](const ast::Function* func) {
312 if (func->IsEntryPoint()) {
313 return EmitEntryPointFunction(func);
314 }
315 return EmitFunction(func);
316 },
dan sinclair41e4d9a2022-05-01 14:40:55 +0000317 [&](Default) {
318 TINT_ICE(Writer, diagnostics_)
319 << "unhandled module-scope declaration: " << decl->TypeInfo().name;
320 return false;
321 });
322
323 if (!ok) {
324 return false;
325 }
326 }
327
328 if (!helpers_.lines.empty()) {
329 current_buffer_->Insert(helpers_, 0, 0);
330 }
331
332 return true;
333}
334
335bool GeneratorImpl::EmitDynamicVectorAssignment(const ast::AssignmentStatement* stmt,
336 const sem::Vector* vec) {
337 auto name = utils::GetOrCreate(dynamic_vector_write_, vec, [&]() -> std::string {
338 std::string fn;
339 {
340 std::ostringstream ss;
Ben Claytond2e0db32022-10-12 18:49:15 +0000341 if (!EmitType(ss, vec, tint::ast::AddressSpace::kUndefined, ast::Access::kUndefined,
342 "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000343 return "";
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000344 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000345 fn = UniqueIdentifier("set_" + ss.str());
346 }
347 {
348 auto out = line(&helpers_);
349 out << "void " << fn << "(inout ";
Ben Claytond2e0db32022-10-12 18:49:15 +0000350 if (!EmitTypeAndName(out, vec, ast::AddressSpace::kUndefined, ast::Access::kUndefined,
dan sinclair41e4d9a2022-05-01 14:40:55 +0000351 "vec")) {
352 return "";
353 }
354 out << ", int idx, ";
Ben Claytond2e0db32022-10-12 18:49:15 +0000355 if (!EmitTypeAndName(out, vec->type(), ast::AddressSpace::kUndefined,
356 ast::Access::kUndefined, "val")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000357 return "";
358 }
359 out << ") {";
360 }
361 {
362 ScopedIndent si(&helpers_);
363 auto out = line(&helpers_);
364 switch (vec->Width()) {
365 case 2:
366 out << "vec = (idx.xx == int2(0, 1)) ? val.xx : vec;";
367 break;
368 case 3:
369 out << "vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;";
370 break;
371 case 4:
372 out << "vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec;";
373 break;
374 default:
Ben Claytondcdf66e2022-06-17 12:48:51 +0000375 TINT_UNREACHABLE(Writer, diagnostics_)
dan sinclair41e4d9a2022-05-01 14:40:55 +0000376 << "invalid vector size " << vec->Width();
377 break;
378 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000379 }
380 line(&helpers_) << "}";
381 line(&helpers_);
382 return fn;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000383 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000384
dan sinclair41e4d9a2022-05-01 14:40:55 +0000385 if (name.empty()) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000386 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000387 }
388
dan sinclair41e4d9a2022-05-01 14:40:55 +0000389 auto* ast_access_expr = stmt->lhs->As<ast::IndexAccessorExpression>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000390
dan sinclair41e4d9a2022-05-01 14:40:55 +0000391 auto out = line();
392 out << name << "(";
393 if (!EmitExpression(out, ast_access_expr->object)) {
394 return false;
395 }
396 out << ", ";
397 if (!EmitExpression(out, ast_access_expr->index)) {
398 return false;
399 }
400 out << ", ";
401 if (!EmitExpression(out, stmt->rhs)) {
402 return false;
403 }
404 out << ");";
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000405
dan sinclair41e4d9a2022-05-01 14:40:55 +0000406 return true;
407}
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000408
dan sinclair41e4d9a2022-05-01 14:40:55 +0000409bool GeneratorImpl::EmitDynamicMatrixVectorAssignment(const ast::AssignmentStatement* stmt,
410 const sem::Matrix* mat) {
411 auto name = utils::GetOrCreate(dynamic_matrix_vector_write_, mat, [&]() -> std::string {
412 std::string fn;
413 {
414 std::ostringstream ss;
Ben Claytond2e0db32022-10-12 18:49:15 +0000415 if (!EmitType(ss, mat, tint::ast::AddressSpace::kUndefined, ast::Access::kUndefined,
416 "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000417 return "";
418 }
419 fn = UniqueIdentifier("set_vector_" + ss.str());
420 }
421 {
422 auto out = line(&helpers_);
423 out << "void " << fn << "(inout ";
Ben Claytond2e0db32022-10-12 18:49:15 +0000424 if (!EmitTypeAndName(out, mat, ast::AddressSpace::kUndefined, ast::Access::kUndefined,
dan sinclair41e4d9a2022-05-01 14:40:55 +0000425 "mat")) {
426 return "";
427 }
428 out << ", int col, ";
Ben Claytond2e0db32022-10-12 18:49:15 +0000429 if (!EmitTypeAndName(out, mat->ColumnType(), ast::AddressSpace::kUndefined,
430 ast::Access::kUndefined, "val")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000431 return "";
432 }
433 out << ") {";
434 }
435 {
436 ScopedIndent si(&helpers_);
437 line(&helpers_) << "switch (col) {";
438 {
439 ScopedIndent si2(&helpers_);
440 for (uint32_t i = 0; i < mat->columns(); ++i) {
441 line(&helpers_) << "case " << i << ": mat[" << i << "] = val; break;";
442 }
443 }
444 line(&helpers_) << "}";
445 }
446 line(&helpers_) << "}";
447 line(&helpers_);
448 return fn;
449 });
450
451 if (name.empty()) {
452 return false;
453 }
454
455 auto* ast_access_expr = stmt->lhs->As<ast::IndexAccessorExpression>();
456
457 auto out = line();
458 out << name << "(";
459 if (!EmitExpression(out, ast_access_expr->object)) {
460 return false;
461 }
462 out << ", ";
463 if (!EmitExpression(out, ast_access_expr->index)) {
464 return false;
465 }
466 out << ", ";
467 if (!EmitExpression(out, stmt->rhs)) {
468 return false;
469 }
470 out << ");";
471
472 return true;
473}
474
475bool GeneratorImpl::EmitDynamicMatrixScalarAssignment(const ast::AssignmentStatement* stmt,
476 const sem::Matrix* mat) {
477 auto* lhs_col_access = stmt->lhs->As<ast::IndexAccessorExpression>();
478 auto* lhs_row_access = lhs_col_access->object->As<ast::IndexAccessorExpression>();
479
480 auto name = utils::GetOrCreate(dynamic_matrix_scalar_write_, mat, [&]() -> std::string {
481 std::string fn;
482 {
483 std::ostringstream ss;
Ben Claytond2e0db32022-10-12 18:49:15 +0000484 if (!EmitType(ss, mat, tint::ast::AddressSpace::kUndefined, ast::Access::kUndefined,
485 "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000486 return "";
487 }
488 fn = UniqueIdentifier("set_scalar_" + ss.str());
489 }
490 {
491 auto out = line(&helpers_);
492 out << "void " << fn << "(inout ";
Ben Claytond2e0db32022-10-12 18:49:15 +0000493 if (!EmitTypeAndName(out, mat, ast::AddressSpace::kUndefined, ast::Access::kUndefined,
dan sinclair41e4d9a2022-05-01 14:40:55 +0000494 "mat")) {
495 return "";
496 }
497 out << ", int col, int row, ";
Ben Claytond2e0db32022-10-12 18:49:15 +0000498 if (!EmitTypeAndName(out, mat->type(), ast::AddressSpace::kUndefined,
499 ast::Access::kUndefined, "val")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000500 return "";
501 }
502 out << ") {";
503 }
504 {
505 ScopedIndent si(&helpers_);
506 line(&helpers_) << "switch (col) {";
507 {
508 ScopedIndent si2(&helpers_);
509 auto* vec = TypeOf(lhs_row_access->object)->UnwrapRef()->As<sem::Vector>();
510 for (uint32_t i = 0; i < mat->columns(); ++i) {
511 line(&helpers_) << "case " << i << ":";
512 {
513 auto vec_name = "mat[" + std::to_string(i) + "]";
514 ScopedIndent si3(&helpers_);
515 {
516 auto out = line(&helpers_);
517 switch (mat->rows()) {
518 case 2:
519 out << vec_name
520 << " = (row.xx == int2(0, 1)) ? val.xx : " << vec_name
521 << ";";
522 break;
523 case 3:
524 out << vec_name
525 << " = (row.xxx == int3(0, 1, 2)) ? val.xxx : " << vec_name
526 << ";";
527 break;
528 case 4:
529 out << vec_name
530 << " = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : "
531 << vec_name << ";";
532 break;
533 default:
Ben Claytondcdf66e2022-06-17 12:48:51 +0000534 TINT_UNREACHABLE(Writer, diagnostics_)
dan sinclair41e4d9a2022-05-01 14:40:55 +0000535 << "invalid vector size " << vec->Width();
536 break;
537 }
538 }
539 line(&helpers_) << "break;";
540 }
541 }
542 }
543 line(&helpers_) << "}";
544 }
545 line(&helpers_) << "}";
546 line(&helpers_);
547 return fn;
548 });
549
550 if (name.empty()) {
551 return false;
552 }
553
554 auto out = line();
555 out << name << "(";
556 if (!EmitExpression(out, lhs_row_access->object)) {
557 return false;
558 }
559 out << ", ";
560 if (!EmitExpression(out, lhs_col_access->index)) {
561 return false;
562 }
563 out << ", ";
564 if (!EmitExpression(out, lhs_row_access->index)) {
565 return false;
566 }
567 out << ", ";
568 if (!EmitExpression(out, stmt->rhs)) {
569 return false;
570 }
571 out << ");";
572
573 return true;
574}
575
576bool GeneratorImpl::EmitIndexAccessor(std::ostream& out, const ast::IndexAccessorExpression* expr) {
577 if (!EmitExpression(out, expr->object)) {
578 return false;
579 }
580 out << "[";
581
582 if (!EmitExpression(out, expr->index)) {
583 return false;
584 }
585 out << "]";
586
587 return true;
588}
589
590bool GeneratorImpl::EmitBitcast(std::ostream& out, const ast::BitcastExpression* expr) {
591 auto* type = TypeOf(expr);
592 if (auto* vec = type->UnwrapRef()->As<sem::Vector>()) {
593 type = vec->type();
594 }
595
596 if (!type->is_integer_scalar() && !type->is_float_scalar()) {
597 diagnostics_.add_error(diag::System::Writer, "Unable to do bitcast to type " +
598 type->FriendlyName(builder_.Symbols()));
599 return false;
600 }
601
602 out << "as";
dan sinclairff7cf212022-10-03 14:05:23 +0000603 if (!EmitType(out, type, ast::AddressSpace::kNone, ast::Access::kReadWrite, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000604 return false;
605 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000606 out << "(";
dan sinclair41e4d9a2022-05-01 14:40:55 +0000607 if (!EmitExpression(out, expr->expr)) {
608 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000609 }
610 out << ")";
611 return true;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000612}
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000613
dan sinclair41e4d9a2022-05-01 14:40:55 +0000614bool GeneratorImpl::EmitAssign(const ast::AssignmentStatement* stmt) {
615 if (auto* lhs_access = stmt->lhs->As<ast::IndexAccessorExpression>()) {
616 // BUG(crbug.com/tint/1333): work around assignment of scalar to matrices
617 // with at least one dynamic index
618 if (auto* lhs_sub_access = lhs_access->object->As<ast::IndexAccessorExpression>()) {
619 if (auto* mat = TypeOf(lhs_sub_access->object)->UnwrapRef()->As<sem::Matrix>()) {
620 auto* rhs_col_idx_sem = builder_.Sem().Get(lhs_access->index);
621 auto* rhs_row_idx_sem = builder_.Sem().Get(lhs_sub_access->index);
Ben Claytonaa037ac2022-06-29 19:07:30 +0000622 if (!rhs_col_idx_sem->ConstantValue() || !rhs_row_idx_sem->ConstantValue()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000623 return EmitDynamicMatrixScalarAssignment(stmt, mat);
624 }
625 }
626 }
627 // BUG(crbug.com/tint/1333): work around assignment of vector to matrices
628 // with dynamic indices
629 const auto* lhs_access_type = TypeOf(lhs_access->object)->UnwrapRef();
630 if (auto* mat = lhs_access_type->As<sem::Matrix>()) {
631 auto* lhs_index_sem = builder_.Sem().Get(lhs_access->index);
Ben Claytonaa037ac2022-06-29 19:07:30 +0000632 if (!lhs_index_sem->ConstantValue()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000633 return EmitDynamicMatrixVectorAssignment(stmt, mat);
634 }
635 }
636 // BUG(crbug.com/tint/534): work around assignment to vectors with dynamic
637 // indices
638 if (auto* vec = lhs_access_type->As<sem::Vector>()) {
639 auto* rhs_sem = builder_.Sem().Get(lhs_access->index);
Ben Claytonaa037ac2022-06-29 19:07:30 +0000640 if (!rhs_sem->ConstantValue()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000641 return EmitDynamicVectorAssignment(stmt, vec);
642 }
643 }
644 }
645
646 auto out = line();
647 if (!EmitExpression(out, stmt->lhs)) {
648 return false;
649 }
650 out << " = ";
651 if (!EmitExpression(out, stmt->rhs)) {
652 return false;
653 }
654 out << ";";
655 return true;
656}
657
658bool GeneratorImpl::EmitExpressionOrOneIfZero(std::ostream& out, const ast::Expression* expr) {
659 // For constants, replace literal 0 with 1.
Ben Claytonaa037ac2022-06-29 19:07:30 +0000660 if (const auto* val = builder_.Sem().Get(expr)->ConstantValue()) {
661 if (!val->AnyZero()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000662 return EmitExpression(out, expr);
663 }
664
Ben Claytonaa037ac2022-06-29 19:07:30 +0000665 auto* ty = val->Type();
666
667 if (ty->IsAnyOf<sem::I32, sem::U32>()) {
668 return EmitValue(out, ty, 1);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000669 }
670
Ben Claytonaa037ac2022-06-29 19:07:30 +0000671 if (auto* vec = ty->As<sem::Vector>()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000672 auto* elem_ty = vec->type();
673
Ben Claytond2e0db32022-10-12 18:49:15 +0000674 if (!EmitType(out, ty, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000675 return false;
676 }
677
678 out << "(";
Ben Claytonaa037ac2022-06-29 19:07:30 +0000679 for (size_t i = 0; i < vec->Width(); ++i) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000680 if (i != 0) {
681 out << ", ";
682 }
Ben Claytonaa037ac2022-06-29 19:07:30 +0000683 auto s = val->Index(i)->As<AInt>();
Ben Claytonaaa9ba32022-05-17 20:51:04 +0000684 if (!EmitValue(out, elem_ty, (s == 0) ? 1 : static_cast<int>(s))) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000685 return false;
686 }
687 }
688 out << ")";
689 return true;
690 }
691
692 TINT_ICE(Writer, diagnostics_)
693 << "EmitExpressionOrOneIfZero expects integer scalar or vector";
694 return false;
695 }
696
697 auto* ty = TypeOf(expr)->UnwrapRef();
698
699 // For non-constants, we need to emit runtime code to check if the value is 0,
700 // and return 1 in that case.
701 std::string zero;
702 {
703 std::ostringstream ss;
704 EmitValue(ss, ty, 0);
705 zero = ss.str();
706 }
707 std::string one;
708 {
709 std::ostringstream ss;
710 EmitValue(ss, ty, 1);
711 one = ss.str();
712 }
713
714 // For identifiers, no need for a function call as it's fine to evaluate
715 // `expr` more than once.
716 if (expr->Is<ast::IdentifierExpression>()) {
717 out << "(";
718 if (!EmitExpression(out, expr)) {
719 return false;
720 }
721 out << " == " << zero << " ? " << one << " : ";
722 if (!EmitExpression(out, expr)) {
723 return false;
724 }
725 out << ")";
726 return true;
727 }
728
729 // For non-identifier expressions, call a function to make sure `expr` is only
730 // evaluated once.
731 auto name = utils::GetOrCreate(value_or_one_if_zero_, ty, [&]() -> std::string {
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000732 // Example:
733 // int4 tint_value_or_one_if_zero_int4(int4 value) {
734 // return value == 0 ? 0 : value;
735 // }
736 std::string ty_name;
737 {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000738 std::ostringstream ss;
Ben Claytond2e0db32022-10-12 18:49:15 +0000739 if (!EmitType(ss, ty, tint::ast::AddressSpace::kUndefined, ast::Access::kUndefined,
740 "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000741 return "";
742 }
743 ty_name = ss.str();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000744 }
745
746 std::string fn = UniqueIdentifier("value_or_one_if_zero_" + ty_name);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000747 line(&helpers_) << ty_name << " " << fn << "(" << ty_name << " value) {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000748 {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000749 ScopedIndent si(&helpers_);
750 line(&helpers_) << "return value == " << zero << " ? " << one << " : value;";
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000751 }
752 line(&helpers_) << "}";
753 line(&helpers_);
754 return fn;
dan sinclair41e4d9a2022-05-01 14:40:55 +0000755 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000756
dan sinclair41e4d9a2022-05-01 14:40:55 +0000757 if (name.empty()) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000758 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000759 }
760
dan sinclair41e4d9a2022-05-01 14:40:55 +0000761 out << name << "(";
762 if (!EmitExpression(out, expr)) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000763 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000764 }
765 out << ")";
dan sinclair41e4d9a2022-05-01 14:40:55 +0000766 return true;
767}
768
769bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* expr) {
770 if (expr->op == ast::BinaryOp::kLogicalAnd || expr->op == ast::BinaryOp::kLogicalOr) {
771 auto name = UniqueIdentifier(kTempNamePrefix);
772
773 {
774 auto pre = line();
775 pre << "bool " << name << " = ";
776 if (!EmitExpression(pre, expr->lhs)) {
777 return false;
778 }
779 pre << ";";
780 }
781
782 if (expr->op == ast::BinaryOp::kLogicalOr) {
783 line() << "if (!" << name << ") {";
784 } else {
785 line() << "if (" << name << ") {";
786 }
787
788 {
789 ScopedIndent si(this);
790 auto pre = line();
791 pre << name << " = ";
792 if (!EmitExpression(pre, expr->rhs)) {
793 return false;
794 }
795 pre << ";";
796 }
797
798 line() << "}";
799
800 out << "(" << name << ")";
801 return true;
802 }
803
804 auto* lhs_type = TypeOf(expr->lhs)->UnwrapRef();
805 auto* rhs_type = TypeOf(expr->rhs)->UnwrapRef();
806 // Multiplying by a matrix requires the use of `mul` in order to get the
807 // type of multiply we desire.
808 if (expr->op == ast::BinaryOp::kMultiply &&
809 ((lhs_type->Is<sem::Vector>() && rhs_type->Is<sem::Matrix>()) ||
810 (lhs_type->Is<sem::Matrix>() && rhs_type->Is<sem::Vector>()) ||
811 (lhs_type->Is<sem::Matrix>() && rhs_type->Is<sem::Matrix>()))) {
812 // Matrices are transposed, so swap LHS and RHS.
813 out << "mul(";
814 if (!EmitExpression(out, expr->rhs)) {
815 return false;
816 }
817 out << ", ";
818 if (!EmitExpression(out, expr->lhs)) {
819 return false;
820 }
821 out << ")";
822
823 return true;
824 }
825
Ben Claytone9f8b092022-06-01 13:14:39 +0000826 ScopedParen sp(out);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000827
828 if (!EmitExpression(out, expr->lhs)) {
829 return false;
830 }
831 out << " ";
832
833 switch (expr->op) {
834 case ast::BinaryOp::kAnd:
835 out << "&";
836 break;
837 case ast::BinaryOp::kOr:
838 out << "|";
839 break;
840 case ast::BinaryOp::kXor:
841 out << "^";
842 break;
843 case ast::BinaryOp::kLogicalAnd:
844 case ast::BinaryOp::kLogicalOr: {
845 // These are both handled above.
846 TINT_UNREACHABLE(Writer, diagnostics_);
847 return false;
848 }
849 case ast::BinaryOp::kEqual:
850 out << "==";
851 break;
852 case ast::BinaryOp::kNotEqual:
853 out << "!=";
854 break;
855 case ast::BinaryOp::kLessThan:
856 out << "<";
857 break;
858 case ast::BinaryOp::kGreaterThan:
859 out << ">";
860 break;
861 case ast::BinaryOp::kLessThanEqual:
862 out << "<=";
863 break;
864 case ast::BinaryOp::kGreaterThanEqual:
865 out << ">=";
866 break;
867 case ast::BinaryOp::kShiftLeft:
868 out << "<<";
869 break;
870 case ast::BinaryOp::kShiftRight:
871 // TODO(dsinclair): MSL is based on C++14, and >> in C++14 has
872 // implementation-defined behaviour for negative LHS. We may have to
873 // generate extra code to implement WGSL-specified behaviour for negative
874 // LHS.
875 out << R"(>>)";
876 break;
877
878 case ast::BinaryOp::kAdd:
879 out << "+";
880 break;
881 case ast::BinaryOp::kSubtract:
882 out << "-";
883 break;
884 case ast::BinaryOp::kMultiply:
885 out << "*";
886 break;
887 case ast::BinaryOp::kDivide:
888 out << "/";
889 // BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC
890 // compile error, and undefined behavior in WGSL.
891 if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) {
892 out << " ";
893 return EmitExpressionOrOneIfZero(out, expr->rhs);
894 }
895 break;
896 case ast::BinaryOp::kModulo:
897 out << "%";
898 // BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC
899 // compile error, and undefined behavior in WGSL.
900 if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) {
901 out << " ";
902 return EmitExpressionOrOneIfZero(out, expr->rhs);
903 }
904 break;
905 case ast::BinaryOp::kNone:
906 diagnostics_.add_error(diag::System::Writer, "missing binary operation type");
907 return false;
908 }
909 out << " ";
910
911 if (!EmitExpression(out, expr->rhs)) {
912 return false;
913 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000914
915 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000916}
917
Ben Clayton783b1692022-08-02 17:03:35 +0000918bool GeneratorImpl::EmitStatements(utils::VectorRef<const ast::Statement*> stmts) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000919 for (auto* s : stmts) {
920 if (!EmitStatement(s)) {
921 return false;
922 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000923 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000924 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000925}
926
Ben Clayton783b1692022-08-02 17:03:35 +0000927bool GeneratorImpl::EmitStatementsWithIndent(utils::VectorRef<const ast::Statement*> stmts) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000928 ScopedIndent si(this);
929 return EmitStatements(stmts);
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000930}
931
932bool GeneratorImpl::EmitBlock(const ast::BlockStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000933 line() << "{";
934 if (!EmitStatementsWithIndent(stmt->statements)) {
935 return false;
936 }
937 line() << "}";
938 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000939}
940
941bool GeneratorImpl::EmitBreak(const ast::BreakStatement*) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000942 line() << "break;";
943 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000944}
945
dan sinclairb8b0c212022-10-20 22:45:50 +0000946bool GeneratorImpl::EmitBreakIf(const ast::BreakIfStatement* b) {
947 auto out = line();
948 out << "if (";
949 if (!EmitExpression(out, b->condition)) {
950 return false;
951 }
952 out << ") { break; }";
953 return true;
954}
955
dan sinclair41e4d9a2022-05-01 14:40:55 +0000956bool GeneratorImpl::EmitCall(std::ostream& out, const ast::CallExpression* expr) {
Ben Claytone9f8b092022-06-01 13:14:39 +0000957 auto* call = builder_.Sem().Get<sem::Call>(expr);
dan sinclair41e4d9a2022-05-01 14:40:55 +0000958 auto* target = call->Target();
959 return Switch(
960 target, [&](const sem::Function* func) { return EmitFunctionCall(out, call, func); },
961 [&](const sem::Builtin* builtin) { return EmitBuiltinCall(out, call, builtin); },
962 [&](const sem::TypeConversion* conv) { return EmitTypeConversion(out, call, conv); },
dan sinclair6e77b472022-10-20 13:38:28 +0000963 [&](const sem::TypeInitializer* ctor) { return EmitTypeInitializer(out, call, ctor); },
dan sinclair41e4d9a2022-05-01 14:40:55 +0000964 [&](Default) {
965 TINT_ICE(Writer, diagnostics_) << "unhandled call target: " << target->TypeInfo().name;
966 return false;
967 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000968}
969
970bool GeneratorImpl::EmitFunctionCall(std::ostream& out,
971 const sem::Call* call,
972 const sem::Function* func) {
dan sinclair41e4d9a2022-05-01 14:40:55 +0000973 auto* expr = call->Declaration();
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000974
dan sinclair41e4d9a2022-05-01 14:40:55 +0000975 if (ast::HasAttribute<transform::CalculateArrayLength::BufferSizeIntrinsic>(
976 func->Declaration()->attributes)) {
977 // Special function generated by the CalculateArrayLength transform for
978 // calling X.GetDimensions(Y)
979 if (!EmitExpression(out, call->Arguments()[0]->Declaration())) {
980 return false;
981 }
982 out << ".GetDimensions(";
983 if (!EmitExpression(out, call->Arguments()[1]->Declaration())) {
984 return false;
985 }
986 out << ")";
987 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +0000988 }
dan sinclair41e4d9a2022-05-01 14:40:55 +0000989
990 if (auto* intrinsic = ast::GetAttribute<transform::DecomposeMemoryAccess::Intrinsic>(
991 func->Declaration()->attributes)) {
dan sinclairff7cf212022-10-03 14:05:23 +0000992 switch (intrinsic->address_space) {
993 case ast::AddressSpace::kUniform:
dan sinclair41e4d9a2022-05-01 14:40:55 +0000994 return EmitUniformBufferAccess(out, expr, intrinsic);
dan sinclairff7cf212022-10-03 14:05:23 +0000995 case ast::AddressSpace::kStorage:
Antonio Maiorano08f4b552022-05-31 13:20:28 +0000996 if (!intrinsic->IsAtomic()) {
997 return EmitStorageBufferAccess(out, expr, intrinsic);
998 }
999 break;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001000 default:
1001 TINT_UNREACHABLE(Writer, diagnostics_)
dan sinclairff7cf212022-10-03 14:05:23 +00001002 << "unsupported DecomposeMemoryAccess::Intrinsic address space:"
1003 << intrinsic->address_space;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001004 return false;
1005 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001006 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001007
1008 out << builder_.Symbols().NameFor(func->Declaration()->symbol) << "(";
1009
1010 bool first = true;
1011 for (auto* arg : call->Arguments()) {
1012 if (!first) {
1013 out << ", ";
1014 }
1015 first = false;
1016
1017 if (!EmitExpression(out, arg->Declaration())) {
1018 return false;
1019 }
1020 }
1021
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001022 out << ")";
1023 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001024}
1025
1026bool GeneratorImpl::EmitBuiltinCall(std::ostream& out,
1027 const sem::Call* call,
1028 const sem::Builtin* builtin) {
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001029 const auto type = builtin->Type();
1030
dan sinclair41e4d9a2022-05-01 14:40:55 +00001031 auto* expr = call->Declaration();
1032 if (builtin->IsTexture()) {
1033 return EmitTextureCall(out, call, builtin);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001034 }
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001035 if (type == sem::BuiltinType::kSelect) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001036 return EmitSelectCall(out, expr);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001037 }
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001038 if (type == sem::BuiltinType::kModf) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001039 return EmitModfCall(out, expr, builtin);
1040 }
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001041 if (type == sem::BuiltinType::kFrexp) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001042 return EmitFrexpCall(out, expr, builtin);
1043 }
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001044 if (type == sem::BuiltinType::kDegrees) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001045 return EmitDegreesCall(out, expr, builtin);
1046 }
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001047 if (type == sem::BuiltinType::kRadians) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001048 return EmitRadiansCall(out, expr, builtin);
1049 }
1050 if (builtin->IsDataPacking()) {
1051 return EmitDataPackingCall(out, expr, builtin);
1052 }
1053 if (builtin->IsDataUnpacking()) {
1054 return EmitDataUnpackingCall(out, expr, builtin);
1055 }
1056 if (builtin->IsBarrier()) {
1057 return EmitBarrierCall(out, builtin);
1058 }
1059 if (builtin->IsAtomic()) {
1060 return EmitWorkgroupAtomicCall(out, expr, builtin);
1061 }
Jiawei Shaoab975702022-05-13 00:09:56 +00001062 if (builtin->IsDP4a()) {
1063 return EmitDP4aCall(out, expr, builtin);
1064 }
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001065
dan sinclair41e4d9a2022-05-01 14:40:55 +00001066 auto name = generate_builtin_name(builtin);
1067 if (name.empty()) {
1068 return false;
1069 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001070
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001071 // Handle single argument builtins that only accept and return uint (not int overload). We need
1072 // to explicitly cast the return value (we also cast the arg for good measure). See
1073 // crbug.com/tint/1550
1074 if (type == sem::BuiltinType::kCountOneBits || type == sem::BuiltinType::kReverseBits) {
1075 auto* arg = call->Arguments()[0];
1076 if (arg->Type()->UnwrapRef()->is_signed_scalar_or_vector()) {
1077 out << "asint(" << name << "(asuint(";
1078 if (!EmitExpression(out, arg->Declaration())) {
1079 return false;
1080 }
1081 out << ")))";
1082 return true;
1083 }
1084 }
1085
dan sinclair41e4d9a2022-05-01 14:40:55 +00001086 out << name << "(";
1087
1088 bool first = true;
1089 for (auto* arg : call->Arguments()) {
1090 if (!first) {
1091 out << ", ";
1092 }
1093 first = false;
1094
1095 if (!EmitExpression(out, arg->Declaration())) {
1096 return false;
1097 }
1098 }
1099
1100 out << ")";
Antonio Maioranoab4c0352022-05-20 01:58:40 +00001101
dan sinclair41e4d9a2022-05-01 14:40:55 +00001102 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001103}
1104
1105bool GeneratorImpl::EmitTypeConversion(std::ostream& out,
1106 const sem::Call* call,
1107 const sem::TypeConversion* conv) {
dan sinclairff7cf212022-10-03 14:05:23 +00001108 if (!EmitType(out, conv->Target(), ast::AddressSpace::kNone, ast::Access::kReadWrite, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001109 return false;
1110 }
1111 out << "(";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001112
dan sinclair41e4d9a2022-05-01 14:40:55 +00001113 if (!EmitExpression(out, call->Arguments()[0]->Declaration())) {
1114 return false;
1115 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001116
dan sinclair41e4d9a2022-05-01 14:40:55 +00001117 out << ")";
1118 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001119}
1120
dan sinclair6e77b472022-10-20 13:38:28 +00001121bool GeneratorImpl::EmitTypeInitializer(std::ostream& out,
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001122 const sem::Call* call,
dan sinclair6e77b472022-10-20 13:38:28 +00001123 const sem::TypeInitializer* ctor) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001124 auto* type = call->Type();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001125
dan sinclair6e77b472022-10-20 13:38:28 +00001126 // If the type initializer is empty then we need to construct with the zero
dan sinclair41e4d9a2022-05-01 14:40:55 +00001127 // value for all components.
Ben Clayton958a4642022-07-26 07:55:24 +00001128 if (call->Arguments().IsEmpty()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001129 return EmitZeroValue(out, type);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001130 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001131
dan sinclair6e77b472022-10-20 13:38:28 +00001132 // Single parameter matrix initializers must be identity initializer.
Zhaoming Jiangc5f7e8f2022-06-24 17:21:59 +00001133 // It could also be conversions between f16 and f32 matrix when f16 is properly supported.
Ben Clayton958a4642022-07-26 07:55:24 +00001134 if (type->Is<sem::Matrix>() && call->Arguments().Length() == 1) {
Zhaoming Jiangc5f7e8f2022-06-24 17:21:59 +00001135 if (!ctor->Parameters()[0]->Type()->UnwrapRef()->is_float_matrix()) {
1136 TINT_UNREACHABLE(Writer, diagnostics_)
dan sinclair6e77b472022-10-20 13:38:28 +00001137 << "found a single-parameter matrix initializer that is not identity initializer";
Zhaoming Jiangc5f7e8f2022-06-24 17:21:59 +00001138 return false;
Ben Clayton3b5edf12022-05-16 21:14:11 +00001139 }
1140 }
1141
dan sinclair41e4d9a2022-05-01 14:40:55 +00001142 bool brackets = type->IsAnyOf<sem::Array, sem::Struct>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001143
dan sinclair41e4d9a2022-05-01 14:40:55 +00001144 // For single-value vector initializers, swizzle the scalar to the right
1145 // vector dimension using .x
1146 const bool is_single_value_vector_init = type->is_scalar_vector() &&
Ben Clayton958a4642022-07-26 07:55:24 +00001147 call->Arguments().Length() == 1 &&
dan sinclair41e4d9a2022-05-01 14:40:55 +00001148 ctor->Parameters()[0]->Type()->is_scalar();
1149
Ben Clayton6c098ba2022-07-14 20:46:39 +00001150 if (brackets) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001151 out << "{";
1152 } else {
dan sinclairff7cf212022-10-03 14:05:23 +00001153 if (!EmitType(out, type, ast::AddressSpace::kNone, ast::Access::kReadWrite, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001154 return false;
1155 }
1156 out << "(";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001157 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001158
dan sinclair41e4d9a2022-05-01 14:40:55 +00001159 if (is_single_value_vector_init) {
1160 out << "(";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001161 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001162
dan sinclair41e4d9a2022-05-01 14:40:55 +00001163 bool first = true;
1164 for (auto* e : call->Arguments()) {
1165 if (!first) {
1166 out << ", ";
1167 }
1168 first = false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001169
dan sinclair41e4d9a2022-05-01 14:40:55 +00001170 if (!EmitExpression(out, e->Declaration())) {
1171 return false;
1172 }
1173 }
1174
1175 if (is_single_value_vector_init) {
1176 out << ")." << std::string(type->As<sem::Vector>()->Width(), 'x');
1177 }
1178
1179 out << (brackets ? "}" : ")");
1180 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001181}
1182
1183bool GeneratorImpl::EmitUniformBufferAccess(
1184 std::ostream& out,
1185 const ast::CallExpression* expr,
1186 const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001187 const auto& args = expr->args;
1188 auto* offset_arg = builder_.Sem().Get(args[1]);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001189
dan sinclair41e4d9a2022-05-01 14:40:55 +00001190 uint32_t scalar_offset_value = 0;
1191 std::string scalar_offset_expr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001192
dan sinclair41e4d9a2022-05-01 14:40:55 +00001193 // If true, use scalar_offset_value, otherwise use scalar_offset_expr
1194 bool scalar_offset_constant = false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001195
Ben Claytonaa037ac2022-06-29 19:07:30 +00001196 if (auto* val = offset_arg->ConstantValue()) {
1197 TINT_ASSERT(Writer, val->Type()->Is<sem::U32>());
1198 scalar_offset_value = static_cast<uint32_t>(std::get<AInt>(val->Value()));
dan sinclair41e4d9a2022-05-01 14:40:55 +00001199 scalar_offset_value /= 4; // bytes -> scalar index
1200 scalar_offset_constant = true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001201 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001202
dan sinclair41e4d9a2022-05-01 14:40:55 +00001203 if (!scalar_offset_constant) {
1204 // UBO offset not compile-time known.
1205 // Calculate the scalar offset into a temporary.
1206 scalar_offset_expr = UniqueIdentifier("scalar_offset");
1207 auto pre = line();
1208 pre << "const uint " << scalar_offset_expr << " = (";
1209 if (!EmitExpression(pre, args[1])) { // offset
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001210 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001211 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001212 pre << ") / 4;";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001213 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001214
1215 using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
1216 using DataType = transform::DecomposeMemoryAccess::Intrinsic::DataType;
1217 switch (intrinsic->op) {
1218 case Op::kLoad: {
1219 auto cast = [&](const char* to, auto&& load) {
1220 out << to << "(";
1221 auto result = load();
1222 out << ")";
1223 return result;
1224 };
1225 auto load_scalar = [&]() {
1226 if (!EmitExpression(out, args[0])) { // buffer
1227 return false;
1228 }
1229 if (scalar_offset_constant) {
1230 char swizzle[] = {'x', 'y', 'z', 'w'};
1231 out << "[" << (scalar_offset_value / 4) << "]."
1232 << swizzle[scalar_offset_value & 3];
1233 } else {
1234 out << "[" << scalar_offset_expr << " / 4][" << scalar_offset_expr << " % 4]";
1235 }
1236 return true;
1237 };
1238 // Has a minimum alignment of 8 bytes, so is either .xy or .zw
1239 auto load_vec2 = [&] {
1240 if (scalar_offset_constant) {
1241 if (!EmitExpression(out, args[0])) { // buffer
1242 return false;
1243 }
1244 out << "[" << (scalar_offset_value / 4) << "]";
1245 out << ((scalar_offset_value & 2) == 0 ? ".xy" : ".zw");
1246 } else {
1247 std::string ubo_load = UniqueIdentifier("ubo_load");
1248 {
1249 auto pre = line();
1250 pre << "uint4 " << ubo_load << " = ";
1251 if (!EmitExpression(pre, args[0])) { // buffer
1252 return false;
1253 }
1254 pre << "[" << scalar_offset_expr << " / 4];";
1255 }
1256 out << "((" << scalar_offset_expr << " & 2) ? " << ubo_load
1257 << ".zw : " << ubo_load << ".xy)";
1258 }
1259 return true;
1260 };
1261 // vec4 has a minimum alignment of 16 bytes, easiest case
1262 auto load_vec4 = [&] {
1263 if (!EmitExpression(out, args[0])) { // buffer
1264 return false;
1265 }
1266 if (scalar_offset_constant) {
1267 out << "[" << (scalar_offset_value / 4) << "]";
1268 } else {
1269 out << "[" << scalar_offset_expr << " / 4]";
1270 }
1271 return true;
1272 };
1273 // vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle
1274 auto load_vec3 = [&] {
1275 if (!load_vec4()) {
1276 return false;
1277 }
1278 out << ".xyz";
1279 return true;
1280 };
1281 switch (intrinsic->type) {
1282 case DataType::kU32:
1283 return load_scalar();
1284 case DataType::kF32:
1285 return cast("asfloat", load_scalar);
1286 case DataType::kI32:
1287 return cast("asint", load_scalar);
1288 case DataType::kVec2U32:
1289 return load_vec2();
1290 case DataType::kVec2F32:
1291 return cast("asfloat", load_vec2);
1292 case DataType::kVec2I32:
1293 return cast("asint", load_vec2);
1294 case DataType::kVec3U32:
1295 return load_vec3();
1296 case DataType::kVec3F32:
1297 return cast("asfloat", load_vec3);
1298 case DataType::kVec3I32:
1299 return cast("asint", load_vec3);
1300 case DataType::kVec4U32:
1301 return load_vec4();
1302 case DataType::kVec4F32:
1303 return cast("asfloat", load_vec4);
1304 case DataType::kVec4I32:
1305 return cast("asint", load_vec4);
1306 }
1307 TINT_UNREACHABLE(Writer, diagnostics_)
1308 << "unsupported DecomposeMemoryAccess::Intrinsic::DataType: "
1309 << static_cast<int>(intrinsic->type);
1310 return false;
1311 }
1312 default:
1313 break;
1314 }
1315 TINT_UNREACHABLE(Writer, diagnostics_)
1316 << "unsupported DecomposeMemoryAccess::Intrinsic::Op: " << static_cast<int>(intrinsic->op);
1317 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001318}
1319
1320bool GeneratorImpl::EmitStorageBufferAccess(
1321 std::ostream& out,
1322 const ast::CallExpression* expr,
1323 const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001324 const auto& args = expr->args;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001325
dan sinclair41e4d9a2022-05-01 14:40:55 +00001326 using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
1327 using DataType = transform::DecomposeMemoryAccess::Intrinsic::DataType;
1328 switch (intrinsic->op) {
1329 case Op::kLoad: {
1330 auto load = [&](const char* cast, int n) {
1331 if (cast) {
1332 out << cast << "(";
1333 }
1334 if (!EmitExpression(out, args[0])) { // buffer
1335 return false;
1336 }
1337 out << ".Load";
1338 if (n > 1) {
1339 out << n;
1340 }
1341 ScopedParen sp(out);
1342 if (!EmitExpression(out, args[1])) { // offset
1343 return false;
1344 }
1345 if (cast) {
1346 out << ")";
1347 }
1348 return true;
1349 };
1350 switch (intrinsic->type) {
1351 case DataType::kU32:
1352 return load(nullptr, 1);
1353 case DataType::kF32:
1354 return load("asfloat", 1);
1355 case DataType::kI32:
1356 return load("asint", 1);
1357 case DataType::kVec2U32:
1358 return load(nullptr, 2);
1359 case DataType::kVec2F32:
1360 return load("asfloat", 2);
1361 case DataType::kVec2I32:
1362 return load("asint", 2);
1363 case DataType::kVec3U32:
1364 return load(nullptr, 3);
1365 case DataType::kVec3F32:
1366 return load("asfloat", 3);
1367 case DataType::kVec3I32:
1368 return load("asint", 3);
1369 case DataType::kVec4U32:
1370 return load(nullptr, 4);
1371 case DataType::kVec4F32:
1372 return load("asfloat", 4);
1373 case DataType::kVec4I32:
1374 return load("asint", 4);
1375 }
1376 TINT_UNREACHABLE(Writer, diagnostics_)
1377 << "unsupported DecomposeMemoryAccess::Intrinsic::DataType: "
1378 << static_cast<int>(intrinsic->type);
1379 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001380 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001381
1382 case Op::kStore: {
1383 auto store = [&](int n) {
1384 if (!EmitExpression(out, args[0])) { // buffer
1385 return false;
1386 }
1387 out << ".Store";
1388 if (n > 1) {
1389 out << n;
1390 }
1391 ScopedParen sp1(out);
1392 if (!EmitExpression(out, args[1])) { // offset
1393 return false;
1394 }
1395 out << ", asuint";
1396 ScopedParen sp2(out);
1397 if (!EmitExpression(out, args[2])) { // value
1398 return false;
1399 }
1400 return true;
1401 };
1402 switch (intrinsic->type) {
1403 case DataType::kU32:
1404 return store(1);
1405 case DataType::kF32:
1406 return store(1);
1407 case DataType::kI32:
1408 return store(1);
1409 case DataType::kVec2U32:
1410 return store(2);
1411 case DataType::kVec2F32:
1412 return store(2);
1413 case DataType::kVec2I32:
1414 return store(2);
1415 case DataType::kVec3U32:
1416 return store(3);
1417 case DataType::kVec3F32:
1418 return store(3);
1419 case DataType::kVec3I32:
1420 return store(3);
1421 case DataType::kVec4U32:
1422 return store(4);
1423 case DataType::kVec4F32:
1424 return store(4);
1425 case DataType::kVec4I32:
1426 return store(4);
1427 }
1428 TINT_UNREACHABLE(Writer, diagnostics_)
1429 << "unsupported DecomposeMemoryAccess::Intrinsic::DataType: "
1430 << static_cast<int>(intrinsic->type);
1431 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001432 }
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001433 default:
1434 // Break out to error case below/
1435 // Note that atomic intrinsics are generated as functions.
1436 break;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001437 }
1438
dan sinclair41e4d9a2022-05-01 14:40:55 +00001439 TINT_UNREACHABLE(Writer, diagnostics_)
1440 << "unsupported DecomposeMemoryAccess::Intrinsic::Op: " << static_cast<int>(intrinsic->op);
1441 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001442}
1443
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001444bool GeneratorImpl::EmitStorageAtomicIntrinsic(
1445 const ast::Function* func,
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001446 const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001447 using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001448
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001449 const sem::Function* sem_func = builder_.Sem().Get(func);
1450 auto* result_ty = sem_func->ReturnType();
1451 const auto& params = sem_func->Parameters();
1452 const auto name = builder_.Symbols().NameFor(func->symbol);
1453 auto& buf = *current_buffer_;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001454
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001455 auto rmw = [&](const char* hlsl) -> bool {
1456 {
1457 auto fn = line(&buf);
Ben Claytond2e0db32022-10-12 18:49:15 +00001458 if (!EmitTypeAndName(fn, result_ty, ast::AddressSpace::kNone, ast::Access::kUndefined,
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001459 name)) {
1460 return false;
1461 }
1462 fn << "(RWByteAddressBuffer buffer, uint offset, ";
Ben Claytond2e0db32022-10-12 18:49:15 +00001463 if (!EmitTypeAndName(fn, result_ty, ast::AddressSpace::kNone, ast::Access::kUndefined,
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001464 "value")) {
1465 return false;
1466 }
1467 fn << ") {";
1468 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001469
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001470 buf.IncrementIndent();
1471 TINT_DEFER({
1472 buf.DecrementIndent();
1473 line(&buf) << "}";
1474 line(&buf);
1475 });
1476
1477 {
1478 auto l = line(&buf);
Ben Claytond2e0db32022-10-12 18:49:15 +00001479 if (!EmitTypeAndName(l, result_ty, ast::AddressSpace::kNone, ast::Access::kUndefined,
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001480 "original_value")) {
1481 return false;
1482 }
1483 l << " = 0;";
1484 }
1485 {
1486 auto l = line(&buf);
1487 l << "buffer." << hlsl << "(offset, ";
1488 if (intrinsic->op == Op::kAtomicSub) {
1489 l << "-";
1490 }
1491 l << "value, original_value);";
1492 }
1493 line(&buf) << "return original_value;";
1494 return true;
1495 };
1496
1497 switch (intrinsic->op) {
1498 case Op::kAtomicAdd:
1499 return rmw("InterlockedAdd");
1500
1501 case Op::kAtomicSub:
1502 // Use add with the operand negated.
1503 return rmw("InterlockedAdd");
1504
1505 case Op::kAtomicMax:
1506 return rmw("InterlockedMax");
1507
1508 case Op::kAtomicMin:
1509 return rmw("InterlockedMin");
1510
1511 case Op::kAtomicAnd:
1512 return rmw("InterlockedAnd");
1513
1514 case Op::kAtomicOr:
1515 return rmw("InterlockedOr");
1516
1517 case Op::kAtomicXor:
1518 return rmw("InterlockedXor");
1519
1520 case Op::kAtomicExchange:
1521 return rmw("InterlockedExchange");
1522
1523 case Op::kAtomicLoad: {
1524 // HLSL does not have an InterlockedLoad, so we emulate it with
1525 // InterlockedOr using 0 as the OR value
dan sinclair41e4d9a2022-05-01 14:40:55 +00001526 {
1527 auto fn = line(&buf);
Ben Claytond2e0db32022-10-12 18:49:15 +00001528 if (!EmitTypeAndName(fn, result_ty, ast::AddressSpace::kNone,
1529 ast::Access::kUndefined, name)) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001530 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001531 }
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001532 fn << "(RWByteAddressBuffer buffer, uint offset) {";
1533 }
1534
1535 buf.IncrementIndent();
1536 TINT_DEFER({
1537 buf.DecrementIndent();
1538 line(&buf) << "}";
1539 line(&buf);
1540 });
1541
1542 {
1543 auto l = line(&buf);
Ben Claytond2e0db32022-10-12 18:49:15 +00001544 if (!EmitTypeAndName(l, result_ty, ast::AddressSpace::kNone,
1545 ast::Access::kUndefined, "value")) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001546 return false;
1547 }
1548 l << " = 0;";
1549 }
1550
1551 line(&buf) << "buffer.InterlockedOr(offset, 0, value);";
1552 line(&buf) << "return value;";
1553 return true;
1554 }
1555 case Op::kAtomicStore: {
1556 // HLSL does not have an InterlockedStore, so we emulate it with
1557 // InterlockedExchange and discard the returned value
1558 auto* value_ty = params[2]->Type()->UnwrapRef();
1559 {
1560 auto fn = line(&buf);
1561 fn << "void " << name << "(RWByteAddressBuffer buffer, uint offset, ";
Ben Claytond2e0db32022-10-12 18:49:15 +00001562 if (!EmitTypeAndName(fn, value_ty, ast::AddressSpace::kNone,
1563 ast::Access::kUndefined, "value")) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001564 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001565 }
1566 fn << ") {";
1567 }
1568
1569 buf.IncrementIndent();
1570 TINT_DEFER({
1571 buf.DecrementIndent();
1572 line(&buf) << "}";
1573 line(&buf);
1574 });
1575
1576 {
1577 auto l = line(&buf);
Ben Claytond2e0db32022-10-12 18:49:15 +00001578 if (!EmitTypeAndName(l, value_ty, ast::AddressSpace::kNone, ast::Access::kUndefined,
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001579 "ignored")) {
1580 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001581 }
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001582 l << ";";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001583 }
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001584 line(&buf) << "buffer.InterlockedExchange(offset, value, ignored);";
1585 return true;
1586 }
1587 case Op::kAtomicCompareExchangeWeak: {
1588 // NOTE: We don't need to emit the return type struct here as DecomposeMemoryAccess
1589 // already added it to the AST, and it should have already been emitted by now.
1590 auto* value_ty = params[2]->Type()->UnwrapRef();
dan sinclair41e4d9a2022-05-01 14:40:55 +00001591 {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001592 auto fn = line(&buf);
Ben Claytond2e0db32022-10-12 18:49:15 +00001593 if (!EmitTypeAndName(fn, result_ty, ast::AddressSpace::kNone,
1594 ast::Access::kUndefined, name)) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001595 return false;
1596 }
1597 fn << "(RWByteAddressBuffer buffer, uint offset, ";
Ben Claytond2e0db32022-10-12 18:49:15 +00001598 if (!EmitTypeAndName(fn, value_ty, ast::AddressSpace::kNone,
1599 ast::Access::kUndefined, "compare")) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001600 return false;
1601 }
1602 fn << ", ";
Ben Claytond2e0db32022-10-12 18:49:15 +00001603 if (!EmitTypeAndName(fn, value_ty, ast::AddressSpace::kNone,
1604 ast::Access::kUndefined, "value")) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001605 return false;
1606 }
1607 fn << ") {";
1608 }
1609
1610 buf.IncrementIndent();
1611 TINT_DEFER({
1612 buf.DecrementIndent();
1613 line(&buf) << "}";
1614 line(&buf);
1615 });
1616
1617 { // T result = {0};
dan sinclair41e4d9a2022-05-01 14:40:55 +00001618 auto l = line(&buf);
Ben Claytond2e0db32022-10-12 18:49:15 +00001619 if (!EmitTypeAndName(l, result_ty, ast::AddressSpace::kNone,
1620 ast::Access::kUndefined, "result")) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001621 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00001622 }
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001623 l << "=";
1624 if (!EmitZeroValue(l, result_ty)) {
1625 return false;
1626 }
1627 l << ";";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001628 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001629
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001630 line(&buf) << "buffer.InterlockedCompareExchange(offset, compare, value, "
1631 "result.old_value);";
1632 line(&buf) << "result.exchanged = result.old_value == compare;";
1633 line(&buf) << "return result;";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001634
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001635 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001636 }
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001637 default:
1638 break;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001639 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001640
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001641 TINT_UNREACHABLE(Writer, diagnostics_)
1642 << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
1643 << static_cast<int>(intrinsic->op);
1644 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001645}
1646
1647bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
1648 const ast::CallExpression* expr,
1649 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001650 std::string result = UniqueIdentifier("atomic_result");
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001651
dan sinclair41e4d9a2022-05-01 14:40:55 +00001652 if (!builtin->ReturnType()->Is<sem::Void>()) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001653 auto pre = line();
dan sinclairff7cf212022-10-03 14:05:23 +00001654 if (!EmitTypeAndName(pre, builtin->ReturnType(), ast::AddressSpace::kNone,
Ben Claytond2e0db32022-10-12 18:49:15 +00001655 ast::Access::kUndefined, result)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001656 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001657 }
1658 pre << " = ";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001659 if (!EmitZeroValue(pre, builtin->ReturnType())) {
1660 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001661 }
1662 pre << ";";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001663 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001664
dan sinclair41e4d9a2022-05-01 14:40:55 +00001665 auto call = [&](const char* name) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001666 auto pre = line();
dan sinclair41e4d9a2022-05-01 14:40:55 +00001667 pre << name;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001668
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001669 {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001670 ScopedParen sp(pre);
Ben Clayton783b1692022-08-02 17:03:35 +00001671 for (size_t i = 0; i < expr->args.Length(); i++) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001672 auto* arg = expr->args[i];
1673 if (i > 0) {
1674 pre << ", ";
1675 }
1676 if (i == 1 && builtin->Type() == sem::BuiltinType::kAtomicSub) {
1677 // Sub uses InterlockedAdd with the operand negated.
1678 pre << "-";
1679 }
1680 if (!EmitExpression(pre, arg)) {
1681 return false;
1682 }
1683 }
1684
1685 pre << ", " << result;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001686 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001687
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001688 pre << ";";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001689
dan sinclair41e4d9a2022-05-01 14:40:55 +00001690 out << result;
1691 return true;
1692 };
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001693
dan sinclair41e4d9a2022-05-01 14:40:55 +00001694 switch (builtin->Type()) {
1695 case sem::BuiltinType::kAtomicLoad: {
1696 // HLSL does not have an InterlockedLoad, so we emulate it with
1697 // InterlockedOr using 0 as the OR value
1698 auto pre = line();
1699 pre << "InterlockedOr";
1700 {
1701 ScopedParen sp(pre);
1702 if (!EmitExpression(pre, expr->args[0])) {
1703 return false;
1704 }
1705 pre << ", 0, " << result;
1706 }
1707 pre << ";";
1708
1709 out << result;
1710 return true;
1711 }
1712 case sem::BuiltinType::kAtomicStore: {
1713 // HLSL does not have an InterlockedStore, so we emulate it with
1714 // InterlockedExchange and discard the returned value
1715 { // T result = 0;
1716 auto pre = line();
1717 auto* value_ty = builtin->Parameters()[1]->Type()->UnwrapRef();
Ben Claytond2e0db32022-10-12 18:49:15 +00001718 if (!EmitTypeAndName(pre, value_ty, ast::AddressSpace::kNone,
1719 ast::Access::kUndefined, result)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001720 return false;
1721 }
1722 pre << " = ";
1723 if (!EmitZeroValue(pre, value_ty)) {
1724 return false;
1725 }
1726 pre << ";";
1727 }
1728
1729 out << "InterlockedExchange";
1730 {
1731 ScopedParen sp(out);
1732 if (!EmitExpression(out, expr->args[0])) {
1733 return false;
1734 }
1735 out << ", ";
1736 if (!EmitExpression(out, expr->args[1])) {
1737 return false;
1738 }
1739 out << ", " << result;
1740 }
1741 return true;
1742 }
1743 case sem::BuiltinType::kAtomicCompareExchangeWeak: {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001744 // Emit the builtin return type unique to this overload. This does not
1745 // exist in the AST, so it will not be generated in Generate().
Antonio Maioranof25140f2022-06-03 14:47:01 +00001746 if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001747 return false;
1748 }
1749
dan sinclair41e4d9a2022-05-01 14:40:55 +00001750 auto* dest = expr->args[0];
1751 auto* compare_value = expr->args[1];
1752 auto* value = expr->args[2];
1753
1754 std::string compare = UniqueIdentifier("atomic_compare_value");
1755
1756 { // T compare_value = <compare_value>;
1757 auto pre = line();
Antonio Maioranof99671b2022-06-23 13:14:54 +00001758 if (!EmitTypeAndName(pre, TypeOf(compare_value)->UnwrapRef(),
Ben Claytond2e0db32022-10-12 18:49:15 +00001759 ast::AddressSpace::kNone, ast::Access::kUndefined, compare)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001760 return false;
1761 }
1762 pre << " = ";
1763 if (!EmitExpression(pre, compare_value)) {
1764 return false;
1765 }
1766 pre << ";";
1767 }
1768
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001769 { // InterlockedCompareExchange(dst, compare, value, result.old_value);
dan sinclair41e4d9a2022-05-01 14:40:55 +00001770 auto pre = line();
1771 pre << "InterlockedCompareExchange";
1772 {
1773 ScopedParen sp(pre);
1774 if (!EmitExpression(pre, dest)) {
1775 return false;
1776 }
1777 pre << ", " << compare << ", ";
1778 if (!EmitExpression(pre, value)) {
1779 return false;
1780 }
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001781 pre << ", " << result << ".old_value";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001782 }
1783 pre << ";";
1784 }
1785
Antonio Maiorano08f4b552022-05-31 13:20:28 +00001786 // result.exchanged = result.old_value == compare;
1787 line() << result << ".exchanged = " << result << ".old_value == " << compare << ";";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001788
1789 out << result;
1790 return true;
1791 }
1792
1793 case sem::BuiltinType::kAtomicAdd:
1794 case sem::BuiltinType::kAtomicSub:
1795 return call("InterlockedAdd");
1796
1797 case sem::BuiltinType::kAtomicMax:
1798 return call("InterlockedMax");
1799
1800 case sem::BuiltinType::kAtomicMin:
1801 return call("InterlockedMin");
1802
1803 case sem::BuiltinType::kAtomicAnd:
1804 return call("InterlockedAnd");
1805
1806 case sem::BuiltinType::kAtomicOr:
1807 return call("InterlockedOr");
1808
1809 case sem::BuiltinType::kAtomicXor:
1810 return call("InterlockedXor");
1811
1812 case sem::BuiltinType::kAtomicExchange:
1813 return call("InterlockedExchange");
1814
1815 default:
1816 break;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001817 }
1818
dan sinclair41e4d9a2022-05-01 14:40:55 +00001819 TINT_UNREACHABLE(Writer, diagnostics_) << "unsupported atomic builtin: " << builtin->Type();
1820 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001821}
1822
dan sinclair41e4d9a2022-05-01 14:40:55 +00001823bool GeneratorImpl::EmitSelectCall(std::ostream& out, const ast::CallExpression* expr) {
1824 auto* expr_false = expr->args[0];
1825 auto* expr_true = expr->args[1];
1826 auto* expr_cond = expr->args[2];
1827 ScopedParen paren(out);
1828 if (!EmitExpression(out, expr_cond)) {
1829 return false;
1830 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001831
dan sinclair41e4d9a2022-05-01 14:40:55 +00001832 out << " ? ";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001833
dan sinclair41e4d9a2022-05-01 14:40:55 +00001834 if (!EmitExpression(out, expr_true)) {
1835 return false;
1836 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001837
dan sinclair41e4d9a2022-05-01 14:40:55 +00001838 out << " : ";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001839
dan sinclair41e4d9a2022-05-01 14:40:55 +00001840 if (!EmitExpression(out, expr_false)) {
1841 return false;
1842 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001843
dan sinclair41e4d9a2022-05-01 14:40:55 +00001844 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001845}
1846
1847bool GeneratorImpl::EmitModfCall(std::ostream& out,
1848 const ast::CallExpression* expr,
1849 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001850 return CallBuiltinHelper(
1851 out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
1852 auto* ty = builtin->Parameters()[0]->Type();
1853 auto in = params[0];
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001854
dan sinclair41e4d9a2022-05-01 14:40:55 +00001855 std::string width;
1856 if (auto* vec = ty->As<sem::Vector>()) {
1857 width = std::to_string(vec->Width());
1858 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001859
dan sinclair41e4d9a2022-05-01 14:40:55 +00001860 // Emit the builtin return type unique to this overload. This does not
1861 // exist in the AST, so it will not be generated in Generate().
1862 if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
1863 return false;
1864 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001865
dan sinclair41e4d9a2022-05-01 14:40:55 +00001866 {
1867 auto l = line(b);
dan sinclairff7cf212022-10-03 14:05:23 +00001868 if (!EmitType(l, builtin->ReturnType(), ast::AddressSpace::kNone,
Ben Claytond2e0db32022-10-12 18:49:15 +00001869 ast::Access::kUndefined, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001870 return false;
1871 }
Zhaoming Jiang20cddbf2022-08-05 15:11:44 +00001872 l << " result;";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001873 }
Zhaoming Jiang20cddbf2022-08-05 15:11:44 +00001874 line(b) << "result.fract = modf(" << params[0] << ", result.whole);";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001875 line(b) << "return result;";
1876 return true;
1877 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001878}
1879
1880bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
1881 const ast::CallExpression* expr,
1882 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001883 return CallBuiltinHelper(
1884 out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
1885 auto* ty = builtin->Parameters()[0]->Type();
1886 auto in = params[0];
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001887
dan sinclair41e4d9a2022-05-01 14:40:55 +00001888 std::string width;
1889 if (auto* vec = ty->As<sem::Vector>()) {
1890 width = std::to_string(vec->Width());
1891 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001892
dan sinclair41e4d9a2022-05-01 14:40:55 +00001893 // Emit the builtin return type unique to this overload. This does not
1894 // exist in the AST, so it will not be generated in Generate().
1895 if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
1896 return false;
1897 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001898
Zhaoming Jiang20cddbf2022-08-05 15:11:44 +00001899 std::string member_type;
1900 if (Is<sem::F16>(sem::Type::DeepestElementOf(ty))) {
1901 member_type = width.empty() ? "float16_t" : ("vector<float16_t, " + width + ">");
1902 } else {
1903 member_type = "float" + width;
1904 }
1905
1906 line(b) << member_type << " exp;";
1907 line(b) << member_type << " sig = frexp(" << in << ", exp);";
dan sinclair41e4d9a2022-05-01 14:40:55 +00001908 {
1909 auto l = line(b);
dan sinclairff7cf212022-10-03 14:05:23 +00001910 if (!EmitType(l, builtin->ReturnType(), ast::AddressSpace::kNone,
Ben Claytond2e0db32022-10-12 18:49:15 +00001911 ast::Access::kUndefined, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001912 return false;
1913 }
1914 l << " result = {sig, int" << width << "(exp)};";
1915 }
1916 line(b) << "return result;";
1917 return true;
1918 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001919}
1920
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001921bool GeneratorImpl::EmitDegreesCall(std::ostream& out,
1922 const ast::CallExpression* expr,
1923 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001924 return CallBuiltinHelper(out, expr, builtin,
1925 [&](TextBuffer* b, const std::vector<std::string>& params) {
1926 line(b) << "return " << params[0] << " * " << std::setprecision(20)
1927 << sem::kRadToDeg << ";";
1928 return true;
1929 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001930}
1931
1932bool GeneratorImpl::EmitRadiansCall(std::ostream& out,
1933 const ast::CallExpression* expr,
1934 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001935 return CallBuiltinHelper(out, expr, builtin,
1936 [&](TextBuffer* b, const std::vector<std::string>& params) {
1937 line(b) << "return " << params[0] << " * " << std::setprecision(20)
1938 << sem::kDegToRad << ";";
1939 return true;
1940 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001941}
1942
1943bool GeneratorImpl::EmitDataPackingCall(std::ostream& out,
1944 const ast::CallExpression* expr,
1945 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001946 return CallBuiltinHelper(
1947 out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
1948 uint32_t dims = 2;
1949 bool is_signed = false;
1950 uint32_t scale = 65535;
Ben Clayton73683022022-10-06 19:23:29 +00001951 if (builtin->Type() == sem::BuiltinType::kPack4X8Snorm ||
1952 builtin->Type() == sem::BuiltinType::kPack4X8Unorm) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001953 dims = 4;
1954 scale = 255;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001955 }
Ben Clayton73683022022-10-06 19:23:29 +00001956 if (builtin->Type() == sem::BuiltinType::kPack4X8Snorm ||
1957 builtin->Type() == sem::BuiltinType::kPack2X16Snorm) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001958 is_signed = true;
1959 scale = (scale - 1) / 2;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00001960 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00001961 switch (builtin->Type()) {
Ben Clayton73683022022-10-06 19:23:29 +00001962 case sem::BuiltinType::kPack4X8Snorm:
1963 case sem::BuiltinType::kPack4X8Unorm:
1964 case sem::BuiltinType::kPack2X16Snorm:
1965 case sem::BuiltinType::kPack2X16Unorm: {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001966 {
1967 auto l = line(b);
1968 l << (is_signed ? "" : "u") << "int" << dims
1969 << " i = " << (is_signed ? "" : "u") << "int" << dims << "(round(clamp("
1970 << params[0] << ", " << (is_signed ? "-1.0" : "0.0") << ", 1.0) * "
1971 << scale << ".0))";
1972 if (is_signed) {
1973 l << " & " << (dims == 4 ? "0xff" : "0xffff");
1974 }
1975 l << ";";
1976 }
1977 {
1978 auto l = line(b);
1979 l << "return ";
1980 if (is_signed) {
1981 l << "asuint";
1982 }
1983 l << "(i.x | i.y << " << (32 / dims);
1984 if (dims == 4) {
1985 l << " | i.z << 16 | i.w << 24";
1986 }
1987 l << ");";
1988 }
1989 break;
1990 }
Ben Clayton73683022022-10-06 19:23:29 +00001991 case sem::BuiltinType::kPack2X16Float: {
dan sinclair41e4d9a2022-05-01 14:40:55 +00001992 line(b) << "uint2 i = f32tof16(" << params[0] << ");";
1993 line(b) << "return i.x | (i.y << 16);";
1994 break;
1995 }
1996 default:
1997 diagnostics_.add_error(diag::System::Writer,
1998 "Internal error: unhandled data packing builtin");
1999 return false;
2000 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002001
dan sinclair41e4d9a2022-05-01 14:40:55 +00002002 return true;
2003 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002004}
2005
2006bool GeneratorImpl::EmitDataUnpackingCall(std::ostream& out,
2007 const ast::CallExpression* expr,
2008 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002009 return CallBuiltinHelper(
2010 out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
2011 uint32_t dims = 2;
2012 bool is_signed = false;
2013 uint32_t scale = 65535;
Ben Clayton73683022022-10-06 19:23:29 +00002014 if (builtin->Type() == sem::BuiltinType::kUnpack4X8Snorm ||
2015 builtin->Type() == sem::BuiltinType::kUnpack4X8Unorm) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002016 dims = 4;
2017 scale = 255;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002018 }
Ben Clayton73683022022-10-06 19:23:29 +00002019 if (builtin->Type() == sem::BuiltinType::kUnpack4X8Snorm ||
2020 builtin->Type() == sem::BuiltinType::kUnpack2X16Snorm) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002021 is_signed = true;
2022 scale = (scale - 1) / 2;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002023 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002024 switch (builtin->Type()) {
Ben Clayton73683022022-10-06 19:23:29 +00002025 case sem::BuiltinType::kUnpack4X8Snorm:
2026 case sem::BuiltinType::kUnpack2X16Snorm: {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002027 line(b) << "int j = int(" << params[0] << ");";
2028 { // Perform sign extension on the converted values.
2029 auto l = line(b);
2030 l << "int" << dims << " i = int" << dims << "(";
2031 if (dims == 2) {
2032 l << "j << 16, j) >> 16";
2033 } else {
2034 l << "j << 24, j << 16, j << 8, j) >> 24";
2035 }
2036 l << ";";
2037 }
2038 line(b) << "return clamp(float" << dims << "(i) / " << scale << ".0, "
2039 << (is_signed ? "-1.0" : "0.0") << ", 1.0);";
2040 break;
2041 }
Ben Clayton73683022022-10-06 19:23:29 +00002042 case sem::BuiltinType::kUnpack4X8Unorm:
2043 case sem::BuiltinType::kUnpack2X16Unorm: {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002044 line(b) << "uint j = " << params[0] << ";";
2045 {
2046 auto l = line(b);
2047 l << "uint" << dims << " i = uint" << dims << "(";
2048 l << "j & " << (dims == 2 ? "0xffff" : "0xff") << ", ";
2049 if (dims == 4) {
2050 l << "(j >> " << (32 / dims) << ") & 0xff, (j >> 16) & 0xff, j >> 24";
2051 } else {
2052 l << "j >> " << (32 / dims);
2053 }
2054 l << ");";
2055 }
2056 line(b) << "return float" << dims << "(i) / " << scale << ".0;";
2057 break;
2058 }
Ben Clayton73683022022-10-06 19:23:29 +00002059 case sem::BuiltinType::kUnpack2X16Float:
dan sinclair41e4d9a2022-05-01 14:40:55 +00002060 line(b) << "uint i = " << params[0] << ";";
2061 line(b) << "return f16tof32(uint2(i & 0xffff, i >> 16));";
2062 break;
2063 default:
2064 diagnostics_.add_error(diag::System::Writer,
2065 "Internal error: unhandled data packing builtin");
2066 return false;
2067 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002068
dan sinclair41e4d9a2022-05-01 14:40:55 +00002069 return true;
2070 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002071}
2072
Jiawei Shaoab975702022-05-13 00:09:56 +00002073bool GeneratorImpl::EmitDP4aCall(std::ostream& out,
2074 const ast::CallExpression* expr,
2075 const sem::Builtin* builtin) {
2076 // TODO(crbug.com/tint/1497): support the polyfill version of DP4a functions.
2077 return CallBuiltinHelper(
2078 out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
2079 std::string functionName;
2080 switch (builtin->Type()) {
2081 case sem::BuiltinType::kDot4I8Packed:
Jiawei Shao1c759212022-05-15 13:53:21 +00002082 line(b) << "int accumulator = 0;";
Jiawei Shaoab975702022-05-13 00:09:56 +00002083 functionName = "dot4add_i8packed";
2084 break;
2085 case sem::BuiltinType::kDot4U8Packed:
Jiawei Shao1c759212022-05-15 13:53:21 +00002086 line(b) << "uint accumulator = 0u;";
Jiawei Shaoab975702022-05-13 00:09:56 +00002087 functionName = "dot4add_u8packed";
2088 break;
2089 default:
2090 diagnostics_.add_error(diag::System::Writer,
2091 "Internal error: unhandled DP4a builtin");
2092 return false;
2093 }
2094 line(b) << "return " << functionName << "(" << params[0] << ", " << params[1]
Jiawei Shao1c759212022-05-15 13:53:21 +00002095 << ", accumulator);";
Jiawei Shaoab975702022-05-13 00:09:56 +00002096
2097 return true;
2098 });
2099}
2100
dan sinclair41e4d9a2022-05-01 14:40:55 +00002101bool GeneratorImpl::EmitBarrierCall(std::ostream& out, const sem::Builtin* builtin) {
2102 // TODO(crbug.com/tint/661): Combine sequential barriers to a single
2103 // instruction.
2104 if (builtin->Type() == sem::BuiltinType::kWorkgroupBarrier) {
2105 out << "GroupMemoryBarrierWithGroupSync()";
2106 } else if (builtin->Type() == sem::BuiltinType::kStorageBarrier) {
2107 out << "DeviceMemoryBarrierWithGroupSync()";
2108 } else {
2109 TINT_UNREACHABLE(Writer, diagnostics_)
2110 << "unexpected barrier builtin type " << sem::str(builtin->Type());
2111 return false;
2112 }
2113 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002114}
2115
2116bool GeneratorImpl::EmitTextureCall(std::ostream& out,
2117 const sem::Call* call,
2118 const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002119 using Usage = sem::ParameterUsage;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002120
dan sinclair41e4d9a2022-05-01 14:40:55 +00002121 auto& signature = builtin->Signature();
2122 auto* expr = call->Declaration();
2123 auto arguments = expr->args;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002124
dan sinclair41e4d9a2022-05-01 14:40:55 +00002125 // Returns the argument with the given usage
2126 auto arg = [&](Usage usage) {
2127 int idx = signature.IndexOf(usage);
dan sinclair3a2a2792022-06-29 14:38:15 +00002128 return (idx >= 0) ? arguments[static_cast<size_t>(idx)] : nullptr;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002129 };
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002130
dan sinclair41e4d9a2022-05-01 14:40:55 +00002131 auto* texture = arg(Usage::kTexture);
2132 if (!texture) {
2133 TINT_ICE(Writer, diagnostics_) << "missing texture argument";
2134 return false;
2135 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002136
dan sinclair41e4d9a2022-05-01 14:40:55 +00002137 auto* texture_type = TypeOf(texture)->UnwrapRef()->As<sem::Texture>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002138
dan sinclair41e4d9a2022-05-01 14:40:55 +00002139 switch (builtin->Type()) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002140 case sem::BuiltinType::kTextureDimensions:
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002141 case sem::BuiltinType::kTextureNumLayers:
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002142 case sem::BuiltinType::kTextureNumLevels:
dan sinclair41e4d9a2022-05-01 14:40:55 +00002143 case sem::BuiltinType::kTextureNumSamples: {
2144 // All of these builtins use the GetDimensions() method on the texture
2145 bool is_ms =
2146 texture_type->IsAnyOf<sem::MultisampledTexture, sem::DepthMultisampledTexture>();
2147 int num_dimensions = 0;
2148 std::string swizzle;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002149
dan sinclair41e4d9a2022-05-01 14:40:55 +00002150 switch (builtin->Type()) {
2151 case sem::BuiltinType::kTextureDimensions:
2152 switch (texture_type->dim()) {
2153 case ast::TextureDimension::kNone:
2154 TINT_ICE(Writer, diagnostics_) << "texture dimension is kNone";
2155 return false;
2156 case ast::TextureDimension::k1d:
2157 num_dimensions = 1;
2158 break;
2159 case ast::TextureDimension::k2d:
2160 num_dimensions = is_ms ? 3 : 2;
2161 swizzle = is_ms ? ".xy" : "";
2162 break;
2163 case ast::TextureDimension::k2dArray:
2164 num_dimensions = is_ms ? 4 : 3;
2165 swizzle = ".xy";
2166 break;
2167 case ast::TextureDimension::k3d:
2168 num_dimensions = 3;
2169 break;
2170 case ast::TextureDimension::kCube:
2171 num_dimensions = 2;
2172 break;
2173 case ast::TextureDimension::kCubeArray:
2174 num_dimensions = 3;
2175 swizzle = ".xy";
2176 break;
2177 }
2178 break;
2179 case sem::BuiltinType::kTextureNumLayers:
2180 switch (texture_type->dim()) {
2181 default:
2182 TINT_ICE(Writer, diagnostics_) << "texture dimension is not arrayed";
2183 return false;
2184 case ast::TextureDimension::k2dArray:
2185 num_dimensions = is_ms ? 4 : 3;
2186 swizzle = ".z";
2187 break;
2188 case ast::TextureDimension::kCubeArray:
2189 num_dimensions = 3;
2190 swizzle = ".z";
2191 break;
2192 }
2193 break;
2194 case sem::BuiltinType::kTextureNumLevels:
2195 switch (texture_type->dim()) {
2196 default:
2197 TINT_ICE(Writer, diagnostics_)
2198 << "texture dimension does not support mips";
2199 return false;
2200 case ast::TextureDimension::k1d:
2201 num_dimensions = 2;
2202 swizzle = ".y";
2203 break;
2204 case ast::TextureDimension::k2d:
2205 case ast::TextureDimension::kCube:
2206 num_dimensions = 3;
2207 swizzle = ".z";
2208 break;
2209 case ast::TextureDimension::k2dArray:
2210 case ast::TextureDimension::k3d:
2211 case ast::TextureDimension::kCubeArray:
2212 num_dimensions = 4;
2213 swizzle = ".w";
2214 break;
2215 }
2216 break;
2217 case sem::BuiltinType::kTextureNumSamples:
2218 switch (texture_type->dim()) {
2219 default:
2220 TINT_ICE(Writer, diagnostics_)
2221 << "texture dimension does not support multisampling";
2222 return false;
2223 case ast::TextureDimension::k2d:
2224 num_dimensions = 3;
2225 swizzle = ".z";
2226 break;
2227 case ast::TextureDimension::k2dArray:
2228 num_dimensions = 4;
2229 swizzle = ".w";
2230 break;
2231 }
2232 break;
2233 default:
2234 TINT_ICE(Writer, diagnostics_) << "unexpected builtin";
2235 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002236 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002237
2238 auto* level_arg = arg(Usage::kLevel);
2239
2240 if (level_arg) {
2241 // `NumberOfLevels` is a non-optional argument if `MipLevel` was passed.
2242 // Increment the number of dimensions for the temporary vector to
2243 // accommodate this.
2244 num_dimensions++;
2245
2246 // If the swizzle was empty, the expression will evaluate to the whole
2247 // vector. As we've grown the vector by one element, we now need to
2248 // swizzle to keep the result expression equivalent.
2249 if (swizzle.empty()) {
2250 static constexpr const char* swizzles[] = {"", ".x", ".xy", ".xyz"};
2251 swizzle = swizzles[num_dimensions - 1];
2252 }
2253 }
2254
2255 if (num_dimensions > 4) {
2256 TINT_ICE(Writer, diagnostics_) << "Texture query builtin temporary vector has "
2257 << num_dimensions << " dimensions";
2258 return false;
2259 }
2260
2261 // Declare a variable to hold the queried texture info
2262 auto dims = UniqueIdentifier(kTempNamePrefix);
2263 if (num_dimensions == 1) {
2264 line() << "int " << dims << ";";
2265 } else {
2266 line() << "int" << num_dimensions << " " << dims << ";";
2267 }
2268
2269 { // texture.GetDimensions(...)
2270 auto pre = line();
2271 if (!EmitExpression(pre, texture)) {
2272 return false;
2273 }
2274 pre << ".GetDimensions(";
2275
2276 if (level_arg) {
2277 if (!EmitExpression(pre, level_arg)) {
2278 return false;
2279 }
2280 pre << ", ";
2281 } else if (builtin->Type() == sem::BuiltinType::kTextureNumLevels) {
2282 pre << "0, ";
2283 }
2284
2285 if (num_dimensions == 1) {
2286 pre << dims;
2287 } else {
2288 static constexpr char xyzw[] = {'x', 'y', 'z', 'w'};
2289 if (num_dimensions < 0 || num_dimensions > 4) {
2290 TINT_ICE(Writer, diagnostics_)
2291 << "vector dimensions are " << num_dimensions;
2292 return false;
2293 }
2294 for (int i = 0; i < num_dimensions; i++) {
2295 if (i > 0) {
2296 pre << ", ";
2297 }
2298 pre << dims << "." << xyzw[i];
2299 }
2300 }
2301
2302 pre << ");";
2303 }
2304
2305 // The out parameters of the GetDimensions() call is now in temporary
2306 // `dims` variable. This may be packed with other data, so the final
2307 // expression may require a swizzle.
2308 out << dims << swizzle;
2309 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002310 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002311 default:
2312 break;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002313 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002314
Austin Eng86a617f2022-05-19 20:08:19 +00002315 if (!EmitExpression(out, texture)) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002316 return false;
Austin Eng86a617f2022-05-19 20:08:19 +00002317 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002318
2319 // If pack_level_in_coords is true, then the mip level will be appended as the
2320 // last value of the coordinates argument. If the WGSL builtin overload does
2321 // not have a level parameter and pack_level_in_coords is true, then a zero
2322 // mip level will be inserted.
2323 bool pack_level_in_coords = false;
2324
2325 uint32_t hlsl_ret_width = 4u;
2326
2327 switch (builtin->Type()) {
2328 case sem::BuiltinType::kTextureSample:
2329 out << ".Sample(";
2330 break;
2331 case sem::BuiltinType::kTextureSampleBias:
2332 out << ".SampleBias(";
2333 break;
2334 case sem::BuiltinType::kTextureSampleLevel:
2335 out << ".SampleLevel(";
2336 break;
2337 case sem::BuiltinType::kTextureSampleGrad:
2338 out << ".SampleGrad(";
2339 break;
2340 case sem::BuiltinType::kTextureSampleCompare:
2341 out << ".SampleCmp(";
2342 hlsl_ret_width = 1;
2343 break;
2344 case sem::BuiltinType::kTextureSampleCompareLevel:
2345 out << ".SampleCmpLevelZero(";
2346 hlsl_ret_width = 1;
2347 break;
2348 case sem::BuiltinType::kTextureLoad:
2349 out << ".Load(";
2350 // Multisampled textures do not support mip-levels.
2351 if (!texture_type->Is<sem::MultisampledTexture>()) {
2352 pack_level_in_coords = true;
2353 }
2354 break;
2355 case sem::BuiltinType::kTextureGather:
2356 out << ".Gather";
2357 if (builtin->Parameters()[0]->Usage() == sem::ParameterUsage::kComponent) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00002358 switch (call->Arguments()[0]->ConstantValue()->As<AInt>()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002359 case 0:
2360 out << "Red";
2361 break;
2362 case 1:
2363 out << "Green";
2364 break;
2365 case 2:
2366 out << "Blue";
2367 break;
2368 case 3:
2369 out << "Alpha";
2370 break;
2371 }
2372 }
2373 out << "(";
2374 break;
2375 case sem::BuiltinType::kTextureGatherCompare:
2376 out << ".GatherCmp(";
2377 break;
2378 case sem::BuiltinType::kTextureStore:
2379 out << "[";
2380 break;
2381 default:
2382 diagnostics_.add_error(diag::System::Writer,
2383 "Internal compiler error: Unhandled texture builtin '" +
2384 std::string(builtin->str()) + "'");
2385 return false;
2386 }
2387
2388 if (auto* sampler = arg(Usage::kSampler)) {
Austin Eng86a617f2022-05-19 20:08:19 +00002389 if (!EmitExpression(out, sampler)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002390 return false;
Austin Eng86a617f2022-05-19 20:08:19 +00002391 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002392 out << ", ";
2393 }
2394
2395 auto* param_coords = arg(Usage::kCoords);
2396 if (!param_coords) {
2397 TINT_ICE(Writer, diagnostics_) << "missing coords argument";
2398 return false;
2399 }
2400
2401 auto emit_vector_appended_with_i32_zero = [&](const ast::Expression* vector) {
2402 auto* i32 = builder_.create<sem::I32>();
Ben Clayton0ce9ab02022-05-05 20:23:40 +00002403 auto* zero = builder_.Expr(0_i);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002404 auto* stmt = builder_.Sem().Get(vector)->Stmt();
Ben Claytonaa037ac2022-06-29 19:07:30 +00002405 builder_.Sem().Add(
Ben Clayton83bd7382022-07-15 23:46:31 +00002406 zero, builder_.create<sem::Expression>(zero, i32, sem::EvaluationStage::kRuntime, stmt,
2407 /* constant_value */ nullptr,
Ben Claytonaa037ac2022-06-29 19:07:30 +00002408 /* has_side_effects */ false));
dan sinclair41e4d9a2022-05-01 14:40:55 +00002409 auto* packed = AppendVector(&builder_, vector, zero);
2410 return EmitExpression(out, packed->Declaration());
2411 };
2412
2413 auto emit_vector_appended_with_level = [&](const ast::Expression* vector) {
2414 if (auto* level = arg(Usage::kLevel)) {
2415 auto* packed = AppendVector(&builder_, vector, level);
2416 return EmitExpression(out, packed->Declaration());
2417 }
2418 return emit_vector_appended_with_i32_zero(vector);
2419 };
2420
2421 if (auto* array_index = arg(Usage::kArrayIndex)) {
2422 // Array index needs to be appended to the coordinates.
2423 auto* packed = AppendVector(&builder_, param_coords, array_index);
2424 if (pack_level_in_coords) {
2425 // Then mip level needs to be appended to the coordinates.
2426 if (!emit_vector_appended_with_level(packed->Declaration())) {
2427 return false;
2428 }
2429 } else {
2430 if (!EmitExpression(out, packed->Declaration())) {
2431 return false;
2432 }
2433 }
2434 } else if (pack_level_in_coords) {
2435 // Mip level needs to be appended to the coordinates.
2436 if (!emit_vector_appended_with_level(param_coords)) {
2437 return false;
2438 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002439 } else {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002440 if (!EmitExpression(out, param_coords)) {
2441 return false;
2442 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002443 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002444
dan sinclair41e4d9a2022-05-01 14:40:55 +00002445 for (auto usage : {Usage::kDepthRef, Usage::kBias, Usage::kLevel, Usage::kDdx, Usage::kDdy,
2446 Usage::kSampleIndex, Usage::kOffset}) {
2447 if (usage == Usage::kLevel && pack_level_in_coords) {
2448 continue; // mip level already packed in coordinates.
2449 }
2450 if (auto* e = arg(usage)) {
2451 out << ", ";
2452 if (!EmitExpression(out, e)) {
2453 return false;
2454 }
2455 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002456 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002457
dan sinclair41e4d9a2022-05-01 14:40:55 +00002458 if (builtin->Type() == sem::BuiltinType::kTextureStore) {
2459 out << "] = ";
2460 if (!EmitExpression(out, arg(Usage::kValue))) {
2461 return false;
2462 }
2463 } else {
2464 out << ")";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002465
dan sinclair41e4d9a2022-05-01 14:40:55 +00002466 // If the builtin return type does not match the number of elements of the
2467 // HLSL builtin, we need to swizzle the expression to generate the correct
2468 // number of components.
2469 uint32_t wgsl_ret_width = 1;
2470 if (auto* vec = builtin->ReturnType()->As<sem::Vector>()) {
2471 wgsl_ret_width = vec->Width();
2472 }
2473 if (wgsl_ret_width < hlsl_ret_width) {
2474 out << ".";
2475 for (uint32_t i = 0; i < wgsl_ret_width; i++) {
2476 out << "xyz"[i];
2477 }
2478 }
2479 if (wgsl_ret_width > hlsl_ret_width) {
2480 TINT_ICE(Writer, diagnostics_)
2481 << "WGSL return width (" << wgsl_ret_width << ") is wider than HLSL return width ("
2482 << hlsl_ret_width << ") for " << builtin->Type();
2483 return false;
2484 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002485 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002486
dan sinclair41e4d9a2022-05-01 14:40:55 +00002487 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002488}
2489
2490std::string GeneratorImpl::generate_builtin_name(const sem::Builtin* builtin) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002491 switch (builtin->Type()) {
2492 case sem::BuiltinType::kAbs:
2493 case sem::BuiltinType::kAcos:
2494 case sem::BuiltinType::kAll:
2495 case sem::BuiltinType::kAny:
2496 case sem::BuiltinType::kAsin:
2497 case sem::BuiltinType::kAtan:
2498 case sem::BuiltinType::kAtan2:
2499 case sem::BuiltinType::kCeil:
2500 case sem::BuiltinType::kClamp:
2501 case sem::BuiltinType::kCos:
2502 case sem::BuiltinType::kCosh:
2503 case sem::BuiltinType::kCross:
2504 case sem::BuiltinType::kDeterminant:
2505 case sem::BuiltinType::kDistance:
2506 case sem::BuiltinType::kDot:
2507 case sem::BuiltinType::kExp:
2508 case sem::BuiltinType::kExp2:
2509 case sem::BuiltinType::kFloor:
2510 case sem::BuiltinType::kFrexp:
2511 case sem::BuiltinType::kLdexp:
2512 case sem::BuiltinType::kLength:
2513 case sem::BuiltinType::kLog:
2514 case sem::BuiltinType::kLog2:
2515 case sem::BuiltinType::kMax:
2516 case sem::BuiltinType::kMin:
2517 case sem::BuiltinType::kModf:
2518 case sem::BuiltinType::kNormalize:
2519 case sem::BuiltinType::kPow:
2520 case sem::BuiltinType::kReflect:
2521 case sem::BuiltinType::kRefract:
2522 case sem::BuiltinType::kRound:
Ben Clayton751e6682022-09-13 22:57:52 +00002523 case sem::BuiltinType::kSaturate:
dan sinclair41e4d9a2022-05-01 14:40:55 +00002524 case sem::BuiltinType::kSign:
2525 case sem::BuiltinType::kSin:
2526 case sem::BuiltinType::kSinh:
2527 case sem::BuiltinType::kSqrt:
2528 case sem::BuiltinType::kStep:
2529 case sem::BuiltinType::kTan:
2530 case sem::BuiltinType::kTanh:
2531 case sem::BuiltinType::kTranspose:
2532 case sem::BuiltinType::kTrunc:
2533 return builtin->str();
Antonio Maioranoab4c0352022-05-20 01:58:40 +00002534 case sem::BuiltinType::kCountOneBits: // uint
dan sinclair41e4d9a2022-05-01 14:40:55 +00002535 return "countbits";
2536 case sem::BuiltinType::kDpdx:
2537 return "ddx";
2538 case sem::BuiltinType::kDpdxCoarse:
2539 return "ddx_coarse";
2540 case sem::BuiltinType::kDpdxFine:
2541 return "ddx_fine";
2542 case sem::BuiltinType::kDpdy:
2543 return "ddy";
2544 case sem::BuiltinType::kDpdyCoarse:
2545 return "ddy_coarse";
2546 case sem::BuiltinType::kDpdyFine:
2547 return "ddy_fine";
2548 case sem::BuiltinType::kFaceForward:
2549 return "faceforward";
2550 case sem::BuiltinType::kFract:
2551 return "frac";
2552 case sem::BuiltinType::kFma:
2553 return "mad";
2554 case sem::BuiltinType::kFwidth:
2555 case sem::BuiltinType::kFwidthCoarse:
2556 case sem::BuiltinType::kFwidthFine:
2557 return "fwidth";
2558 case sem::BuiltinType::kInverseSqrt:
2559 return "rsqrt";
2560 case sem::BuiltinType::kMix:
2561 return "lerp";
Antonio Maioranoab4c0352022-05-20 01:58:40 +00002562 case sem::BuiltinType::kReverseBits: // uint
dan sinclair41e4d9a2022-05-01 14:40:55 +00002563 return "reversebits";
2564 case sem::BuiltinType::kSmoothstep:
dan sinclair41e4d9a2022-05-01 14:40:55 +00002565 return "smoothstep";
2566 default:
2567 diagnostics_.add_error(diag::System::Writer,
2568 "Unknown builtin method: " + std::string(builtin->str()));
2569 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002570
dan sinclair41e4d9a2022-05-01 14:40:55 +00002571 return "";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002572}
2573
2574bool GeneratorImpl::EmitCase(const ast::SwitchStatement* s, size_t case_idx) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002575 auto* stmt = s->body[case_idx];
dan sinclairf148f082022-10-19 15:55:02 +00002576 auto* sem = builder_.Sem().Get<sem::CaseStatement>(stmt);
2577 for (auto* selector : sem->Selectors()) {
2578 auto out = line();
2579 if (selector->IsDefault()) {
2580 out << "default";
2581 } else {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002582 out << "case ";
dan sinclairf148f082022-10-19 15:55:02 +00002583 if (!EmitConstant(out, selector->Value())) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002584 return false;
2585 }
dan sinclairf148f082022-10-19 15:55:02 +00002586 }
2587 out << ":";
2588 if (selector == sem->Selectors().back()) {
2589 out << " {";
dan sinclair41e4d9a2022-05-01 14:40:55 +00002590 }
2591 }
2592
2593 increment_indent();
2594 TINT_DEFER({
2595 decrement_indent();
2596 line() << "}";
2597 });
2598
2599 // Emit the case statement
2600 if (!EmitStatements(stmt->body->statements)) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002601 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002602 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002603
dan sinclair41e4d9a2022-05-01 14:40:55 +00002604 // Inline all fallthrough case statements. FXC cannot handle fallthroughs.
2605 while (tint::Is<ast::FallthroughStatement>(stmt->body->Last())) {
2606 case_idx++;
2607 stmt = s->body[case_idx];
2608 // Generate each fallthrough case statement in a new block. This is done to
2609 // prevent symbol collision of variables declared in these cases statements.
2610 if (!EmitBlock(stmt->body)) {
2611 return false;
2612 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002613 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002614
dan sinclair41e4d9a2022-05-01 14:40:55 +00002615 if (!tint::IsAnyOf<ast::BreakStatement, ast::FallthroughStatement>(stmt->body->Last())) {
2616 line() << "break;";
2617 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002618
dan sinclair41e4d9a2022-05-01 14:40:55 +00002619 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002620}
2621
2622bool GeneratorImpl::EmitContinue(const ast::ContinueStatement*) {
dan sinclair4b88dbc2022-06-16 15:27:38 +00002623 if (!emit_continuing_ || !emit_continuing_()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002624 return false;
2625 }
2626 line() << "continue;";
2627 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002628}
2629
2630bool GeneratorImpl::EmitDiscard(const ast::DiscardStatement*) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002631 // TODO(dsinclair): Verify this is correct when the discard semantics are
2632 // defined for WGSL (https://github.com/gpuweb/gpuweb/issues/361)
2633 line() << "discard;";
2634 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002635}
2636
dan sinclair41e4d9a2022-05-01 14:40:55 +00002637bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
Ben Claytone9f8b092022-06-01 13:14:39 +00002638 if (auto* sem = builder_.Sem().Get(expr)) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00002639 if (auto* constant = sem->ConstantValue()) {
Ben Claytonc64ca232022-06-29 00:55:36 +00002640 return EmitConstant(out, constant);
Ben Claytone9f8b092022-06-01 13:14:39 +00002641 }
2642 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002643 return Switch(
Ben Claytonb90b6bf2022-08-23 16:23:05 +00002644 expr, //
2645 [&](const ast::IndexAccessorExpression* a) { return EmitIndexAccessor(out, a); },
2646 [&](const ast::BinaryExpression* b) { return EmitBinary(out, b); },
2647 [&](const ast::BitcastExpression* b) { return EmitBitcast(out, b); },
2648 [&](const ast::CallExpression* c) { return EmitCall(out, c); },
2649 [&](const ast::IdentifierExpression* i) { return EmitIdentifier(out, i); },
2650 [&](const ast::LiteralExpression* l) { return EmitLiteral(out, l); },
2651 [&](const ast::MemberAccessorExpression* m) { return EmitMemberAccessor(out, m); },
2652 [&](const ast::UnaryOpExpression* u) { return EmitUnaryOp(out, u); },
2653 [&](Default) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002654 diagnostics_.add_error(diag::System::Writer, "unknown expression type: " +
2655 std::string(expr->TypeInfo().name));
2656 return false;
2657 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002658}
2659
dan sinclair41e4d9a2022-05-01 14:40:55 +00002660bool GeneratorImpl::EmitIdentifier(std::ostream& out, const ast::IdentifierExpression* expr) {
2661 out << builder_.Symbols().NameFor(expr->symbol);
2662 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002663}
2664
2665bool GeneratorImpl::EmitIf(const ast::IfStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002666 {
2667 auto out = line();
2668 out << "if (";
2669 if (!EmitExpression(out, stmt->condition)) {
2670 return false;
2671 }
2672 out << ") {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002673 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002674
dan sinclair41e4d9a2022-05-01 14:40:55 +00002675 if (!EmitStatementsWithIndent(stmt->body->statements)) {
James Price26ebe5e2022-04-29 00:14:53 +00002676 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002677 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002678
dan sinclair41e4d9a2022-05-01 14:40:55 +00002679 if (stmt->else_statement) {
2680 line() << "} else {";
2681 if (auto* block = stmt->else_statement->As<ast::BlockStatement>()) {
2682 if (!EmitStatementsWithIndent(block->statements)) {
2683 return false;
2684 }
2685 } else {
Ben Clayton783b1692022-08-02 17:03:35 +00002686 if (!EmitStatementsWithIndent(utils::Vector{stmt->else_statement})) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002687 return false;
2688 }
2689 }
2690 }
2691 line() << "}";
2692
2693 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002694}
2695
2696bool GeneratorImpl::EmitFunction(const ast::Function* func) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002697 auto* sem = builder_.Sem().Get(func);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002698
Antonio Maiorano08f4b552022-05-31 13:20:28 +00002699 // Emit storage atomic helpers
2700 if (auto* intrinsic =
2701 ast::GetAttribute<transform::DecomposeMemoryAccess::Intrinsic>(func->attributes)) {
dan sinclairff7cf212022-10-03 14:05:23 +00002702 if (intrinsic->address_space == ast::AddressSpace::kStorage && intrinsic->IsAtomic()) {
Antonio Maiorano08f4b552022-05-31 13:20:28 +00002703 if (!EmitStorageAtomicIntrinsic(func, intrinsic)) {
2704 return false;
2705 }
2706 }
2707 return true;
2708 }
2709
dan sinclair41e4d9a2022-05-01 14:40:55 +00002710 if (ast::HasAttribute<ast::InternalAttribute>(func->attributes)) {
2711 // An internal function. Do not emit.
2712 return true;
2713 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002714
dan sinclair41e4d9a2022-05-01 14:40:55 +00002715 {
2716 auto out = line();
2717 auto name = builder_.Symbols().NameFor(func->symbol);
2718 // If the function returns an array, then we need to declare a typedef for
2719 // this.
2720 if (sem->ReturnType()->Is<sem::Array>()) {
2721 auto typedef_name = UniqueIdentifier(name + "_ret");
2722 auto pre = line();
2723 pre << "typedef ";
dan sinclairff7cf212022-10-03 14:05:23 +00002724 if (!EmitTypeAndName(pre, sem->ReturnType(), ast::AddressSpace::kNone,
dan sinclair41e4d9a2022-05-01 14:40:55 +00002725 ast::Access::kReadWrite, typedef_name)) {
2726 return false;
2727 }
2728 pre << ";";
2729 out << typedef_name;
2730 } else {
dan sinclairff7cf212022-10-03 14:05:23 +00002731 if (!EmitType(out, sem->ReturnType(), ast::AddressSpace::kNone, ast::Access::kReadWrite,
dan sinclair41e4d9a2022-05-01 14:40:55 +00002732 "")) {
2733 return false;
2734 }
2735 }
2736
2737 out << " " << name << "(";
2738
2739 bool first = true;
2740
2741 for (auto* v : sem->Parameters()) {
2742 if (!first) {
2743 out << ", ";
2744 }
2745 first = false;
2746
2747 auto const* type = v->Type();
dan sinclairff7cf212022-10-03 14:05:23 +00002748 auto address_space = ast::AddressSpace::kNone;
Ben Claytond2e0db32022-10-12 18:49:15 +00002749 auto access = ast::Access::kUndefined;
dan sinclair41e4d9a2022-05-01 14:40:55 +00002750
2751 if (auto* ptr = type->As<sem::Pointer>()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002752 type = ptr->StoreType();
dan sinclairff7cf212022-10-03 14:05:23 +00002753 switch (ptr->AddressSpace()) {
2754 case ast::AddressSpace::kStorage:
2755 case ast::AddressSpace::kUniform:
Ben Clayton2032d032022-06-15 19:32:37 +00002756 // Not allowed by WGSL, but is used by certain transforms (e.g. DMA) to pass
2757 // storage buffers and uniform buffers down into transform-generated
2758 // functions. In this situation we want to generate the parameter without an
dan sinclairff7cf212022-10-03 14:05:23 +00002759 // 'inout', using the address space and access from the pointer.
2760 address_space = ptr->AddressSpace();
Ben Clayton2032d032022-06-15 19:32:37 +00002761 access = ptr->Access();
2762 break;
2763 default:
2764 // Transform regular WGSL pointer parameters in to `inout` parameters.
2765 out << "inout ";
2766 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00002767 }
2768
dan sinclairff7cf212022-10-03 14:05:23 +00002769 // Note: WGSL only allows for AddressSpace::kNone on parameters, however
dan sinclair41e4d9a2022-05-01 14:40:55 +00002770 // the sanitizer transforms generates load / store functions for storage
2771 // or uniform buffers. These functions have a buffer parameter with
dan sinclairff7cf212022-10-03 14:05:23 +00002772 // AddressSpace::kStorage or AddressSpace::kUniform. This is required to
dan sinclair41e4d9a2022-05-01 14:40:55 +00002773 // correctly translate the parameter to a [RW]ByteAddressBuffer for
2774 // storage buffers and a uint4[N] for uniform buffers.
dan sinclairff7cf212022-10-03 14:05:23 +00002775 if (!EmitTypeAndName(out, type, address_space, access,
dan sinclair41e4d9a2022-05-01 14:40:55 +00002776 builder_.Symbols().NameFor(v->Declaration()->symbol))) {
2777 return false;
2778 }
2779 }
2780 out << ") {";
2781 }
2782
Ben Claytond9222f42022-10-14 13:44:54 +00002783 if (sem->DiscardStatement() && !sem->ReturnType()->Is<sem::Void>()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002784 // BUG(crbug.com/tint/1081): work around non-void functions with discard
2785 // failing compilation sometimes
2786 if (!EmitFunctionBodyWithDiscard(func)) {
2787 return false;
2788 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002789 } else {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002790 if (!EmitStatementsWithIndent(func->body->statements)) {
2791 return false;
2792 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002793 }
2794
dan sinclair41e4d9a2022-05-01 14:40:55 +00002795 line() << "}";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002796
dan sinclair41e4d9a2022-05-01 14:40:55 +00002797 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002798}
2799
2800bool GeneratorImpl::EmitFunctionBodyWithDiscard(const ast::Function* func) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002801 // FXC sometimes fails to compile functions that discard with 'Not all control
2802 // paths return a value'. We work around this by wrapping the function body
2803 // within an "if (true) { <body> } return <default return type obj>;" so that
2804 // there is always an (unused) return statement.
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002805
dan sinclair41e4d9a2022-05-01 14:40:55 +00002806 auto* sem = builder_.Sem().Get(func);
Ben Claytond9222f42022-10-14 13:44:54 +00002807 TINT_ASSERT(Writer, sem->DiscardStatement() && !sem->ReturnType()->Is<sem::Void>());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002808
dan sinclair41e4d9a2022-05-01 14:40:55 +00002809 ScopedIndent si(this);
2810 line() << "if (true) {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002811
dan sinclair41e4d9a2022-05-01 14:40:55 +00002812 if (!EmitStatementsWithIndent(func->body->statements)) {
2813 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002814 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002815
dan sinclair41e4d9a2022-05-01 14:40:55 +00002816 line() << "}";
2817
2818 // Return an unused result that matches the type of the return value
2819 auto name = builder_.Symbols().NameFor(builder_.Symbols().New("unused"));
2820 {
2821 auto out = line();
dan sinclairff7cf212022-10-03 14:05:23 +00002822 if (!EmitTypeAndName(out, sem->ReturnType(), ast::AddressSpace::kNone,
dan sinclair41e4d9a2022-05-01 14:40:55 +00002823 ast::Access::kReadWrite, name)) {
2824 return false;
2825 }
2826 out << ";";
2827 }
2828 line() << "return " << name << ";";
2829
2830 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002831}
2832
2833bool GeneratorImpl::EmitGlobalVariable(const ast::Variable* global) {
Ben Claytondcdf66e2022-06-17 12:48:51 +00002834 return Switch(
2835 global, //
2836 [&](const ast::Var* var) {
2837 auto* sem = builder_.Sem().Get(global);
dan sinclairff7cf212022-10-03 14:05:23 +00002838 switch (sem->AddressSpace()) {
2839 case ast::AddressSpace::kUniform:
Ben Claytondcdf66e2022-06-17 12:48:51 +00002840 return EmitUniformVariable(var, sem);
dan sinclairff7cf212022-10-03 14:05:23 +00002841 case ast::AddressSpace::kStorage:
Ben Claytondcdf66e2022-06-17 12:48:51 +00002842 return EmitStorageVariable(var, sem);
dan sinclairff7cf212022-10-03 14:05:23 +00002843 case ast::AddressSpace::kHandle:
Ben Claytondcdf66e2022-06-17 12:48:51 +00002844 return EmitHandleVariable(var, sem);
dan sinclairff7cf212022-10-03 14:05:23 +00002845 case ast::AddressSpace::kPrivate:
Ben Claytondcdf66e2022-06-17 12:48:51 +00002846 return EmitPrivateVariable(sem);
dan sinclairff7cf212022-10-03 14:05:23 +00002847 case ast::AddressSpace::kWorkgroup:
Ben Claytondcdf66e2022-06-17 12:48:51 +00002848 return EmitWorkgroupVariable(sem);
dan sinclairff7cf212022-10-03 14:05:23 +00002849 case ast::AddressSpace::kPushConstant:
dan sinclair4abf28e2022-08-02 15:55:35 +00002850 diagnostics_.add_error(
2851 diag::System::Writer,
dan sinclairff7cf212022-10-03 14:05:23 +00002852 "unhandled address space " + utils::ToString(sem->AddressSpace()));
dan sinclair4abf28e2022-08-02 15:55:35 +00002853 return false;
dan sinclair8dbd4d02022-07-27 18:54:05 +00002854 default: {
Ben Claytondcdf66e2022-06-17 12:48:51 +00002855 TINT_ICE(Writer, diagnostics_)
dan sinclairff7cf212022-10-03 14:05:23 +00002856 << "unhandled address space " << sem->AddressSpace();
Ben Claytondcdf66e2022-06-17 12:48:51 +00002857 return false;
dan sinclair8dbd4d02022-07-27 18:54:05 +00002858 }
Ben Claytondcdf66e2022-06-17 12:48:51 +00002859 }
2860 },
dan sinclairf6a94042022-09-09 16:16:19 +00002861 [&](const ast::Override*) {
2862 // Override is removed with SubstituteOverride
Ben Clayton490d9882022-09-21 21:05:45 +00002863 diagnostics_.add_error(diag::System::Writer,
Ben Claytonf10a5792022-10-13 13:47:39 +00002864 "override-expressions should have been removed with the "
Ben Clayton490d9882022-09-21 21:05:45 +00002865 "SubstituteOverride transform");
dan sinclairf6a94042022-09-09 16:16:19 +00002866 return false;
2867 },
Ben Clayton19576e92022-06-28 12:44:16 +00002868 [&](const ast::Const*) {
2869 return true; // Constants are embedded at their use
2870 },
Ben Claytondcdf66e2022-06-17 12:48:51 +00002871 [&](Default) {
2872 TINT_ICE(Writer, diagnostics_)
2873 << "unhandled global variable type " << global->TypeInfo().name;
dan sinclair4abf28e2022-08-02 15:55:35 +00002874
Ben Claytondcdf66e2022-06-17 12:48:51 +00002875 return false;
2876 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002877}
2878
Ben Claytondcdf66e2022-06-17 12:48:51 +00002879bool GeneratorImpl::EmitUniformVariable(const ast::Var* var, const sem::Variable* sem) {
dan sinclairacdf6e12022-08-24 15:47:25 +00002880 auto binding_point = sem->As<sem::GlobalVariable>()->BindingPoint();
Ben Claytondcdf66e2022-06-17 12:48:51 +00002881 auto* type = sem->Type()->UnwrapRef();
2882 auto name = builder_.Symbols().NameFor(var->symbol);
dan sinclair41e4d9a2022-05-01 14:40:55 +00002883 line() << "cbuffer cbuffer_" << name << RegisterAndSpace('b', binding_point) << " {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002884
dan sinclair41e4d9a2022-05-01 14:40:55 +00002885 {
2886 ScopedIndent si(this);
2887 auto out = line();
dan sinclairff7cf212022-10-03 14:05:23 +00002888 if (!EmitTypeAndName(out, type, ast::AddressSpace::kUniform, sem->Access(), name)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002889 return false;
2890 }
2891 out << ";";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002892 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002893
dan sinclair41e4d9a2022-05-01 14:40:55 +00002894 line() << "};";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002895
dan sinclair41e4d9a2022-05-01 14:40:55 +00002896 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002897}
2898
Ben Claytondcdf66e2022-06-17 12:48:51 +00002899bool GeneratorImpl::EmitStorageVariable(const ast::Var* var, const sem::Variable* sem) {
2900 auto* type = sem->Type()->UnwrapRef();
dan sinclair41e4d9a2022-05-01 14:40:55 +00002901 auto out = line();
dan sinclairff7cf212022-10-03 14:05:23 +00002902 if (!EmitTypeAndName(out, type, ast::AddressSpace::kStorage, sem->Access(),
Ben Claytondcdf66e2022-06-17 12:48:51 +00002903 builder_.Symbols().NameFor(var->symbol))) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002904 return false;
2905 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002906
dan sinclairacdf6e12022-08-24 15:47:25 +00002907 auto* global_sem = sem->As<sem::GlobalVariable>();
2908 out << RegisterAndSpace(sem->Access() == ast::Access::kRead ? 't' : 'u',
2909 global_sem->BindingPoint())
dan sinclair41e4d9a2022-05-01 14:40:55 +00002910 << ";";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002911
dan sinclair41e4d9a2022-05-01 14:40:55 +00002912 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002913}
2914
Ben Claytondcdf66e2022-06-17 12:48:51 +00002915bool GeneratorImpl::EmitHandleVariable(const ast::Var* var, const sem::Variable* sem) {
2916 auto* unwrapped_type = sem->Type()->UnwrapRef();
dan sinclair41e4d9a2022-05-01 14:40:55 +00002917 auto out = line();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002918
Ben Claytondcdf66e2022-06-17 12:48:51 +00002919 auto name = builder_.Symbols().NameFor(var->symbol);
2920 auto* type = sem->Type()->UnwrapRef();
dan sinclairff7cf212022-10-03 14:05:23 +00002921 if (!EmitTypeAndName(out, type, sem->AddressSpace(), sem->Access(), name)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002922 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002923 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002924
dan sinclair41e4d9a2022-05-01 14:40:55 +00002925 const char* register_space = nullptr;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002926
dan sinclair41e4d9a2022-05-01 14:40:55 +00002927 if (unwrapped_type->Is<sem::Texture>()) {
2928 register_space = "t";
2929 if (unwrapped_type->Is<sem::StorageTexture>()) {
2930 register_space = "u";
2931 }
2932 } else if (unwrapped_type->Is<sem::Sampler>()) {
2933 register_space = "s";
2934 }
2935
2936 if (register_space) {
dan sinclairacdf6e12022-08-24 15:47:25 +00002937 auto bp = sem->As<sem::GlobalVariable>()->BindingPoint();
2938 out << " : register(" << register_space << bp.binding << ", space" << bp.group << ")";
dan sinclair41e4d9a2022-05-01 14:40:55 +00002939 }
2940
2941 out << ";";
2942 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002943}
2944
2945bool GeneratorImpl::EmitPrivateVariable(const sem::Variable* var) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002946 auto* decl = var->Declaration();
2947 auto out = line();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002948
dan sinclair41e4d9a2022-05-01 14:40:55 +00002949 out << "static ";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002950
dan sinclair41e4d9a2022-05-01 14:40:55 +00002951 auto name = builder_.Symbols().NameFor(decl->symbol);
2952 auto* type = var->Type()->UnwrapRef();
dan sinclairff7cf212022-10-03 14:05:23 +00002953 if (!EmitTypeAndName(out, type, var->AddressSpace(), var->Access(), name)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002954 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002955 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002956
dan sinclair41e4d9a2022-05-01 14:40:55 +00002957 out << " = ";
dan sinclair6e77b472022-10-20 13:38:28 +00002958 if (auto* initializer = decl->initializer) {
2959 if (!EmitExpression(out, initializer)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002960 return false;
2961 }
2962 } else {
2963 if (!EmitZeroValue(out, var->Type()->UnwrapRef())) {
2964 return false;
2965 }
2966 }
2967
2968 out << ";";
2969 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002970}
2971
2972bool GeneratorImpl::EmitWorkgroupVariable(const sem::Variable* var) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002973 auto* decl = var->Declaration();
2974 auto out = line();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002975
dan sinclair41e4d9a2022-05-01 14:40:55 +00002976 out << "groupshared ";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002977
dan sinclair41e4d9a2022-05-01 14:40:55 +00002978 auto name = builder_.Symbols().NameFor(decl->symbol);
2979 auto* type = var->Type()->UnwrapRef();
dan sinclairff7cf212022-10-03 14:05:23 +00002980 if (!EmitTypeAndName(out, type, var->AddressSpace(), var->Access(), name)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002981 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002982 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002983
dan sinclair6e77b472022-10-20 13:38:28 +00002984 if (auto* initializer = decl->initializer) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002985 out << " = ";
dan sinclair6e77b472022-10-20 13:38:28 +00002986 if (!EmitExpression(out, initializer)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002987 return false;
2988 }
2989 }
2990
2991 out << ";";
2992 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00002993}
2994
Ben Claytonf3302292022-07-27 18:48:06 +00002995std::string GeneratorImpl::builtin_to_attribute(ast::BuiltinValue builtin) const {
dan sinclair41e4d9a2022-05-01 14:40:55 +00002996 switch (builtin) {
Ben Claytonf3302292022-07-27 18:48:06 +00002997 case ast::BuiltinValue::kPosition:
dan sinclair41e4d9a2022-05-01 14:40:55 +00002998 return "SV_Position";
Ben Claytonf3302292022-07-27 18:48:06 +00002999 case ast::BuiltinValue::kVertexIndex:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003000 return "SV_VertexID";
Ben Claytonf3302292022-07-27 18:48:06 +00003001 case ast::BuiltinValue::kInstanceIndex:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003002 return "SV_InstanceID";
Ben Claytonf3302292022-07-27 18:48:06 +00003003 case ast::BuiltinValue::kFrontFacing:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003004 return "SV_IsFrontFace";
Ben Claytonf3302292022-07-27 18:48:06 +00003005 case ast::BuiltinValue::kFragDepth:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003006 return "SV_Depth";
Ben Claytonf3302292022-07-27 18:48:06 +00003007 case ast::BuiltinValue::kLocalInvocationId:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003008 return "SV_GroupThreadID";
Ben Claytonf3302292022-07-27 18:48:06 +00003009 case ast::BuiltinValue::kLocalInvocationIndex:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003010 return "SV_GroupIndex";
Ben Claytonf3302292022-07-27 18:48:06 +00003011 case ast::BuiltinValue::kGlobalInvocationId:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003012 return "SV_DispatchThreadID";
Ben Claytonf3302292022-07-27 18:48:06 +00003013 case ast::BuiltinValue::kWorkgroupId:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003014 return "SV_GroupID";
Ben Claytonf3302292022-07-27 18:48:06 +00003015 case ast::BuiltinValue::kSampleIndex:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003016 return "SV_SampleIndex";
Ben Claytonf3302292022-07-27 18:48:06 +00003017 case ast::BuiltinValue::kSampleMask:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003018 return "SV_Coverage";
3019 default:
3020 break;
3021 }
3022 return "";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003023}
3024
dan sinclair41e4d9a2022-05-01 14:40:55 +00003025std::string GeneratorImpl::interpolation_to_modifiers(ast::InterpolationType type,
3026 ast::InterpolationSampling sampling) const {
3027 std::string modifiers;
3028 switch (type) {
3029 case ast::InterpolationType::kPerspective:
3030 modifiers += "linear ";
3031 break;
3032 case ast::InterpolationType::kLinear:
3033 modifiers += "noperspective ";
3034 break;
3035 case ast::InterpolationType::kFlat:
3036 modifiers += "nointerpolation ";
3037 break;
Ben Claytond2e0db32022-10-12 18:49:15 +00003038 case ast::InterpolationType::kUndefined:
Ben Claytonf9ed9d32022-10-11 19:49:17 +00003039 break;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003040 }
3041 switch (sampling) {
3042 case ast::InterpolationSampling::kCentroid:
3043 modifiers += "centroid ";
3044 break;
3045 case ast::InterpolationSampling::kSample:
3046 modifiers += "sample ";
3047 break;
3048 case ast::InterpolationSampling::kCenter:
Ben Claytond2e0db32022-10-12 18:49:15 +00003049 case ast::InterpolationSampling::kUndefined:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003050 break;
3051 }
3052 return modifiers;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003053}
3054
3055bool GeneratorImpl::EmitEntryPointFunction(const ast::Function* func) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003056 auto* func_sem = builder_.Sem().Get(func);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003057
dan sinclair41e4d9a2022-05-01 14:40:55 +00003058 {
3059 auto out = line();
3060 if (func->PipelineStage() == ast::PipelineStage::kCompute) {
3061 // Emit the workgroup_size attribute.
3062 auto wgsize = func_sem->WorkgroupSize();
3063 out << "[numthreads(";
dan sinclair3a2a2792022-06-29 14:38:15 +00003064 for (size_t i = 0; i < 3; i++) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003065 if (i > 0) {
3066 out << ", ";
3067 }
Ben Clayton490d9882022-09-21 21:05:45 +00003068 if (!wgsize[i].has_value()) {
3069 diagnostics_.add_error(
3070 diag::System::Writer,
Ben Claytonf10a5792022-10-13 13:47:39 +00003071 "override-expressions should have been removed with the SubstituteOverride "
Ben Clayton490d9882022-09-21 21:05:45 +00003072 "transform");
3073 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003074 }
Ben Clayton490d9882022-09-21 21:05:45 +00003075 out << std::to_string(wgsize[i].value());
dan sinclair41e4d9a2022-05-01 14:40:55 +00003076 }
3077 out << ")]" << std::endl;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003078 }
3079
dan sinclair41e4d9a2022-05-01 14:40:55 +00003080 out << func->return_type->FriendlyName(builder_.Symbols());
3081
3082 out << " " << builder_.Symbols().NameFor(func->symbol) << "(";
3083
3084 bool first = true;
3085
3086 // Emit entry point parameters.
3087 for (auto* var : func->params) {
3088 auto* sem = builder_.Sem().Get(var);
3089 auto* type = sem->Type();
3090 if (!type->Is<sem::Struct>()) {
3091 // ICE likely indicates that the CanonicalizeEntryPointIO transform was
3092 // not run, or a builtin parameter was added after it was run.
3093 TINT_ICE(Writer, diagnostics_) << "Unsupported non-struct entry point parameter";
3094 }
3095
3096 if (!first) {
3097 out << ", ";
3098 }
3099 first = false;
3100
dan sinclairff7cf212022-10-03 14:05:23 +00003101 if (!EmitTypeAndName(out, type, sem->AddressSpace(), sem->Access(),
dan sinclair41e4d9a2022-05-01 14:40:55 +00003102 builder_.Symbols().NameFor(var->symbol))) {
3103 return false;
3104 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003105 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00003106
3107 out << ") {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003108 }
3109
dan sinclair41e4d9a2022-05-01 14:40:55 +00003110 {
3111 ScopedIndent si(this);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003112
dan sinclair41e4d9a2022-05-01 14:40:55 +00003113 if (!EmitStatements(func->body->statements)) {
3114 return false;
3115 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003116
dan sinclair41e4d9a2022-05-01 14:40:55 +00003117 if (!Is<ast::ReturnStatement>(func->body->Last())) {
Ben Clayton4a92a3c2022-07-18 20:50:02 +00003118 ast::ReturnStatement ret(ProgramID(), ast::NodeID{}, Source{});
dan sinclair41e4d9a2022-05-01 14:40:55 +00003119 if (!EmitStatement(&ret)) {
3120 return false;
3121 }
3122 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003123 }
3124
dan sinclair41e4d9a2022-05-01 14:40:55 +00003125 line() << "}";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003126
dan sinclair41e4d9a2022-05-01 14:40:55 +00003127 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003128}
3129
Ben Claytonaa037ac2022-06-29 19:07:30 +00003130bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constant) {
Ben Clayton50414802022-06-24 08:06:19 +00003131 return Switch(
Ben Claytonaa037ac2022-06-29 19:07:30 +00003132 constant->Type(), //
Ben Clayton50414802022-06-24 08:06:19 +00003133 [&](const sem::Bool*) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00003134 out << (constant->As<AInt>() ? "true" : "false");
Ben Claytone9f8b092022-06-01 13:14:39 +00003135 return true;
Ben Clayton50414802022-06-24 08:06:19 +00003136 },
3137 [&](const sem::F32*) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00003138 PrintF32(out, constant->As<float>());
Ben Clayton50414802022-06-24 08:06:19 +00003139 return true;
3140 },
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003141 [&](const sem::F16*) {
3142 // emit a f16 scalar with explicit float16_t type declaration.
3143 out << "float16_t(";
Antonio Maiorano679cf4f2022-09-03 21:43:01 +00003144 PrintF16(out, constant->As<float>());
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003145 out << ")";
Antonio Maiorano679cf4f2022-09-03 21:43:01 +00003146 return true;
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003147 },
Ben Clayton50414802022-06-24 08:06:19 +00003148 [&](const sem::I32*) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00003149 out << constant->As<AInt>();
Ben Clayton50414802022-06-24 08:06:19 +00003150 return true;
3151 },
3152 [&](const sem::U32*) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00003153 out << constant->As<AInt>() << "u";
Ben Clayton50414802022-06-24 08:06:19 +00003154 return true;
3155 },
3156 [&](const sem::Vector* v) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00003157 if (constant->AllEqual()) {
Ben Clayton50414802022-06-24 08:06:19 +00003158 {
3159 ScopedParen sp(out);
Ben Claytonaa037ac2022-06-29 19:07:30 +00003160 if (!EmitConstant(out, constant->Index(0))) {
Ben Clayton50414802022-06-24 08:06:19 +00003161 return false;
3162 }
3163 }
3164 out << ".";
Ben Claytonaa037ac2022-06-29 19:07:30 +00003165 for (size_t i = 0; i < v->Width(); i++) {
Ben Clayton50414802022-06-24 08:06:19 +00003166 out << "x";
3167 }
3168 return true;
3169 }
Ben Claytone9f8b092022-06-01 13:14:39 +00003170
Ben Claytond2e0db32022-10-12 18:49:15 +00003171 if (!EmitType(out, v, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
Ben Clayton50414802022-06-24 08:06:19 +00003172 return false;
3173 }
Ben Claytone9f8b092022-06-01 13:14:39 +00003174
Ben Clayton50414802022-06-24 08:06:19 +00003175 ScopedParen sp(out);
Ben Claytone9f8b092022-06-01 13:14:39 +00003176
Ben Claytonaa037ac2022-06-29 19:07:30 +00003177 for (size_t i = 0; i < v->Width(); i++) {
3178 if (i > 0) {
Ben Claytone9f8b092022-06-01 13:14:39 +00003179 out << ", ";
3180 }
Ben Claytonaa037ac2022-06-29 19:07:30 +00003181 if (!EmitConstant(out, constant->Index(i))) {
Ben Claytone9f8b092022-06-01 13:14:39 +00003182 return false;
3183 }
3184 }
3185 return true;
Ben Clayton50414802022-06-24 08:06:19 +00003186 },
3187 [&](const sem::Matrix* m) {
Ben Claytond2e0db32022-10-12 18:49:15 +00003188 if (!EmitType(out, m, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
Ben Claytone9f8b092022-06-01 13:14:39 +00003189 return false;
3190 }
Ben Clayton50414802022-06-24 08:06:19 +00003191
3192 ScopedParen sp(out);
3193
Ben Claytonaa037ac2022-06-29 19:07:30 +00003194 for (size_t i = 0; i < m->columns(); i++) {
3195 if (i > 0) {
Ben Clayton50414802022-06-24 08:06:19 +00003196 out << ", ";
3197 }
Ben Claytonaa037ac2022-06-29 19:07:30 +00003198 if (!EmitConstant(out, constant->Index(i))) {
Ben Clayton50414802022-06-24 08:06:19 +00003199 return false;
3200 }
3201 }
3202 return true;
3203 },
Ben Clayton19576e92022-06-28 12:44:16 +00003204 [&](const sem::Array* a) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00003205 if (constant->AllZero()) {
Ben Clayton19576e92022-06-28 12:44:16 +00003206 out << "(";
Ben Claytond2e0db32022-10-12 18:49:15 +00003207 if (!EmitType(out, a, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
Ben Clayton19576e92022-06-28 12:44:16 +00003208 return false;
3209 }
3210 out << ")0";
3211 return true;
3212 }
3213
3214 out << "{";
3215 TINT_DEFER(out << "}");
3216
dan sinclair78f80672022-09-22 22:28:21 +00003217 auto count = a->ConstantCount();
3218 if (!count) {
3219 diagnostics_.add_error(diag::System::Writer, sem::Array::kErrExpectedConstantCount);
3220 return false;
3221 }
3222
3223 for (size_t i = 0; i < count; i++) {
Ben Claytonaa037ac2022-06-29 19:07:30 +00003224 if (i > 0) {
Ben Clayton19576e92022-06-28 12:44:16 +00003225 out << ", ";
3226 }
Ben Claytonaa037ac2022-06-29 19:07:30 +00003227 if (!EmitConstant(out, constant->Index(i))) {
Ben Clayton19576e92022-06-28 12:44:16 +00003228 return false;
3229 }
3230 }
3231
3232 return true;
3233 },
Ben Clayton6c098ba2022-07-14 20:46:39 +00003234 [&](const sem::Struct* s) {
3235 if (constant->AllZero()) {
3236 out << "(";
Ben Claytond2e0db32022-10-12 18:49:15 +00003237 if (!EmitType(out, s, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
Ben Clayton6c098ba2022-07-14 20:46:39 +00003238 return false;
3239 }
3240 out << ")0";
3241 return true;
3242 }
3243
3244 out << "{";
3245 TINT_DEFER(out << "}");
3246
3247 for (size_t i = 0; i < s->Members().size(); i++) {
3248 if (i > 0) {
3249 out << ", ";
3250 }
3251 if (!EmitConstant(out, constant->Index(i))) {
3252 return false;
3253 }
3254 }
3255
3256 return true;
3257 },
Ben Claytone9f8b092022-06-01 13:14:39 +00003258 [&](Default) {
3259 diagnostics_.add_error(
3260 diag::System::Writer,
Ben Claytonaa037ac2022-06-29 19:07:30 +00003261 "unhandled constant type: " + builder_.FriendlyName(constant->Type()));
Ben Claytone9f8b092022-06-01 13:14:39 +00003262 return false;
3263 });
3264}
3265
dan sinclair41e4d9a2022-05-01 14:40:55 +00003266bool GeneratorImpl::EmitLiteral(std::ostream& out, const ast::LiteralExpression* lit) {
3267 return Switch(
3268 lit,
3269 [&](const ast::BoolLiteralExpression* l) {
3270 out << (l->value ? "true" : "false");
3271 return true;
3272 },
Ben Clayton3ad927c2022-05-25 23:12:14 +00003273 [&](const ast::FloatLiteralExpression* l) {
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003274 if (l->suffix == ast::FloatLiteralExpression::Suffix::kH) {
3275 // Emit f16 literal with explicit float16_t type declaration.
3276 out << "float16_t(";
Antonio Maiorano679cf4f2022-09-03 21:43:01 +00003277 PrintF16(out, static_cast<float>(l->value));
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003278 out << ")";
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003279 }
Ben Claytone9f8b092022-06-01 13:14:39 +00003280 PrintF32(out, static_cast<float>(l->value));
dan sinclair41e4d9a2022-05-01 14:40:55 +00003281 return true;
3282 },
Ben Clayton8822e292022-05-04 22:18:49 +00003283 [&](const ast::IntLiteralExpression* i) {
3284 out << i->value;
3285 switch (i->suffix) {
3286 case ast::IntLiteralExpression::Suffix::kNone:
3287 case ast::IntLiteralExpression::Suffix::kI:
3288 return true;
3289 case ast::IntLiteralExpression::Suffix::kU:
3290 out << "u";
3291 return true;
3292 }
3293 diagnostics_.add_error(diag::System::Writer, "unknown integer literal suffix type");
3294 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003295 },
3296 [&](Default) {
3297 diagnostics_.add_error(diag::System::Writer, "unknown literal type");
3298 return false;
3299 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003300}
3301
dan sinclair41e4d9a2022-05-01 14:40:55 +00003302bool GeneratorImpl::EmitValue(std::ostream& out, const sem::Type* type, int value) {
3303 return Switch(
3304 type,
3305 [&](const sem::Bool*) {
3306 out << (value == 0 ? "false" : "true");
3307 return true;
3308 },
3309 [&](const sem::F32*) {
3310 out << value << ".0f";
3311 return true;
3312 },
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003313 [&](const sem::F16*) {
3314 out << "float16_t(" << value << ".0h)";
3315 return true;
3316 },
dan sinclair41e4d9a2022-05-01 14:40:55 +00003317 [&](const sem::I32*) {
3318 out << value;
3319 return true;
3320 },
3321 [&](const sem::U32*) {
3322 out << value << "u";
3323 return true;
3324 },
3325 [&](const sem::Vector* vec) {
dan sinclairff7cf212022-10-03 14:05:23 +00003326 if (!EmitType(out, type, ast::AddressSpace::kNone, ast::Access::kReadWrite, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003327 return false;
3328 }
3329 ScopedParen sp(out);
3330 for (uint32_t i = 0; i < vec->Width(); i++) {
3331 if (i != 0) {
3332 out << ", ";
3333 }
3334 if (!EmitValue(out, vec->type(), value)) {
3335 return false;
3336 }
3337 }
3338 return true;
3339 },
3340 [&](const sem::Matrix* mat) {
dan sinclairff7cf212022-10-03 14:05:23 +00003341 if (!EmitType(out, type, ast::AddressSpace::kNone, ast::Access::kReadWrite, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003342 return false;
3343 }
3344 ScopedParen sp(out);
3345 for (uint32_t i = 0; i < (mat->rows() * mat->columns()); i++) {
3346 if (i != 0) {
3347 out << ", ";
3348 }
3349 if (!EmitValue(out, mat->type(), value)) {
3350 return false;
3351 }
3352 }
3353 return true;
3354 },
3355 [&](const sem::Struct*) {
3356 out << "(";
3357 TINT_DEFER(out << ")" << value);
Ben Claytond2e0db32022-10-12 18:49:15 +00003358 return EmitType(out, type, ast::AddressSpace::kNone, ast::Access::kUndefined, "");
dan sinclair41e4d9a2022-05-01 14:40:55 +00003359 },
3360 [&](const sem::Array*) {
3361 out << "(";
3362 TINT_DEFER(out << ")" << value);
Ben Claytond2e0db32022-10-12 18:49:15 +00003363 return EmitType(out, type, ast::AddressSpace::kNone, ast::Access::kUndefined, "");
dan sinclair41e4d9a2022-05-01 14:40:55 +00003364 },
3365 [&](Default) {
3366 diagnostics_.add_error(
3367 diag::System::Writer,
3368 "Invalid type for value emission: " + type->FriendlyName(builder_.Symbols()));
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003369 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003370 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003371}
3372
3373bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003374 return EmitValue(out, type, 0);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003375}
3376
3377bool GeneratorImpl::EmitLoop(const ast::LoopStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003378 auto emit_continuing = [this, stmt]() {
3379 if (stmt->continuing && !stmt->continuing->Empty()) {
3380 if (!EmitBlock(stmt->continuing)) {
3381 return false;
3382 }
3383 }
3384 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003385 };
3386
3387 TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
Antonio Maiorano06844a52022-09-29 16:53:58 +00003388 line() << "while (true) {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003389 {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003390 ScopedIndent si(this);
3391 if (!EmitStatements(stmt->body->statements)) {
3392 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003393 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00003394 if (!emit_continuing_()) {
3395 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003396 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003397 }
3398 line() << "}";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003399
dan sinclair41e4d9a2022-05-01 14:40:55 +00003400 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003401}
3402
dan sinclair41e4d9a2022-05-01 14:40:55 +00003403bool GeneratorImpl::EmitForLoop(const ast::ForLoopStatement* stmt) {
3404 // Nest a for loop with a new block. In HLSL the initializer scope is not
3405 // nested by the for-loop, so we may get variable redefinitions.
3406 line() << "{";
3407 increment_indent();
3408 TINT_DEFER({
3409 decrement_indent();
3410 line() << "}";
3411 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003412
dan sinclair41e4d9a2022-05-01 14:40:55 +00003413 TextBuffer init_buf;
3414 if (auto* init = stmt->initializer) {
3415 TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf);
3416 if (!EmitStatement(init)) {
3417 return false;
3418 }
3419 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003420
dan sinclair41e4d9a2022-05-01 14:40:55 +00003421 TextBuffer cond_pre;
3422 std::stringstream cond_buf;
3423 if (auto* cond = stmt->condition) {
3424 TINT_SCOPED_ASSIGNMENT(current_buffer_, &cond_pre);
3425 if (!EmitExpression(cond_buf, cond)) {
3426 return false;
3427 }
3428 }
3429
3430 TextBuffer cont_buf;
3431 if (auto* cont = stmt->continuing) {
3432 TINT_SCOPED_ASSIGNMENT(current_buffer_, &cont_buf);
3433 if (!EmitStatement(cont)) {
3434 return false;
3435 }
3436 }
3437
3438 // If the for-loop has a multi-statement conditional and / or continuing, then
3439 // we cannot emit this as a regular for-loop in HLSL. Instead we need to
3440 // generate a `while(true)` loop.
3441 bool emit_as_loop = cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1;
3442
3443 // If the for-loop has multi-statement initializer, or is going to be emitted
3444 // as a `while(true)` loop, then declare the initializer statement(s) before
3445 // the loop.
3446 if (init_buf.lines.size() > 1 || (stmt->initializer && emit_as_loop)) {
3447 current_buffer_->Append(init_buf);
3448 init_buf.lines.clear(); // Don't emit the initializer again in the 'for'
3449 }
3450
3451 if (emit_as_loop) {
3452 auto emit_continuing = [&]() {
3453 current_buffer_->Append(cont_buf);
3454 return true;
3455 };
3456
3457 TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
Antonio Maiorano06844a52022-09-29 16:53:58 +00003458 line() << "while (true) {";
dan sinclair41e4d9a2022-05-01 14:40:55 +00003459 increment_indent();
3460 TINT_DEFER({
3461 decrement_indent();
3462 line() << "}";
3463 });
3464
3465 if (stmt->condition) {
3466 current_buffer_->Append(cond_pre);
3467 line() << "if (!(" << cond_buf.str() << ")) { break; }";
3468 }
3469
3470 if (!EmitStatements(stmt->body->statements)) {
3471 return false;
3472 }
3473
3474 if (!emit_continuing_()) {
3475 return false;
3476 }
3477 } else {
3478 // For-loop can be generated.
3479 {
3480 auto out = line();
Antonio Maiorano06844a52022-09-29 16:53:58 +00003481 out << "for";
dan sinclair41e4d9a2022-05-01 14:40:55 +00003482 {
3483 ScopedParen sp(out);
3484
3485 if (!init_buf.lines.empty()) {
3486 out << init_buf.lines[0].content << " ";
3487 } else {
3488 out << "; ";
3489 }
3490
3491 out << cond_buf.str() << "; ";
3492
3493 if (!cont_buf.lines.empty()) {
3494 out << TrimSuffix(cont_buf.lines[0].content, ";");
3495 }
3496 }
3497 out << " {";
3498 }
3499 {
3500 auto emit_continuing = [] { return true; };
3501 TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
3502 if (!EmitStatementsWithIndent(stmt->body->statements)) {
3503 return false;
3504 }
3505 }
3506 line() << "}";
3507 }
3508
3509 return true;
3510}
3511
dan sinclair49d1a2d2022-06-16 12:01:27 +00003512bool GeneratorImpl::EmitWhile(const ast::WhileStatement* stmt) {
3513 TextBuffer cond_pre;
3514 std::stringstream cond_buf;
3515 {
3516 auto* cond = stmt->condition;
3517 TINT_SCOPED_ASSIGNMENT(current_buffer_, &cond_pre);
3518 if (!EmitExpression(cond_buf, cond)) {
3519 return false;
3520 }
3521 }
3522
dan sinclair4b88dbc2022-06-16 15:27:38 +00003523 auto emit_continuing = [&]() { return true; };
3524 TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
3525
dan sinclair49d1a2d2022-06-16 12:01:27 +00003526 // If the while has a multi-statement conditional, then we cannot emit this
3527 // as a regular while in HLSL. Instead we need to generate a `while(true)` loop.
3528 bool emit_as_loop = cond_pre.lines.size() > 0;
3529 if (emit_as_loop) {
Antonio Maiorano06844a52022-09-29 16:53:58 +00003530 line() << "while (true) {";
dan sinclair49d1a2d2022-06-16 12:01:27 +00003531 increment_indent();
3532 TINT_DEFER({
3533 decrement_indent();
3534 line() << "}";
3535 });
3536
3537 current_buffer_->Append(cond_pre);
3538 line() << "if (!(" << cond_buf.str() << ")) { break; }";
3539 if (!EmitStatements(stmt->body->statements)) {
3540 return false;
3541 }
3542 } else {
3543 // While can be generated.
3544 {
3545 auto out = line();
Antonio Maiorano06844a52022-09-29 16:53:58 +00003546 out << "while";
dan sinclair49d1a2d2022-06-16 12:01:27 +00003547 {
3548 ScopedParen sp(out);
3549 out << cond_buf.str();
3550 }
3551 out << " {";
3552 }
3553 if (!EmitStatementsWithIndent(stmt->body->statements)) {
3554 return false;
3555 }
3556 line() << "}";
3557 }
3558
3559 return true;
3560}
3561
dan sinclair41e4d9a2022-05-01 14:40:55 +00003562bool GeneratorImpl::EmitMemberAccessor(std::ostream& out,
3563 const ast::MemberAccessorExpression* expr) {
3564 if (!EmitExpression(out, expr->structure)) {
3565 return false;
3566 }
3567 out << ".";
3568
3569 // Swizzles output the name directly
3570 if (builder_.Sem().Get(expr)->Is<sem::Swizzle>()) {
3571 out << builder_.Symbols().NameFor(expr->member->symbol);
3572 } else if (!EmitExpression(out, expr->member)) {
3573 return false;
3574 }
3575
3576 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003577}
3578
3579bool GeneratorImpl::EmitReturn(const ast::ReturnStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003580 if (stmt->value) {
3581 auto out = line();
3582 out << "return ";
3583 if (!EmitExpression(out, stmt->value)) {
3584 return false;
3585 }
3586 out << ";";
3587 } else {
3588 line() << "return;";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003589 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00003590 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003591}
3592
3593bool GeneratorImpl::EmitStatement(const ast::Statement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003594 return Switch(
3595 stmt,
3596 [&](const ast::AssignmentStatement* a) { //
3597 return EmitAssign(a);
3598 },
3599 [&](const ast::BlockStatement* b) { //
3600 return EmitBlock(b);
3601 },
3602 [&](const ast::BreakStatement* b) { //
3603 return EmitBreak(b);
3604 },
dan sinclairb8b0c212022-10-20 22:45:50 +00003605 [&](const ast::BreakIfStatement* b) { //
3606 return EmitBreakIf(b);
3607 },
dan sinclair41e4d9a2022-05-01 14:40:55 +00003608 [&](const ast::CallStatement* c) { //
3609 auto out = line();
3610 if (!EmitCall(out, c->expr)) {
3611 return false;
3612 }
3613 out << ";";
3614 return true;
3615 },
3616 [&](const ast::ContinueStatement* c) { //
3617 return EmitContinue(c);
3618 },
3619 [&](const ast::DiscardStatement* d) { //
3620 return EmitDiscard(d);
3621 },
3622 [&](const ast::FallthroughStatement*) { //
3623 line() << "/* fallthrough */";
3624 return true;
3625 },
3626 [&](const ast::IfStatement* i) { //
3627 return EmitIf(i);
3628 },
3629 [&](const ast::LoopStatement* l) { //
3630 return EmitLoop(l);
3631 },
3632 [&](const ast::ForLoopStatement* l) { //
3633 return EmitForLoop(l);
3634 },
dan sinclair49d1a2d2022-06-16 12:01:27 +00003635 [&](const ast::WhileStatement* l) { //
3636 return EmitWhile(l);
3637 },
dan sinclair41e4d9a2022-05-01 14:40:55 +00003638 [&](const ast::ReturnStatement* r) { //
3639 return EmitReturn(r);
3640 },
3641 [&](const ast::SwitchStatement* s) { //
3642 return EmitSwitch(s);
3643 },
3644 [&](const ast::VariableDeclStatement* v) { //
Ben Claytondcdf66e2022-06-17 12:48:51 +00003645 return Switch(
3646 v->variable, //
3647 [&](const ast::Var* var) { return EmitVar(var); },
3648 [&](const ast::Let* let) { return EmitLet(let); },
Ben Clayton19576e92022-06-28 12:44:16 +00003649 [&](const ast::Const*) {
3650 return true; // Constants are embedded at their use
3651 },
Ben Claytondcdf66e2022-06-17 12:48:51 +00003652 [&](Default) { //
3653 TINT_ICE(Writer, diagnostics_)
3654 << "unknown variable type: " << v->variable->TypeInfo().name;
3655 return false;
3656 });
dan sinclair41e4d9a2022-05-01 14:40:55 +00003657 },
Ben Claytonb4744ac2022-08-03 07:01:08 +00003658 [&](const ast::StaticAssert*) {
3659 return true; // Not emitted
3660 },
dan sinclair41e4d9a2022-05-01 14:40:55 +00003661 [&](Default) { //
3662 diagnostics_.add_error(diag::System::Writer,
3663 "unknown statement type: " + std::string(stmt->TypeInfo().name));
3664 return false;
3665 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003666}
3667
3668bool GeneratorImpl::EmitDefaultOnlySwitch(const ast::SwitchStatement* stmt) {
dan sinclairf148f082022-10-19 15:55:02 +00003669 TINT_ASSERT(Writer, stmt->body.Length() == 1 && stmt->body[0]->ContainsDefault());
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003670
dan sinclair41e4d9a2022-05-01 14:40:55 +00003671 // FXC fails to compile a switch with just a default case, ignoring the
3672 // default case body. We work around this here by emitting the default case
3673 // without the switch.
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003674
dan sinclair41e4d9a2022-05-01 14:40:55 +00003675 // Emit the switch condition as-is in case it has side-effects (e.g.
3676 // function call). Note that's it's fine not to assign the result of the
3677 // expression.
3678 {
3679 auto out = line();
3680 if (!EmitExpression(out, stmt->condition)) {
3681 return false;
3682 }
3683 out << ";";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003684 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003685
dan sinclair41e4d9a2022-05-01 14:40:55 +00003686 // Emit "do { <default case body> } while(false);". We use a 'do' loop so
3687 // that break statements work as expected, and make it 'while (false)' in
3688 // case there isn't a break statement.
3689 line() << "do {";
3690 {
3691 ScopedIndent si(this);
3692 if (!EmitStatements(stmt->body[0]->body->statements)) {
3693 return false;
3694 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003695 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00003696 line() << "} while (false);";
3697 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003698}
3699
3700bool GeneratorImpl::EmitSwitch(const ast::SwitchStatement* stmt) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003701 // BUG(crbug.com/tint/1188): work around default-only switches
dan sinclairf148f082022-10-19 15:55:02 +00003702 if (stmt->body.Length() == 1 && stmt->body[0]->selectors.Length() == 1 &&
3703 stmt->body[0]->ContainsDefault()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003704 return EmitDefaultOnlySwitch(stmt);
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003705 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003706
dan sinclair41e4d9a2022-05-01 14:40:55 +00003707 { // switch(expr) {
3708 auto out = line();
3709 out << "switch(";
3710 if (!EmitExpression(out, stmt->condition)) {
3711 return false;
3712 }
3713 out << ") {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003714 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003715
dan sinclair41e4d9a2022-05-01 14:40:55 +00003716 {
3717 ScopedIndent si(this);
Ben Clayton783b1692022-08-02 17:03:35 +00003718 for (size_t i = 0; i < stmt->body.Length(); i++) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003719 if (!EmitCase(stmt, i)) {
3720 return false;
3721 }
3722 }
3723 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003724
dan sinclair41e4d9a2022-05-01 14:40:55 +00003725 line() << "}";
3726
3727 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003728}
3729
3730bool GeneratorImpl::EmitType(std::ostream& out,
3731 const sem::Type* type,
dan sinclairff7cf212022-10-03 14:05:23 +00003732 ast::AddressSpace address_space,
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003733 ast::Access access,
3734 const std::string& name,
3735 bool* name_printed /* = nullptr */) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003736 if (name_printed) {
3737 *name_printed = false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003738 }
dan sinclairff7cf212022-10-03 14:05:23 +00003739 switch (address_space) {
3740 case ast::AddressSpace::kStorage:
dan sinclair41e4d9a2022-05-01 14:40:55 +00003741 if (access != ast::Access::kRead) {
3742 out << "RW";
3743 }
3744 out << "ByteAddressBuffer";
3745 return true;
dan sinclairff7cf212022-10-03 14:05:23 +00003746 case ast::AddressSpace::kUniform: {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003747 auto array_length = (type->Size() + 15) / 16;
3748 out << "uint4 " << name << "[" << array_length << "]";
3749 if (name_printed) {
3750 *name_printed = true;
3751 }
3752 return true;
3753 }
3754 default:
3755 break;
3756 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003757
dan sinclair41e4d9a2022-05-01 14:40:55 +00003758 return Switch(
3759 type,
3760 [&](const sem::Array* ary) {
3761 const sem::Type* base_type = ary;
3762 std::vector<uint32_t> sizes;
3763 while (auto* arr = base_type->As<sem::Array>()) {
3764 if (arr->IsRuntimeSized()) {
3765 TINT_ICE(Writer, diagnostics_)
dan sinclair78f80672022-09-22 22:28:21 +00003766 << "runtime arrays may only exist in storage buffers, which should have "
Ben Clayton3a68ab42022-06-24 08:30:28 +00003767 "been transformed into a ByteAddressBuffer";
dan sinclair41e4d9a2022-05-01 14:40:55 +00003768 return false;
3769 }
dan sinclair78f80672022-09-22 22:28:21 +00003770 const auto count = arr->ConstantCount();
3771 if (!count) {
3772 diagnostics_.add_error(diag::System::Writer,
3773 sem::Array::kErrExpectedConstantCount);
3774 return false;
3775 }
3776
3777 sizes.push_back(count.value());
dan sinclair41e4d9a2022-05-01 14:40:55 +00003778 base_type = arr->ElemType();
3779 }
dan sinclairff7cf212022-10-03 14:05:23 +00003780 if (!EmitType(out, base_type, address_space, access, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003781 return false;
3782 }
3783 if (!name.empty()) {
3784 out << " " << name;
3785 if (name_printed) {
3786 *name_printed = true;
3787 }
3788 }
3789 for (uint32_t size : sizes) {
3790 out << "[" << size << "]";
3791 }
3792 return true;
3793 },
3794 [&](const sem::Bool*) {
3795 out << "bool";
3796 return true;
3797 },
3798 [&](const sem::F32*) {
3799 out << "float";
3800 return true;
3801 },
Zhaoming Jiang62bfd312022-05-13 12:01:11 +00003802 [&](const sem::F16*) {
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003803 out << "float16_t";
3804 return true;
Zhaoming Jiang62bfd312022-05-13 12:01:11 +00003805 },
dan sinclair41e4d9a2022-05-01 14:40:55 +00003806 [&](const sem::I32*) {
3807 out << "int";
3808 return true;
3809 },
3810 [&](const sem::Matrix* mat) {
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003811 if (mat->type()->Is<sem::F16>()) {
3812 // Use matrix<type, N, M> for f16 matrix
3813 out << "matrix<";
dan sinclairff7cf212022-10-03 14:05:23 +00003814 if (!EmitType(out, mat->type(), address_space, access, "")) {
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003815 return false;
3816 }
3817 out << ", " << mat->columns() << ", " << mat->rows() << ">";
3818 return true;
3819 }
dan sinclairff7cf212022-10-03 14:05:23 +00003820 if (!EmitType(out, mat->type(), address_space, access, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003821 return false;
3822 }
3823 // Note: HLSL's matrices are declared as <type>NxM, where N is the
3824 // number of rows and M is the number of columns. Despite HLSL's
3825 // matrices being column-major by default, the index operator and
dan sinclair6e77b472022-10-20 13:38:28 +00003826 // initializers actually operate on row-vectors, where as WGSL operates
dan sinclair41e4d9a2022-05-01 14:40:55 +00003827 // on column vectors. To simplify everything we use the transpose of the
3828 // matrices. See:
3829 // https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-per-component-math#matrix-ordering
3830 out << mat->columns() << "x" << mat->rows();
3831 return true;
3832 },
3833 [&](const sem::Pointer*) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003834 TINT_ICE(Writer, diagnostics_)
dan sinclair41e4d9a2022-05-01 14:40:55 +00003835 << "Attempting to emit pointer type. These should have been "
3836 "removed with the InlinePointerLets transform";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003837 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003838 },
3839 [&](const sem::Sampler* sampler) {
3840 out << "Sampler";
3841 if (sampler->IsComparison()) {
3842 out << "Comparison";
3843 }
3844 out << "State";
3845 return true;
3846 },
3847 [&](const sem::Struct* str) {
3848 out << StructName(str);
3849 return true;
3850 },
3851 [&](const sem::Texture* tex) {
3852 if (tex->Is<sem::ExternalTexture>()) {
3853 TINT_ICE(Writer, diagnostics_)
3854 << "Multiplanar external texture transform was not run.";
3855 return false;
3856 }
Brandon Jones6661b282022-02-25 20:14:52 +00003857
dan sinclair41e4d9a2022-05-01 14:40:55 +00003858 auto* storage = tex->As<sem::StorageTexture>();
3859 auto* ms = tex->As<sem::MultisampledTexture>();
3860 auto* depth_ms = tex->As<sem::DepthMultisampledTexture>();
3861 auto* sampled = tex->As<sem::SampledTexture>();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003862
dan sinclair41e4d9a2022-05-01 14:40:55 +00003863 if (storage && storage->access() != ast::Access::kRead) {
3864 out << "RW";
3865 }
3866 out << "Texture";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003867
dan sinclair41e4d9a2022-05-01 14:40:55 +00003868 switch (tex->dim()) {
3869 case ast::TextureDimension::k1d:
3870 out << "1D";
3871 break;
3872 case ast::TextureDimension::k2d:
3873 out << ((ms || depth_ms) ? "2DMS" : "2D");
3874 break;
3875 case ast::TextureDimension::k2dArray:
3876 out << ((ms || depth_ms) ? "2DMSArray" : "2DArray");
3877 break;
3878 case ast::TextureDimension::k3d:
3879 out << "3D";
3880 break;
3881 case ast::TextureDimension::kCube:
3882 out << "Cube";
3883 break;
3884 case ast::TextureDimension::kCubeArray:
3885 out << "CubeArray";
3886 break;
3887 default:
3888 TINT_UNREACHABLE(Writer, diagnostics_)
3889 << "unexpected TextureDimension " << tex->dim();
3890 return false;
3891 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003892
dan sinclair41e4d9a2022-05-01 14:40:55 +00003893 if (storage) {
3894 auto* component = image_format_to_rwtexture_type(storage->texel_format());
3895 if (component == nullptr) {
3896 TINT_ICE(Writer, diagnostics_) << "Unsupported StorageTexture TexelFormat: "
3897 << static_cast<int>(storage->texel_format());
3898 return false;
3899 }
3900 out << "<" << component << ">";
3901 } else if (depth_ms) {
3902 out << "<float4>";
3903 } else if (sampled || ms) {
3904 auto* subtype = sampled ? sampled->type() : ms->type();
3905 out << "<";
3906 if (subtype->Is<sem::F32>()) {
3907 out << "float4";
3908 } else if (subtype->Is<sem::I32>()) {
3909 out << "int4";
3910 } else if (subtype->Is<sem::U32>()) {
3911 out << "uint4";
3912 } else {
3913 TINT_ICE(Writer, diagnostics_) << "Unsupported multisampled texture type";
3914 return false;
3915 }
3916 out << ">";
3917 }
3918 return true;
3919 },
3920 [&](const sem::U32*) {
3921 out << "uint";
3922 return true;
3923 },
3924 [&](const sem::Vector* vec) {
3925 auto width = vec->Width();
3926 if (vec->type()->Is<sem::F32>() && width >= 1 && width <= 4) {
3927 out << "float" << width;
3928 } else if (vec->type()->Is<sem::I32>() && width >= 1 && width <= 4) {
3929 out << "int" << width;
3930 } else if (vec->type()->Is<sem::U32>() && width >= 1 && width <= 4) {
3931 out << "uint" << width;
3932 } else if (vec->type()->Is<sem::Bool>() && width >= 1 && width <= 4) {
3933 out << "bool" << width;
3934 } else {
Zhaoming Jianga5988a32022-07-11 15:43:38 +00003935 // For example, use "vector<float16_t, N>" for f16 vector.
dan sinclair41e4d9a2022-05-01 14:40:55 +00003936 out << "vector<";
dan sinclairff7cf212022-10-03 14:05:23 +00003937 if (!EmitType(out, vec->type(), address_space, access, "")) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003938 return false;
3939 }
3940 out << ", " << width << ">";
3941 }
3942 return true;
3943 },
3944 [&](const sem::Atomic* atomic) {
dan sinclairff7cf212022-10-03 14:05:23 +00003945 return EmitType(out, atomic->Type(), address_space, access, name);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003946 },
3947 [&](const sem::Void*) {
3948 out << "void";
3949 return true;
3950 },
3951 [&](Default) {
3952 diagnostics_.add_error(diag::System::Writer, "unknown type in EmitType");
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003953 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003954 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003955}
3956
3957bool GeneratorImpl::EmitTypeAndName(std::ostream& out,
3958 const sem::Type* type,
dan sinclairff7cf212022-10-03 14:05:23 +00003959 ast::AddressSpace address_space,
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003960 ast::Access access,
3961 const std::string& name) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003962 bool name_printed = false;
dan sinclairff7cf212022-10-03 14:05:23 +00003963 if (!EmitType(out, type, address_space, access, name, &name_printed)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003964 return false;
3965 }
3966 if (!name.empty() && !name_printed) {
3967 out << " " << name;
3968 }
3969 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00003970}
3971
3972bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003973 line(b) << "struct " << StructName(str) << " {";
3974 {
3975 ScopedIndent si(b);
3976 for (auto* mem : str->Members()) {
3977 auto mem_name = builder_.Symbols().NameFor(mem->Name());
dan sinclair41e4d9a2022-05-01 14:40:55 +00003978 auto* ty = mem->Type();
dan sinclair41e4d9a2022-05-01 14:40:55 +00003979 auto out = line(b);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003980 std::string pre, post;
dan sinclair41e4d9a2022-05-01 14:40:55 +00003981 if (auto* decl = mem->Declaration()) {
3982 for (auto* attr : decl->attributes) {
dan sinclairf9eeed62022-09-07 22:25:24 +00003983 if (attr->Is<ast::LocationAttribute>()) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00003984 auto& pipeline_stage_uses = str->PipelineStageUses();
3985 if (pipeline_stage_uses.size() != 1) {
3986 TINT_ICE(Writer, diagnostics_) << "invalid entry point IO struct uses";
3987 }
3988
dan sinclairf9eeed62022-09-07 22:25:24 +00003989 auto loc = mem->Location().value();
dan sinclair41e4d9a2022-05-01 14:40:55 +00003990 if (pipeline_stage_uses.count(sem::PipelineStageUsage::kVertexInput)) {
dan sinclairf9eeed62022-09-07 22:25:24 +00003991 post += " : TEXCOORD" + std::to_string(loc);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003992 } else if (pipeline_stage_uses.count(
3993 sem::PipelineStageUsage::kVertexOutput)) {
dan sinclairf9eeed62022-09-07 22:25:24 +00003994 post += " : TEXCOORD" + std::to_string(loc);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003995 } else if (pipeline_stage_uses.count(
3996 sem::PipelineStageUsage::kFragmentInput)) {
dan sinclairf9eeed62022-09-07 22:25:24 +00003997 post += " : TEXCOORD" + std::to_string(loc);
dan sinclair41e4d9a2022-05-01 14:40:55 +00003998 } else if (pipeline_stage_uses.count(
3999 sem::PipelineStageUsage::kFragmentOutput)) {
dan sinclairf9eeed62022-09-07 22:25:24 +00004000 post += " : SV_Target" + std::to_string(loc);
dan sinclair41e4d9a2022-05-01 14:40:55 +00004001 } else {
4002 TINT_ICE(Writer, diagnostics_) << "invalid use of location attribute";
4003 }
4004 } else if (auto* builtin = attr->As<ast::BuiltinAttribute>()) {
4005 auto name = builtin_to_attribute(builtin->builtin);
4006 if (name.empty()) {
4007 diagnostics_.add_error(diag::System::Writer, "unsupported builtin");
4008 return false;
4009 }
4010 post += " : " + name;
4011 } else if (auto* interpolate = attr->As<ast::InterpolateAttribute>()) {
4012 auto mod =
4013 interpolation_to_modifiers(interpolate->type, interpolate->sampling);
4014 if (mod.empty()) {
4015 diagnostics_.add_error(diag::System::Writer,
4016 "unsupported interpolation");
4017 return false;
4018 }
4019 pre += mod;
4020
4021 } else if (attr->Is<ast::InvariantAttribute>()) {
4022 // Note: `precise` is not exactly the same as `invariant`, but is
4023 // stricter and therefore provides the necessary guarantees.
4024 // See discussion here: https://github.com/gpuweb/gpuweb/issues/893
4025 pre += "precise ";
4026 } else if (!attr->IsAnyOf<ast::StructMemberAlignAttribute,
4027 ast::StructMemberOffsetAttribute,
4028 ast::StructMemberSizeAttribute>()) {
4029 TINT_ICE(Writer, diagnostics_)
4030 << "unhandled struct member attribute: " << attr->Name();
4031 return false;
4032 }
4033 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004034 }
4035
dan sinclair41e4d9a2022-05-01 14:40:55 +00004036 out << pre;
dan sinclairff7cf212022-10-03 14:05:23 +00004037 if (!EmitTypeAndName(out, ty, ast::AddressSpace::kNone, ast::Access::kReadWrite,
dan sinclair41e4d9a2022-05-01 14:40:55 +00004038 mem_name)) {
4039 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004040 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00004041 out << post << ";";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004042 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004043 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004044
dan sinclair41e4d9a2022-05-01 14:40:55 +00004045 line(b) << "};";
dan sinclair41e4d9a2022-05-01 14:40:55 +00004046 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004047}
4048
Antonio Maioranof25140f2022-06-03 14:47:01 +00004049bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
4050 auto it = emitted_structs_.emplace(str);
4051 if (!it.second) {
4052 return true;
4053 }
4054 return EmitStructType(buffer, str);
4055}
4056
dan sinclair41e4d9a2022-05-01 14:40:55 +00004057bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
4058 switch (expr->op) {
4059 case ast::UnaryOp::kIndirection:
4060 case ast::UnaryOp::kAddressOf:
4061 return EmitExpression(out, expr->expr);
4062 case ast::UnaryOp::kComplement:
4063 out << "~";
4064 break;
4065 case ast::UnaryOp::kNot:
4066 out << "!";
4067 break;
4068 case ast::UnaryOp::kNegation:
4069 out << "-";
4070 break;
4071 }
4072 out << "(";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004073
dan sinclair41e4d9a2022-05-01 14:40:55 +00004074 if (!EmitExpression(out, expr->expr)) {
4075 return false;
4076 }
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004077
dan sinclair41e4d9a2022-05-01 14:40:55 +00004078 out << ")";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004079
dan sinclair41e4d9a2022-05-01 14:40:55 +00004080 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004081}
4082
Ben Claytondcdf66e2022-06-17 12:48:51 +00004083bool GeneratorImpl::EmitVar(const ast::Var* var) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00004084 auto* sem = builder_.Sem().Get(var);
4085 auto* type = sem->Type()->UnwrapRef();
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004086
dan sinclair41e4d9a2022-05-01 14:40:55 +00004087 auto out = line();
dan sinclairff7cf212022-10-03 14:05:23 +00004088 if (!EmitTypeAndName(out, type, sem->AddressSpace(), sem->Access(),
dan sinclair41e4d9a2022-05-01 14:40:55 +00004089 builder_.Symbols().NameFor(var->symbol))) {
4090 return false;
4091 }
4092
4093 out << " = ";
4094
dan sinclair6e77b472022-10-20 13:38:28 +00004095 if (var->initializer) {
4096 if (!EmitExpression(out, var->initializer)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00004097 return false;
4098 }
4099 } else {
4100 if (!EmitZeroValue(out, type)) {
4101 return false;
4102 }
4103 }
4104 out << ";";
4105
4106 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004107}
4108
Ben Claytondcdf66e2022-06-17 12:48:51 +00004109bool GeneratorImpl::EmitLet(const ast::Let* let) {
4110 auto* sem = builder_.Sem().Get(let);
4111 auto* type = sem->Type()->UnwrapRef();
4112
4113 auto out = line();
4114 out << "const ";
Ben Claytond2e0db32022-10-12 18:49:15 +00004115 if (!EmitTypeAndName(out, type, ast::AddressSpace::kNone, ast::Access::kUndefined,
Ben Claytondcdf66e2022-06-17 12:48:51 +00004116 builder_.Symbols().NameFor(let->symbol))) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004117 return false;
dan sinclair41e4d9a2022-05-01 14:40:55 +00004118 }
Ben Claytondcdf66e2022-06-17 12:48:51 +00004119 out << " = ";
dan sinclair6e77b472022-10-20 13:38:28 +00004120 if (!EmitExpression(out, let->initializer)) {
Ben Claytondcdf66e2022-06-17 12:48:51 +00004121 return false;
4122 }
4123 out << ";";
dan sinclair41e4d9a2022-05-01 14:40:55 +00004124
Ben Claytondcdf66e2022-06-17 12:48:51 +00004125 return true;
4126}
4127
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004128template <typename F>
4129bool GeneratorImpl::CallBuiltinHelper(std::ostream& out,
4130 const ast::CallExpression* call,
4131 const sem::Builtin* builtin,
4132 F&& build) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00004133 // Generate the helper function if it hasn't been created already
4134 auto fn = utils::GetOrCreate(builtins_, builtin, [&]() -> std::string {
4135 TextBuffer b;
4136 TINT_DEFER(helpers_.Append(b));
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004137
dan sinclair41e4d9a2022-05-01 14:40:55 +00004138 auto fn_name = UniqueIdentifier(std::string("tint_") + sem::str(builtin->Type()));
4139 std::vector<std::string> parameter_names;
4140 {
4141 auto decl = line(&b);
dan sinclairff7cf212022-10-03 14:05:23 +00004142 if (!EmitTypeAndName(decl, builtin->ReturnType(), ast::AddressSpace::kNone,
Ben Claytond2e0db32022-10-12 18:49:15 +00004143 ast::Access::kUndefined, fn_name)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00004144 return "";
4145 }
4146 {
4147 ScopedParen sp(decl);
4148 for (auto* param : builtin->Parameters()) {
4149 if (!parameter_names.empty()) {
4150 decl << ", ";
4151 }
4152 auto param_name = "param_" + std::to_string(parameter_names.size());
4153 const auto* ty = param->Type();
4154 if (auto* ptr = ty->As<sem::Pointer>()) {
4155 decl << "inout ";
4156 ty = ptr->StoreType();
4157 }
Ben Claytond2e0db32022-10-12 18:49:15 +00004158 if (!EmitTypeAndName(decl, ty, ast::AddressSpace::kNone,
4159 ast::Access::kUndefined, param_name)) {
dan sinclair41e4d9a2022-05-01 14:40:55 +00004160 return "";
4161 }
4162 parameter_names.emplace_back(std::move(param_name));
4163 }
4164 }
4165 decl << " {";
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004166 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00004167 {
4168 ScopedIndent si(&b);
4169 if (!build(&b, parameter_names)) {
4170 return "";
4171 }
4172 }
4173 line(&b) << "}";
4174 line(&b);
4175 return fn_name;
4176 });
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004177
dan sinclair41e4d9a2022-05-01 14:40:55 +00004178 if (fn.empty()) {
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004179 return false;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004180 }
dan sinclair41e4d9a2022-05-01 14:40:55 +00004181
4182 // Call the helper
4183 out << fn;
4184 {
4185 ScopedParen sp(out);
4186 bool first = true;
4187 for (auto* arg : call->args) {
4188 if (!first) {
4189 out << ", ";
4190 }
4191 first = false;
4192 if (!EmitExpression(out, arg)) {
4193 return false;
4194 }
4195 }
4196 }
4197 return true;
Ryan Harrisondbc13af2022-02-21 15:19:07 +00004198}
4199
dan sinclair6a5bef12022-04-07 14:30:24 +00004200} // namespace tint::writer::hlsl