| // Copyright 2020 The Tint Authors. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| |
| #include "src/writer/msl/generator_impl.h" |
| |
| #include <algorithm> |
| #include <cmath> |
| #include <iomanip> |
| #include <limits> |
| #include <utility> |
| #include <vector> |
| |
| #include "src/ast/alias.h" |
| #include "src/ast/bool_literal_expression.h" |
| #include "src/ast/call_statement.h" |
| #include "src/ast/disable_validation_decoration.h" |
| #include "src/ast/fallthrough_statement.h" |
| #include "src/ast/float_literal_expression.h" |
| #include "src/ast/interpolate_decoration.h" |
| #include "src/ast/module.h" |
| #include "src/ast/override_decoration.h" |
| #include "src/ast/sint_literal_expression.h" |
| #include "src/ast/uint_literal_expression.h" |
| #include "src/ast/variable_decl_statement.h" |
| #include "src/ast/void.h" |
| #include "src/sem/array.h" |
| #include "src/sem/atomic_type.h" |
| #include "src/sem/bool_type.h" |
| #include "src/sem/call.h" |
| #include "src/sem/depth_multisampled_texture_type.h" |
| #include "src/sem/depth_texture_type.h" |
| #include "src/sem/f32_type.h" |
| #include "src/sem/function.h" |
| #include "src/sem/i32_type.h" |
| #include "src/sem/matrix_type.h" |
| #include "src/sem/member_accessor_expression.h" |
| #include "src/sem/multisampled_texture_type.h" |
| #include "src/sem/pointer_type.h" |
| #include "src/sem/reference_type.h" |
| #include "src/sem/sampled_texture_type.h" |
| #include "src/sem/storage_texture_type.h" |
| #include "src/sem/struct.h" |
| #include "src/sem/type_constructor.h" |
| #include "src/sem/type_conversion.h" |
| #include "src/sem/u32_type.h" |
| #include "src/sem/variable.h" |
| #include "src/sem/vector_type.h" |
| #include "src/sem/void_type.h" |
| #include "src/transform/array_length_from_uniform.h" |
| #include "src/transform/canonicalize_entry_point_io.h" |
| #include "src/transform/external_texture_transform.h" |
| #include "src/transform/inline_pointer_lets.h" |
| #include "src/transform/manager.h" |
| #include "src/transform/module_scope_var_to_entry_point_param.h" |
| #include "src/transform/pad_array_elements.h" |
| #include "src/transform/promote_initializers_to_const_var.h" |
| #include "src/transform/remove_phonies.h" |
| #include "src/transform/simplify.h" |
| #include "src/transform/vectorize_scalar_matrix_constructors.h" |
| #include "src/transform/wrap_arrays_in_structs.h" |
| #include "src/transform/zero_init_workgroup_memory.h" |
| #include "src/utils/defer.h" |
| #include "src/utils/get_or_create.h" |
| #include "src/utils/scoped_assignment.h" |
| #include "src/writer/float_to_string.h" |
| |
| namespace tint { |
| namespace writer { |
| namespace msl { |
| namespace { |
| |
| bool last_is_break_or_fallthrough(const ast::BlockStatement* stmts) { |
| return IsAnyOf<ast::BreakStatement, ast::FallthroughStatement>(stmts->Last()); |
| } |
| |
| class ScopedBitCast { |
| public: |
| ScopedBitCast(GeneratorImpl* generator, |
| std::ostream& stream, |
| const sem::Type* curr_type, |
| const sem::Type* target_type) |
| : s(stream) { |
| auto* target_vec_type = target_type->As<sem::Vector>(); |
| |
| // If we need to promote from scalar to vector, bitcast the scalar to the |
| // vector element type. |
| if (curr_type->is_scalar() && target_vec_type) { |
| target_type = target_vec_type->type(); |
| } |
| |
| // Bit cast |
| s << "as_type<"; |
| generator->EmitType(s, target_type, ""); |
| s << ">("; |
| } |
| |
| ~ScopedBitCast() { s << ")"; } |
| |
| private: |
| std::ostream& s; |
| }; |
| } // namespace |
| |
| SanitizedResult Sanitize(const Program* in, |
| uint32_t buffer_size_ubo_index, |
| uint32_t fixed_sample_mask, |
| bool emit_vertex_point_size, |
| bool disable_workgroup_init) { |
| transform::Manager manager; |
| transform::DataMap internal_inputs; |
| |
| // Build the configs for the internal transforms. |
| auto array_length_from_uniform_cfg = |
| transform::ArrayLengthFromUniform::Config( |
| sem::BindingPoint{0, buffer_size_ubo_index}); |
| auto entry_point_io_cfg = transform::CanonicalizeEntryPointIO::Config( |
| transform::CanonicalizeEntryPointIO::ShaderStyle::kMsl, fixed_sample_mask, |
| emit_vertex_point_size); |
| |
| // Use the SSBO binding numbers as the indices for the buffer size lookups. |
| for (auto* var : in->AST().GlobalVariables()) { |
| auto* global = in->Sem().Get<sem::GlobalVariable>(var); |
| if (global && global->StorageClass() == ast::StorageClass::kStorage) { |
| array_length_from_uniform_cfg.bindpoint_to_size_index.emplace( |
| global->BindingPoint(), global->BindingPoint().binding); |
| } |
| } |
| |
| if (!disable_workgroup_init) { |
| // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as |
| // ZeroInitWorkgroupMemory may inject new builtin parameters. |
| manager.Add<transform::ZeroInitWorkgroupMemory>(); |
| } |
| manager.Add<transform::CanonicalizeEntryPointIO>(); |
| manager.Add<transform::ExternalTextureTransform>(); |
| manager.Add<transform::PromoteInitializersToConstVar>(); |
| manager.Add<transform::VectorizeScalarMatrixConstructors>(); |
| manager.Add<transform::WrapArraysInStructs>(); |
| manager.Add<transform::PadArrayElements>(); |
| manager.Add<transform::InlinePointerLets>(); |
| manager.Add<transform::RemovePhonies>(); |
| manager.Add<transform::Simplify>(); |
| // ArrayLengthFromUniform must come after InlinePointerLets and Simplify, as |
| // it assumes that the form of the array length argument is &var.array. |
| manager.Add<transform::ArrayLengthFromUniform>(); |
| manager.Add<transform::ModuleScopeVarToEntryPointParam>(); |
| internal_inputs.Add<transform::ArrayLengthFromUniform::Config>( |
| std::move(array_length_from_uniform_cfg)); |
| internal_inputs.Add<transform::CanonicalizeEntryPointIO::Config>( |
| std::move(entry_point_io_cfg)); |
| auto out = manager.Run(in, internal_inputs); |
| if (!out.program.IsValid()) { |
| return {std::move(out.program)}; |
| } |
| |
| return {std::move(out.program), |
| out.data.Get<transform::ArrayLengthFromUniform::Result>() |
| ->needs_buffer_sizes}; |
| } |
| |
| GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {} |
| |
| GeneratorImpl::~GeneratorImpl() = default; |
| |
| bool GeneratorImpl::Generate() { |
| line() << "#include <metal_stdlib>"; |
| line(); |
| line() << "using namespace metal;"; |
| |
| auto helpers_insertion_point = current_buffer_->lines.size(); |
| |
| for (auto* const type_decl : program_->AST().TypeDecls()) { |
| if (!type_decl->Is<ast::Alias>()) { |
| if (!EmitTypeDecl(TypeOf(type_decl))) { |
| return false; |
| } |
| } |
| } |
| |
| if (!program_->AST().TypeDecls().empty()) { |
| line(); |
| } |
| |
| for (auto* var : program_->AST().GlobalVariables()) { |
| if (var->is_const) { |
| if (!EmitProgramConstVariable(var)) { |
| return false; |
| } |
| } else { |
| // These are pushed into the entry point by sanitizer transforms. |
| TINT_ICE(Writer, diagnostics_) << "module-scope variables should have " |
| "been handled by the MSL sanitizer"; |
| break; |
| } |
| } |
| |
| for (auto* func : program_->AST().Functions()) { |
| if (!func->IsEntryPoint()) { |
| if (!EmitFunction(func)) { |
| return false; |
| } |
| } else { |
| if (!EmitEntryPointFunction(func)) { |
| return false; |
| } |
| } |
| line(); |
| } |
| |
| if (!helpers_.lines.empty()) { |
| current_buffer_->Insert("", helpers_insertion_point++, 0); |
| current_buffer_->Insert(helpers_, helpers_insertion_point++, 0); |
| } |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitTypeDecl(const sem::Type* ty) { |
| if (auto* str = ty->As<sem::Struct>()) { |
| if (!EmitStructType(current_buffer_, str)) { |
| return false; |
| } |
| } else { |
| diagnostics_.add_error(diag::System::Writer, |
| "unknown alias type: " + ty->type_name()); |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitIndexAccessor( |
| std::ostream& out, |
| const ast::IndexAccessorExpression* expr) { |
| bool paren_lhs = |
| !expr->object->IsAnyOf<ast::IndexAccessorExpression, ast::CallExpression, |
| ast::IdentifierExpression, |
| ast::MemberAccessorExpression>(); |
| |
| if (paren_lhs) { |
| out << "("; |
| } |
| if (!EmitExpression(out, expr->object)) { |
| return false; |
| } |
| if (paren_lhs) { |
| out << ")"; |
| } |
| |
| out << "["; |
| |
| if (!EmitExpression(out, expr->index)) { |
| return false; |
| } |
| out << "]"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitBitcast(std::ostream& out, |
| const ast::BitcastExpression* expr) { |
| out << "as_type<"; |
| if (!EmitType(out, TypeOf(expr)->UnwrapRef(), "")) { |
| return false; |
| } |
| |
| out << ">("; |
| if (!EmitExpression(out, expr->expr)) { |
| return false; |
| } |
| |
| out << ")"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitAssign(const ast::AssignmentStatement* stmt) { |
| auto out = line(); |
| |
| if (!EmitExpression(out, stmt->lhs)) { |
| return false; |
| } |
| |
| out << " = "; |
| |
| if (!EmitExpression(out, stmt->rhs)) { |
| return false; |
| } |
| |
| out << ";"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitBinary(std::ostream& out, |
| const ast::BinaryExpression* expr) { |
| auto emit_op = [&] { |
| out << " "; |
| |
| switch (expr->op) { |
| case ast::BinaryOp::kAnd: |
| out << "&"; |
| break; |
| case ast::BinaryOp::kOr: |
| out << "|"; |
| break; |
| case ast::BinaryOp::kXor: |
| out << "^"; |
| break; |
| case ast::BinaryOp::kLogicalAnd: |
| out << "&&"; |
| break; |
| case ast::BinaryOp::kLogicalOr: |
| out << "||"; |
| break; |
| case ast::BinaryOp::kEqual: |
| out << "=="; |
| break; |
| case ast::BinaryOp::kNotEqual: |
| out << "!="; |
| break; |
| case ast::BinaryOp::kLessThan: |
| out << "<"; |
| break; |
| case ast::BinaryOp::kGreaterThan: |
| out << ">"; |
| break; |
| case ast::BinaryOp::kLessThanEqual: |
| out << "<="; |
| break; |
| case ast::BinaryOp::kGreaterThanEqual: |
| out << ">="; |
| break; |
| case ast::BinaryOp::kShiftLeft: |
| out << "<<"; |
| break; |
| case ast::BinaryOp::kShiftRight: |
| // TODO(dsinclair): MSL is based on C++14, and >> in C++14 has |
| // implementation-defined behaviour for negative LHS. We may have to |
| // generate extra code to implement WGSL-specified behaviour for |
| // negative LHS. |
| out << R"(>>)"; |
| break; |
| |
| case ast::BinaryOp::kAdd: |
| out << "+"; |
| break; |
| case ast::BinaryOp::kSubtract: |
| out << "-"; |
| break; |
| case ast::BinaryOp::kMultiply: |
| out << "*"; |
| break; |
| case ast::BinaryOp::kDivide: |
| out << "/"; |
| break; |
| case ast::BinaryOp::kModulo: |
| out << "%"; |
| break; |
| case ast::BinaryOp::kNone: |
| diagnostics_.add_error(diag::System::Writer, |
| "missing binary operation type"); |
| return false; |
| } |
| out << " "; |
| return true; |
| }; |
| |
| auto signed_type_of = [&](const sem::Type* ty) -> const sem::Type* { |
| if (ty->is_integer_scalar()) { |
| return builder_.create<sem::I32>(); |
| } else if (auto* v = ty->As<sem::Vector>()) { |
| return builder_.create<sem::Vector>(builder_.create<sem::I32>(), |
| v->Width()); |
| } |
| return {}; |
| }; |
| |
| auto unsigned_type_of = [&](const sem::Type* ty) -> const sem::Type* { |
| if (ty->is_integer_scalar()) { |
| return builder_.create<sem::U32>(); |
| } else if (auto* v = ty->As<sem::Vector>()) { |
| return builder_.create<sem::Vector>(builder_.create<sem::U32>(), |
| v->Width()); |
| } |
| return {}; |
| }; |
| |
| auto* lhs_type = TypeOf(expr->lhs)->UnwrapRef(); |
| auto* rhs_type = TypeOf(expr->rhs)->UnwrapRef(); |
| |
| // Handle fmod |
| if (expr->op == ast::BinaryOp::kModulo && |
| lhs_type->is_float_scalar_or_vector()) { |
| out << "fmod"; |
| ScopedParen sp(out); |
| if (!EmitExpression(out, expr->lhs)) { |
| return false; |
| } |
| out << ", "; |
| if (!EmitExpression(out, expr->rhs)) { |
| return false; |
| } |
| return true; |
| } |
| |
| // Handle +/-/* of signed values |
| if ((expr->IsAdd() || expr->IsSubtract() || expr->IsMultiply()) && |
| lhs_type->is_signed_scalar_or_vector() && |
| rhs_type->is_signed_scalar_or_vector()) { |
| // If lhs or rhs is a vector, use that type (support implicit scalar to |
| // vector promotion) |
| auto* target_type = |
| lhs_type->Is<sem::Vector>() |
| ? lhs_type |
| : (rhs_type->Is<sem::Vector>() ? rhs_type : lhs_type); |
| |
| // WGSL defines behaviour for signed overflow, MSL does not. For these |
| // cases, bitcast operands to unsigned, then cast result to signed. |
| ScopedBitCast outer_int_cast(this, out, target_type, |
| signed_type_of(target_type)); |
| ScopedParen sp(out); |
| { |
| ScopedBitCast lhs_uint_cast(this, out, lhs_type, |
| unsigned_type_of(target_type)); |
| if (!EmitExpression(out, expr->lhs)) { |
| return false; |
| } |
| } |
| if (!emit_op()) { |
| return false; |
| } |
| { |
| ScopedBitCast rhs_uint_cast(this, out, rhs_type, |
| unsigned_type_of(target_type)); |
| if (!EmitExpression(out, expr->rhs)) { |
| return false; |
| } |
| } |
| return true; |
| } |
| |
| // Handle left bit shifting a signed value |
| // TODO(crbug.com/tint/1077): This may not be necessary. The MSL spec |
| // seems to imply that left shifting a signed value is treated the same as |
| // left shifting an unsigned value, but we need to make sure. |
| if (expr->IsShiftLeft() && lhs_type->is_signed_scalar_or_vector()) { |
| // Shift left: discards top bits, so convert first operand to unsigned |
| // first, then convert result back to signed |
| ScopedBitCast outer_int_cast(this, out, lhs_type, signed_type_of(lhs_type)); |
| ScopedParen sp(out); |
| { |
| ScopedBitCast lhs_uint_cast(this, out, lhs_type, |
| unsigned_type_of(lhs_type)); |
| if (!EmitExpression(out, expr->lhs)) { |
| return false; |
| } |
| } |
| if (!emit_op()) { |
| return false; |
| } |
| if (!EmitExpression(out, expr->rhs)) { |
| return false; |
| } |
| return true; |
| } |
| |
| // Emit as usual |
| ScopedParen sp(out); |
| if (!EmitExpression(out, expr->lhs)) { |
| return false; |
| } |
| if (!emit_op()) { |
| return false; |
| } |
| if (!EmitExpression(out, expr->rhs)) { |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitBreak(const ast::BreakStatement*) { |
| line() << "break;"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitCall(std::ostream& out, |
| const ast::CallExpression* expr) { |
| auto* call = program_->Sem().Get(expr); |
| auto* target = call->Target(); |
| |
| if (auto* func = target->As<sem::Function>()) { |
| return EmitFunctionCall(out, call, func); |
| } |
| if (auto* intrinsic = target->As<sem::Intrinsic>()) { |
| return EmitIntrinsicCall(out, call, intrinsic); |
| } |
| if (auto* conv = target->As<sem::TypeConversion>()) { |
| return EmitTypeConversion(out, call, conv); |
| } |
| if (auto* ctor = target->As<sem::TypeConstructor>()) { |
| return EmitTypeConstructor(out, call, ctor); |
| } |
| |
| TINT_ICE(Writer, diagnostics_) |
| << "unhandled call target: " << target->TypeInfo().name; |
| return false; |
| } |
| |
| bool GeneratorImpl::EmitFunctionCall(std::ostream& out, |
| const sem::Call* call, |
| const sem::Function*) { |
| auto* ident = call->Declaration()->target.name; |
| out << program_->Symbols().NameFor(ident->symbol) << "("; |
| |
| bool first = true; |
| for (auto* arg : call->Arguments()) { |
| if (!first) { |
| out << ", "; |
| } |
| first = false; |
| |
| if (!EmitExpression(out, arg->Declaration())) { |
| return false; |
| } |
| } |
| |
| out << ")"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out, |
| const sem::Call* call, |
| const sem::Intrinsic* intrinsic) { |
| auto* expr = call->Declaration(); |
| if (intrinsic->IsAtomic()) { |
| return EmitAtomicCall(out, expr, intrinsic); |
| } |
| if (intrinsic->IsTexture()) { |
| return EmitTextureCall(out, expr, intrinsic); |
| } |
| |
| auto name = generate_builtin_name(intrinsic); |
| |
| switch (intrinsic->Type()) { |
| case sem::IntrinsicType::kDot: |
| return EmitDotCall(out, expr, intrinsic); |
| case sem::IntrinsicType::kModf: |
| return EmitModfCall(out, expr, intrinsic); |
| case sem::IntrinsicType::kFrexp: |
| return EmitFrexpCall(out, expr, intrinsic); |
| |
| case sem::IntrinsicType::kPack2x16float: |
| case sem::IntrinsicType::kUnpack2x16float: { |
| if (intrinsic->Type() == sem::IntrinsicType::kPack2x16float) { |
| out << "as_type<uint>(half2("; |
| } else { |
| out << "float2(as_type<half2>("; |
| } |
| if (!EmitExpression(out, expr->args[0])) { |
| return false; |
| } |
| out << "))"; |
| return true; |
| } |
| // TODO(crbug.com/tint/661): Combine sequential barriers to a single |
| // instruction. |
| case sem::IntrinsicType::kStorageBarrier: { |
| out << "threadgroup_barrier(mem_flags::mem_device)"; |
| return true; |
| } |
| case sem::IntrinsicType::kWorkgroupBarrier: { |
| out << "threadgroup_barrier(mem_flags::mem_threadgroup)"; |
| return true; |
| } |
| case sem::IntrinsicType::kIgnore: { // [DEPRECATED] |
| out << "(void) "; |
| if (!EmitExpression(out, expr->args[0])) { |
| return false; |
| } |
| return true; |
| } |
| |
| case sem::IntrinsicType::kLength: { |
| auto* sem = builder_.Sem().Get(expr->args[0]); |
| if (sem->Type()->UnwrapRef()->is_scalar()) { |
| // Emulate scalar overload using fabs(x). |
| name = "fabs"; |
| } |
| break; |
| } |
| |
| case sem::IntrinsicType::kDistance: { |
| auto* sem = builder_.Sem().Get(expr->args[0]); |
| if (sem->Type()->UnwrapRef()->is_scalar()) { |
| // Emulate scalar overload using fabs(x - y); |
| out << "fabs"; |
| ScopedParen sp(out); |
| if (!EmitExpression(out, expr->args[0])) { |
| return false; |
| } |
| out << " - "; |
| if (!EmitExpression(out, expr->args[1])) { |
| return false; |
| } |
| return true; |
| } |
| break; |
| } |
| |
| default: |
| break; |
| } |
| |
| if (name.empty()) { |
| return false; |
| } |
| |
| out << name << "("; |
| |
| bool first = true; |
| for (auto* arg : expr->args) { |
| if (!first) { |
| out << ", "; |
| } |
| first = false; |
| |
| if (!EmitExpression(out, arg)) { |
| return false; |
| } |
| } |
| |
| out << ")"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitTypeConversion(std::ostream& out, |
| const sem::Call* call, |
| const sem::TypeConversion* conv) { |
| if (!EmitType(out, conv->Target(), "")) { |
| return false; |
| } |
| out << "("; |
| |
| if (!EmitExpression(out, call->Arguments()[0]->Declaration())) { |
| return false; |
| } |
| |
| out << ")"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitTypeConstructor(std::ostream& out, |
| const sem::Call* call, |
| const sem::TypeConstructor* ctor) { |
| auto* type = ctor->ReturnType(); |
| |
| if (type->IsAnyOf<sem::Array, sem::Struct>()) { |
| out << "{"; |
| } else { |
| if (!EmitType(out, type, "")) { |
| return false; |
| } |
| out << "("; |
| } |
| |
| int i = 0; |
| for (auto* arg : call->Arguments()) { |
| if (i > 0) { |
| out << ", "; |
| } |
| |
| if (auto* struct_ty = type->As<sem::Struct>()) { |
| // Emit field designators for structures to account for padding members. |
| auto* member = struct_ty->Members()[i]->Declaration(); |
| auto name = program_->Symbols().NameFor(member->symbol); |
| out << "." << name << "="; |
| } |
| |
| if (!EmitExpression(out, arg->Declaration())) { |
| return false; |
| } |
| |
| i++; |
| } |
| |
| if (type->IsAnyOf<sem::Array, sem::Struct>()) { |
| out << "}"; |
| } else { |
| out << ")"; |
| } |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitAtomicCall(std::ostream& out, |
| const ast::CallExpression* expr, |
| const sem::Intrinsic* intrinsic) { |
| auto call = [&](const std::string& name, bool append_memory_order_relaxed) { |
| out << name; |
| { |
| ScopedParen sp(out); |
| for (size_t i = 0; i < expr->args.size(); i++) { |
| auto* arg = expr->args[i]; |
| if (i > 0) { |
| out << ", "; |
| } |
| if (!EmitExpression(out, arg)) { |
| return false; |
| } |
| } |
| if (append_memory_order_relaxed) { |
| out << ", memory_order_relaxed"; |
| } |
| } |
| return true; |
| }; |
| |
| switch (intrinsic->Type()) { |
| case sem::IntrinsicType::kAtomicLoad: |
| return call("atomic_load_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicStore: |
| return call("atomic_store_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicAdd: |
| return call("atomic_fetch_add_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicSub: |
| return call("atomic_fetch_sub_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicMax: |
| return call("atomic_fetch_max_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicMin: |
| return call("atomic_fetch_min_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicAnd: |
| return call("atomic_fetch_and_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicOr: |
| return call("atomic_fetch_or_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicXor: |
| return call("atomic_fetch_xor_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicExchange: |
| return call("atomic_exchange_explicit", true); |
| |
| case sem::IntrinsicType::kAtomicCompareExchangeWeak: { |
| auto* ptr_ty = TypeOf(expr->args[0])->UnwrapRef()->As<sem::Pointer>(); |
| auto sc = ptr_ty->StorageClass(); |
| |
| auto func = utils::GetOrCreate( |
| atomicCompareExchangeWeak_, sc, [&]() -> std::string { |
| auto name = UniqueIdentifier("atomicCompareExchangeWeak"); |
| auto& buf = helpers_; |
| |
| line(&buf) << "template <typename A, typename T>"; |
| { |
| auto f = line(&buf); |
| f << "vec<T, 2> " << name << "("; |
| if (!EmitStorageClass(f, sc)) { |
| return ""; |
| } |
| f << " A* atomic, T compare, T value) {"; |
| } |
| |
| buf.IncrementIndent(); |
| TINT_DEFER({ |
| buf.DecrementIndent(); |
| line(&buf) << "}"; |
| line(&buf); |
| }); |
| |
| line(&buf) << "T prev_value = compare;"; |
| line(&buf) << "bool matched = " |
| "atomic_compare_exchange_weak_explicit(atomic, " |
| "&prev_value, value, memory_order_relaxed, " |
| "memory_order_relaxed);"; |
| line(&buf) << "return {prev_value, matched};"; |
| return name; |
| }); |
| |
| return call(func, false); |
| } |
| |
| default: |
| break; |
| } |
| |
| TINT_UNREACHABLE(Writer, diagnostics_) |
| << "unsupported atomic intrinsic: " << intrinsic->Type(); |
| return false; |
| } |
| |
| bool GeneratorImpl::EmitTextureCall(std::ostream& out, |
| const ast::CallExpression* expr, |
| const sem::Intrinsic* intrinsic) { |
| using Usage = sem::ParameterUsage; |
| |
| auto& signature = intrinsic->Signature(); |
| auto arguments = expr->args; |
| |
| // Returns the argument with the given usage |
| auto arg = [&](Usage usage) { |
| int idx = signature.IndexOf(usage); |
| return (idx >= 0) ? arguments[idx] : nullptr; |
| }; |
| |
| auto* texture = arg(Usage::kTexture); |
| if (!texture) { |
| TINT_ICE(Writer, diagnostics_) << "missing texture arg"; |
| return false; |
| } |
| |
| auto* texture_type = TypeOf(texture)->UnwrapRef()->As<sem::Texture>(); |
| |
| // Helper to emit the texture expression, wrapped in parentheses if the |
| // expression includes an operator with lower precedence than the member |
| // accessor used for the function calls. |
| auto texture_expr = [&]() { |
| bool paren_lhs = |
| !texture->IsAnyOf<ast::IndexAccessorExpression, ast::CallExpression, |
| ast::IdentifierExpression, |
| ast::MemberAccessorExpression>(); |
| if (paren_lhs) { |
| out << "("; |
| } |
| if (!EmitExpression(out, texture)) { |
| return false; |
| } |
| if (paren_lhs) { |
| out << ")"; |
| } |
| return true; |
| }; |
| |
| switch (intrinsic->Type()) { |
| case sem::IntrinsicType::kTextureDimensions: { |
| std::vector<const char*> dims; |
| switch (texture_type->dim()) { |
| case ast::TextureDimension::kNone: |
| diagnostics_.add_error(diag::System::Writer, |
| "texture dimension is kNone"); |
| return false; |
| case ast::TextureDimension::k1d: |
| dims = {"width"}; |
| break; |
| case ast::TextureDimension::k2d: |
| case ast::TextureDimension::k2dArray: |
| case ast::TextureDimension::kCube: |
| case ast::TextureDimension::kCubeArray: |
| dims = {"width", "height"}; |
| break; |
| case ast::TextureDimension::k3d: |
| dims = {"width", "height", "depth"}; |
| break; |
| } |
| |
| auto get_dim = [&](const char* name) { |
| if (!texture_expr()) { |
| return false; |
| } |
| out << ".get_" << name << "("; |
| if (auto* level = arg(Usage::kLevel)) { |
| if (!EmitExpression(out, level)) { |
| return false; |
| } |
| } |
| out << ")"; |
| return true; |
| }; |
| |
| if (dims.size() == 1) { |
| out << "int("; |
| get_dim(dims[0]); |
| out << ")"; |
| } else { |
| EmitType(out, TypeOf(expr)->UnwrapRef(), ""); |
| out << "("; |
| for (size_t i = 0; i < dims.size(); i++) { |
| if (i > 0) { |
| out << ", "; |
| } |
| get_dim(dims[i]); |
| } |
| out << ")"; |
| } |
| return true; |
| } |
| case sem::IntrinsicType::kTextureNumLayers: { |
| out << "int("; |
| if (!texture_expr()) { |
| return false; |
| } |
| out << ".get_array_size())"; |
| return true; |
| } |
| case sem::IntrinsicType::kTextureNumLevels: { |
| out << "int("; |
| if (!texture_expr()) { |
| return false; |
| } |
| out << ".get_num_mip_levels())"; |
| return true; |
| } |
| case sem::IntrinsicType::kTextureNumSamples: { |
| out << "int("; |
| if (!texture_expr()) { |
| return false; |
| } |
| out << ".get_num_samples())"; |
| return true; |
| } |
| default: |
| break; |
| } |
| |
| if (!texture_expr()) { |
| return false; |
| } |
| |
| bool lod_param_is_named = true; |
| |
| switch (intrinsic->Type()) { |
| case sem::IntrinsicType::kTextureSample: |
| case sem::IntrinsicType::kTextureSampleBias: |
| case sem::IntrinsicType::kTextureSampleLevel: |
| case sem::IntrinsicType::kTextureSampleGrad: |
| out << ".sample("; |
| break; |
| case sem::IntrinsicType::kTextureSampleCompare: |
| case sem::IntrinsicType::kTextureSampleCompareLevel: |
| out << ".sample_compare("; |
| break; |
| case sem::IntrinsicType::kTextureLoad: |
| out << ".read("; |
| lod_param_is_named = false; |
| break; |
| case sem::IntrinsicType::kTextureStore: |
| out << ".write("; |
| break; |
| default: |
| TINT_UNREACHABLE(Writer, diagnostics_) |
| << "Unhandled texture intrinsic '" << intrinsic->str() << "'"; |
| return false; |
| } |
| |
| bool first_arg = true; |
| auto maybe_write_comma = [&] { |
| if (!first_arg) { |
| out << ", "; |
| } |
| first_arg = false; |
| }; |
| |
| for (auto usage : |
| {Usage::kValue, Usage::kSampler, Usage::kCoords, Usage::kArrayIndex, |
| Usage::kDepthRef, Usage::kSampleIndex}) { |
| if (auto* e = arg(usage)) { |
| auto* sem_e = program_->Sem().Get(e); |
| |
| maybe_write_comma(); |
| |
| // Cast the coordinates to unsigned integers if necessary. |
| bool casted = false; |
| if (usage == Usage::kCoords && |
| sem_e->Type()->UnwrapRef()->is_integer_scalar_or_vector()) { |
| casted = true; |
| switch (texture_type->dim()) { |
| case ast::TextureDimension::k1d: |
| out << "uint("; |
| break; |
| case ast::TextureDimension::k2d: |
| case ast::TextureDimension::k2dArray: |
| out << "uint2("; |
| break; |
| case ast::TextureDimension::k3d: |
| out << "uint3("; |
| break; |
| default: |
| TINT_ICE(Writer, diagnostics_) |
| << "unhandled texture dimensionality"; |
| break; |
| } |
| } |
| |
| if (!EmitExpression(out, e)) |
| return false; |
| |
| if (casted) { |
| out << ")"; |
| } |
| } |
| } |
| |
| if (auto* bias = arg(Usage::kBias)) { |
| maybe_write_comma(); |
| out << "bias("; |
| if (!EmitExpression(out, bias)) { |
| return false; |
| } |
| out << ")"; |
| } |
| if (auto* level = arg(Usage::kLevel)) { |
| maybe_write_comma(); |
| if (lod_param_is_named) { |
| out << "level("; |
| } |
| if (!EmitExpression(out, level)) { |
| return false; |
| } |
| if (lod_param_is_named) { |
| out << ")"; |
| } |
| } |
| if (intrinsic->Type() == sem::IntrinsicType::kTextureSampleCompareLevel) { |
| maybe_write_comma(); |
| out << "level(0)"; |
| } |
| if (auto* ddx = arg(Usage::kDdx)) { |
| auto dim = texture_type->dim(); |
| switch (dim) { |
| case ast::TextureDimension::k2d: |
| case ast::TextureDimension::k2dArray: |
| maybe_write_comma(); |
| out << "gradient2d("; |
| break; |
| case ast::TextureDimension::k3d: |
| maybe_write_comma(); |
| out << "gradient3d("; |
| break; |
| case ast::TextureDimension::kCube: |
| case ast::TextureDimension::kCubeArray: |
| maybe_write_comma(); |
| out << "gradientcube("; |
| break; |
| default: { |
| std::stringstream err; |
| err << "MSL does not support gradients for " << dim << " textures"; |
| diagnostics_.add_error(diag::System::Writer, err.str()); |
| return false; |
| } |
| } |
| if (!EmitExpression(out, ddx)) { |
| return false; |
| } |
| out << ", "; |
| if (!EmitExpression(out, arg(Usage::kDdy))) { |
| return false; |
| } |
| out << ")"; |
| } |
| |
| if (auto* offset = arg(Usage::kOffset)) { |
| maybe_write_comma(); |
| if (!EmitExpression(out, offset)) { |
| return false; |
| } |
| } |
| |
| out << ")"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitDotCall(std::ostream& out, |
| const ast::CallExpression* expr, |
| const sem::Intrinsic* intrinsic) { |
| auto* vec_ty = intrinsic->Parameters()[0]->Type()->As<sem::Vector>(); |
| std::string fn = "dot"; |
| if (vec_ty->type()->is_integer_scalar()) { |
| // MSL does not have a builtin for dot() with integer vector types. |
| // Generate the helper function if it hasn't been created already |
| fn = utils::GetOrCreate( |
| int_dot_funcs_, vec_ty->Width(), [&]() -> std::string { |
| TextBuffer b; |
| TINT_DEFER(helpers_.Append(b)); |
| |
| auto fn_name = |
| UniqueIdentifier("tint_dot" + std::to_string(vec_ty->Width())); |
| auto v = "vec<T," + std::to_string(vec_ty->Width()) + ">"; |
| |
| line(&b) << "template<typename T>"; |
| line(&b) << "T " << fn_name << "(" << v << " a, " << v << " b) {"; |
| { |
| auto l = line(&b); |
| l << " return "; |
| for (uint32_t i = 0; i < vec_ty->Width(); i++) { |
| if (i > 0) { |
| l << " + "; |
| } |
| l << "a[" << i << "]*b[" << i << "]"; |
| } |
| l << ";"; |
| } |
| line(&b) << "}"; |
| return fn_name; |
| }); |
| } |
| |
| out << fn << "("; |
| if (!EmitExpression(out, expr->args[0])) { |
| return false; |
| } |
| out << ", "; |
| if (!EmitExpression(out, expr->args[1])) { |
| return false; |
| } |
| out << ")"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitModfCall(std::ostream& out, |
| const ast::CallExpression* expr, |
| const sem::Intrinsic* intrinsic) { |
| return CallIntrinsicHelper( |
| out, expr, intrinsic, |
| [&](TextBuffer* b, const std::vector<std::string>& params) { |
| auto* ty = intrinsic->Parameters()[0]->Type(); |
| auto in = params[0]; |
| |
| std::string width; |
| if (auto* vec = ty->As<sem::Vector>()) { |
| width = std::to_string(vec->Width()); |
| } |
| |
| // Emit the builtin return type unique to this overload. This does not |
| // exist in the AST, so it will not be generated in Generate(). |
| if (!EmitStructType(&helpers_, |
| intrinsic->ReturnType()->As<sem::Struct>())) { |
| return false; |
| } |
| |
| line(b) << "float" << width << " whole;"; |
| line(b) << "float" << width << " fract = modf(" << in << ", whole);"; |
| line(b) << "return {fract, whole};"; |
| return true; |
| }); |
| } |
| |
| bool GeneratorImpl::EmitFrexpCall(std::ostream& out, |
| const ast::CallExpression* expr, |
| const sem::Intrinsic* intrinsic) { |
| return CallIntrinsicHelper( |
| out, expr, intrinsic, |
| [&](TextBuffer* b, const std::vector<std::string>& params) { |
| auto* ty = intrinsic->Parameters()[0]->Type(); |
| auto in = params[0]; |
| |
| std::string width; |
| if (auto* vec = ty->As<sem::Vector>()) { |
| width = std::to_string(vec->Width()); |
| } |
| |
| // Emit the builtin return type unique to this overload. This does not |
| // exist in the AST, so it will not be generated in Generate(). |
| if (!EmitStructType(&helpers_, |
| intrinsic->ReturnType()->As<sem::Struct>())) { |
| return false; |
| } |
| |
| line(b) << "int" << width << " exp;"; |
| line(b) << "float" << width << " sig = frexp(" << in << ", exp);"; |
| line(b) << "return {sig, exp};"; |
| return true; |
| }); |
| } |
| |
| std::string GeneratorImpl::generate_builtin_name( |
| const sem::Intrinsic* intrinsic) { |
| std::string out = ""; |
| switch (intrinsic->Type()) { |
| case sem::IntrinsicType::kAcos: |
| case sem::IntrinsicType::kAll: |
| case sem::IntrinsicType::kAny: |
| case sem::IntrinsicType::kAsin: |
| case sem::IntrinsicType::kAtan: |
| case sem::IntrinsicType::kAtan2: |
| case sem::IntrinsicType::kCeil: |
| case sem::IntrinsicType::kCos: |
| case sem::IntrinsicType::kCosh: |
| case sem::IntrinsicType::kCross: |
| case sem::IntrinsicType::kDeterminant: |
| case sem::IntrinsicType::kDistance: |
| case sem::IntrinsicType::kDot: |
| case sem::IntrinsicType::kExp: |
| case sem::IntrinsicType::kExp2: |
| case sem::IntrinsicType::kFloor: |
| case sem::IntrinsicType::kFma: |
| case sem::IntrinsicType::kFract: |
| case sem::IntrinsicType::kFrexp: |
| case sem::IntrinsicType::kLength: |
| case sem::IntrinsicType::kLdexp: |
| case sem::IntrinsicType::kLog: |
| case sem::IntrinsicType::kLog2: |
| case sem::IntrinsicType::kMix: |
| case sem::IntrinsicType::kModf: |
| case sem::IntrinsicType::kNormalize: |
| case sem::IntrinsicType::kPow: |
| case sem::IntrinsicType::kReflect: |
| case sem::IntrinsicType::kRefract: |
| case sem::IntrinsicType::kSelect: |
| case sem::IntrinsicType::kSin: |
| case sem::IntrinsicType::kSinh: |
| case sem::IntrinsicType::kSqrt: |
| case sem::IntrinsicType::kStep: |
| case sem::IntrinsicType::kTan: |
| case sem::IntrinsicType::kTanh: |
| case sem::IntrinsicType::kTranspose: |
| case sem::IntrinsicType::kTrunc: |
| case sem::IntrinsicType::kSign: |
| case sem::IntrinsicType::kClamp: |
| out += intrinsic->str(); |
| break; |
| case sem::IntrinsicType::kAbs: |
| if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { |
| out += "fabs"; |
| } else { |
| out += "abs"; |
| } |
| break; |
| case sem::IntrinsicType::kCountOneBits: |
| out += "popcount"; |
| break; |
| case sem::IntrinsicType::kDpdx: |
| case sem::IntrinsicType::kDpdxCoarse: |
| case sem::IntrinsicType::kDpdxFine: |
| out += "dfdx"; |
| break; |
| case sem::IntrinsicType::kDpdy: |
| case sem::IntrinsicType::kDpdyCoarse: |
| case sem::IntrinsicType::kDpdyFine: |
| out += "dfdy"; |
| break; |
| case sem::IntrinsicType::kFwidth: |
| case sem::IntrinsicType::kFwidthCoarse: |
| case sem::IntrinsicType::kFwidthFine: |
| out += "fwidth"; |
| break; |
| case sem::IntrinsicType::kIsFinite: |
| out += "isfinite"; |
| break; |
| case sem::IntrinsicType::kIsInf: |
| out += "isinf"; |
| break; |
| case sem::IntrinsicType::kIsNan: |
| out += "isnan"; |
| break; |
| case sem::IntrinsicType::kIsNormal: |
| out += "isnormal"; |
| break; |
| case sem::IntrinsicType::kMax: |
| if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { |
| out += "fmax"; |
| } else { |
| out += "max"; |
| } |
| break; |
| case sem::IntrinsicType::kMin: |
| if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { |
| out += "fmin"; |
| } else { |
| out += "min"; |
| } |
| break; |
| case sem::IntrinsicType::kFaceForward: |
| out += "faceforward"; |
| break; |
| case sem::IntrinsicType::kPack4x8snorm: |
| out += "pack_float_to_snorm4x8"; |
| break; |
| case sem::IntrinsicType::kPack4x8unorm: |
| out += "pack_float_to_unorm4x8"; |
| break; |
| case sem::IntrinsicType::kPack2x16snorm: |
| out += "pack_float_to_snorm2x16"; |
| break; |
| case sem::IntrinsicType::kPack2x16unorm: |
| out += "pack_float_to_unorm2x16"; |
| break; |
| case sem::IntrinsicType::kReverseBits: |
| out += "reverse_bits"; |
| break; |
| case sem::IntrinsicType::kRound: |
| out += "rint"; |
| break; |
| case sem::IntrinsicType::kSmoothStep: |
| out += "smoothstep"; |
| break; |
| case sem::IntrinsicType::kInverseSqrt: |
| out += "rsqrt"; |
| break; |
| case sem::IntrinsicType::kUnpack4x8snorm: |
| out += "unpack_snorm4x8_to_float"; |
| break; |
| case sem::IntrinsicType::kUnpack4x8unorm: |
| out += "unpack_unorm4x8_to_float"; |
| break; |
| case sem::IntrinsicType::kUnpack2x16snorm: |
| out += "unpack_snorm2x16_to_float"; |
| break; |
| case sem::IntrinsicType::kUnpack2x16unorm: |
| out += "unpack_unorm2x16_to_float"; |
| break; |
| default: |
| diagnostics_.add_error( |
| diag::System::Writer, |
| "Unknown import method: " + std::string(intrinsic->str())); |
| return ""; |
| } |
| return out; |
| } |
| |
| bool GeneratorImpl::EmitCase(const ast::CaseStatement* stmt) { |
| if (stmt->IsDefault()) { |
| line() << "default: {"; |
| } else { |
| for (auto* selector : stmt->selectors) { |
| auto out = line(); |
| out << "case "; |
| if (!EmitLiteral(out, selector)) { |
| return false; |
| } |
| out << ":"; |
| if (selector == stmt->selectors.back()) { |
| out << " {"; |
| } |
| } |
| } |
| |
| { |
| ScopedIndent si(this); |
| |
| for (auto* s : stmt->body->statements) { |
| if (!EmitStatement(s)) { |
| return false; |
| } |
| } |
| |
| if (!last_is_break_or_fallthrough(stmt->body)) { |
| line() << "break;"; |
| } |
| } |
| |
| line() << "}"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitContinue(const ast::ContinueStatement*) { |
| if (!emit_continuing_()) { |
| return false; |
| } |
| |
| line() << "continue;"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) { |
| if (type->Is<sem::Bool>()) { |
| out << "false"; |
| } else if (type->Is<sem::F32>()) { |
| out << "0.0f"; |
| } else if (type->Is<sem::I32>()) { |
| out << "0"; |
| } else if (type->Is<sem::U32>()) { |
| out << "0u"; |
| } else if (auto* vec = type->As<sem::Vector>()) { |
| return EmitZeroValue(out, vec->type()); |
| } else if (auto* mat = type->As<sem::Matrix>()) { |
| if (!EmitType(out, mat, "")) { |
| return false; |
| } |
| out << "("; |
| if (!EmitZeroValue(out, mat->type())) { |
| return false; |
| } |
| out << ")"; |
| } else if (auto* arr = type->As<sem::Array>()) { |
| out << "{"; |
| if (!EmitZeroValue(out, arr->ElemType())) { |
| return false; |
| } |
| out << "}"; |
| } else if (type->As<sem::Struct>()) { |
| out << "{}"; |
| } else { |
| diagnostics_.add_error( |
| diag::System::Writer, |
| "Invalid type for zero emission: " + type->type_name()); |
| return false; |
| } |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitLiteral(std::ostream& out, |
| const ast::LiteralExpression* lit) { |
| if (auto* l = lit->As<ast::BoolLiteralExpression>()) { |
| out << (l->value ? "true" : "false"); |
| } else if (auto* fl = lit->As<ast::FloatLiteralExpression>()) { |
| if (std::isinf(fl->value)) { |
| out << (fl->value >= 0 ? "INFINITY" : "-INFINITY"); |
| } else if (std::isnan(fl->value)) { |
| out << "NAN"; |
| } else { |
| out << FloatToString(fl->value) << "f"; |
| } |
| } else if (auto* sl = lit->As<ast::SintLiteralExpression>()) { |
| // MSL (and C++) parse `-2147483648` as a `long` because it parses unary |
| // minus and `2147483648` as separate tokens, and the latter doesn't |
| // fit into an (32-bit) `int`. WGSL, OTOH, parses this as an `i32`. To avoid |
| // issues with `long` to `int` casts, emit `(2147483647 - 1)` instead, which |
| // ensures the expression type is `int`. |
| const auto int_min = std::numeric_limits<int32_t>::min(); |
| if (sl->ValueAsI32() == int_min) { |
| out << "(" << int_min + 1 << " - 1)"; |
| } else { |
| out << sl->value; |
| } |
| } else if (auto* ul = lit->As<ast::UintLiteralExpression>()) { |
| out << ul->value << "u"; |
| } else { |
| diagnostics_.add_error(diag::System::Writer, "unknown literal type"); |
| return false; |
| } |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitExpression(std::ostream& out, |
| const ast::Expression* expr) { |
| if (auto* a = expr->As<ast::IndexAccessorExpression>()) { |
| return EmitIndexAccessor(out, a); |
| } |
| if (auto* b = expr->As<ast::BinaryExpression>()) { |
| return EmitBinary(out, b); |
| } |
| if (auto* b = expr->As<ast::BitcastExpression>()) { |
| return EmitBitcast(out, b); |
| } |
| if (auto* c = expr->As<ast::CallExpression>()) { |
| return EmitCall(out, c); |
| } |
| if (auto* i = expr->As<ast::IdentifierExpression>()) { |
| return EmitIdentifier(out, i); |
| } |
| if (auto* l = expr->As<ast::LiteralExpression>()) { |
| return EmitLiteral(out, l); |
| } |
| if (auto* m = expr->As<ast::MemberAccessorExpression>()) { |
| return EmitMemberAccessor(out, m); |
| } |
| if (auto* u = expr->As<ast::UnaryOpExpression>()) { |
| return EmitUnaryOp(out, u); |
| } |
| |
| diagnostics_.add_error( |
| diag::System::Writer, |
| "unknown expression type: " + std::string(expr->TypeInfo().name)); |
| return false; |
| } |
| |
| void GeneratorImpl::EmitStage(std::ostream& out, ast::PipelineStage stage) { |
| switch (stage) { |
| case ast::PipelineStage::kFragment: |
| out << "fragment"; |
| break; |
| case ast::PipelineStage::kVertex: |
| out << "vertex"; |
| break; |
| case ast::PipelineStage::kCompute: |
| out << "kernel"; |
| break; |
| case ast::PipelineStage::kNone: |
| break; |
| } |
| return; |
| } |
| |
| bool GeneratorImpl::EmitFunction(const ast::Function* func) { |
| auto* func_sem = program_->Sem().Get(func); |
| |
| { |
| auto out = line(); |
| if (!EmitType(out, func_sem->ReturnType(), "")) { |
| return false; |
| } |
| out << " " << program_->Symbols().NameFor(func->symbol) << "("; |
| |
| bool first = true; |
| for (auto* v : func->params) { |
| if (!first) { |
| out << ", "; |
| } |
| first = false; |
| |
| auto* type = program_->Sem().Get(v)->Type(); |
| |
| std::string param_name = |
| "const " + program_->Symbols().NameFor(v->symbol); |
| if (!EmitType(out, type, param_name)) { |
| return false; |
| } |
| // Parameter name is output as part of the type for arrays and pointers. |
| if (!type->Is<sem::Array>() && !type->Is<sem::Pointer>()) { |
| out << " " << program_->Symbols().NameFor(v->symbol); |
| } |
| } |
| |
| out << ") {"; |
| } |
| |
| if (!EmitStatementsWithIndent(func->body->statements)) { |
| return false; |
| } |
| |
| line() << "}"; |
| |
| return true; |
| } |
| |
| std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const { |
| switch (builtin) { |
| case ast::Builtin::kPosition: |
| return "position"; |
| case ast::Builtin::kVertexIndex: |
| return "vertex_id"; |
| case ast::Builtin::kInstanceIndex: |
| return "instance_id"; |
| case ast::Builtin::kFrontFacing: |
| return "front_facing"; |
| case ast::Builtin::kFragDepth: |
| return "depth(any)"; |
| case ast::Builtin::kLocalInvocationId: |
| return "thread_position_in_threadgroup"; |
| case ast::Builtin::kLocalInvocationIndex: |
| return "thread_index_in_threadgroup"; |
| case ast::Builtin::kGlobalInvocationId: |
| return "thread_position_in_grid"; |
| case ast::Builtin::kWorkgroupId: |
| return "threadgroup_position_in_grid"; |
| case ast::Builtin::kNumWorkgroups: |
| return "threadgroups_per_grid"; |
| case ast::Builtin::kSampleIndex: |
| return "sample_id"; |
| case ast::Builtin::kSampleMask: |
| return "sample_mask"; |
| case ast::Builtin::kPointSize: |
| return "point_size"; |
| default: |
| break; |
| } |
| return ""; |
| } |
| |
| std::string GeneratorImpl::interpolation_to_attribute( |
| ast::InterpolationType type, |
| ast::InterpolationSampling sampling) const { |
| std::string attr; |
| switch (sampling) { |
| case ast::InterpolationSampling::kCenter: |
| attr = "center_"; |
| break; |
| case ast::InterpolationSampling::kCentroid: |
| attr = "centroid_"; |
| break; |
| case ast::InterpolationSampling::kSample: |
| attr = "sample_"; |
| break; |
| case ast::InterpolationSampling::kNone: |
| break; |
| } |
| switch (type) { |
| case ast::InterpolationType::kPerspective: |
| attr += "perspective"; |
| break; |
| case ast::InterpolationType::kLinear: |
| attr += "no_perspective"; |
| break; |
| case ast::InterpolationType::kFlat: |
| attr += "flat"; |
| break; |
| } |
| return attr; |
| } |
| |
| bool GeneratorImpl::EmitEntryPointFunction(const ast::Function* func) { |
| auto func_name = program_->Symbols().NameFor(func->symbol); |
| |
| // Returns the binding index of a variable, requiring that the group attribute |
| // have a value of zero. |
| const uint32_t kInvalidBindingIndex = std::numeric_limits<uint32_t>::max(); |
| auto get_binding_index = [&](const ast::Variable* var) -> uint32_t { |
| auto bp = var->BindingPoint(); |
| if (bp.group == nullptr || bp.binding == nullptr) { |
| TINT_ICE(Writer, diagnostics_) |
| << "missing binding attributes for entry point parameter"; |
| return kInvalidBindingIndex; |
| } |
| if (bp.group->value != 0) { |
| TINT_ICE(Writer, diagnostics_) |
| << "encountered non-zero resource group index (use " |
| "BindingRemapper to fix)"; |
| return kInvalidBindingIndex; |
| } |
| return bp.binding->value; |
| }; |
| |
| { |
| auto out = line(); |
| |
| EmitStage(out, func->PipelineStage()); |
| out << " " << func->return_type->FriendlyName(program_->Symbols()); |
| out << " " << func_name << "("; |
| |
| // Emit entry point parameters. |
| bool first = true; |
| for (auto* var : func->params) { |
| if (!first) { |
| out << ", "; |
| } |
| first = false; |
| |
| auto* type = program_->Sem().Get(var)->Type()->UnwrapRef(); |
| |
| auto param_name = program_->Symbols().NameFor(var->symbol); |
| if (!EmitType(out, type, param_name)) { |
| return false; |
| } |
| // Parameter name is output as part of the type for arrays and pointers. |
| if (!type->Is<sem::Array>() && !type->Is<sem::Pointer>()) { |
| out << " " << param_name; |
| } |
| |
| if (type->Is<sem::Struct>()) { |
| out << " [[stage_in]]"; |
| } else if (type->is_handle()) { |
| uint32_t binding = get_binding_index(var); |
| if (binding == kInvalidBindingIndex) { |
| return false; |
| } |
| if (var->type->Is<ast::Sampler>()) { |
| out << " [[sampler(" << binding << ")]]"; |
| } else if (var->type->Is<ast::Texture>()) { |
| out << " [[texture(" << binding << ")]]"; |
| } else { |
| TINT_ICE(Writer, diagnostics_) |
| << "invalid handle type entry point parameter"; |
| return false; |
| } |
| } else if (auto* ptr = var->type->As<ast::Pointer>()) { |
| auto sc = ptr->storage_class; |
| if (sc == ast::StorageClass::kWorkgroup) { |
| auto& allocations = workgroup_allocations_[func_name]; |
| out << " [[threadgroup(" << allocations.size() << ")]]"; |
| allocations.push_back(program_->Sem().Get(ptr->type)->Size()); |
| } else if (sc == ast::StorageClass::kStorage || |
| sc == ast::StorageClass::kUniform) { |
| uint32_t binding = get_binding_index(var); |
| if (binding == kInvalidBindingIndex) { |
| return false; |
| } |
| out << " [[buffer(" << binding << ")]]"; |
| } else { |
| TINT_ICE(Writer, diagnostics_) |
| << "invalid pointer storage class for entry point parameter"; |
| return false; |
| } |
| } else { |
| auto& decos = var->decorations; |
| bool builtin_found = false; |
| for (auto* deco : decos) { |
| auto* builtin = deco->As<ast::BuiltinDecoration>(); |
| if (!builtin) { |
| continue; |
| } |
| |
| builtin_found = true; |
| |
| auto attr = builtin_to_attribute(builtin->builtin); |
| if (attr.empty()) { |
| diagnostics_.add_error(diag::System::Writer, "unknown builtin"); |
| return false; |
| } |
| out << " [[" << attr << "]]"; |
| } |
| if (!builtin_found) { |
| TINT_ICE(Writer, diagnostics_) << "Unsupported entry point parameter"; |
| } |
| } |
| } |
| out << ") {"; |
| } |
| |
| { |
| ScopedIndent si(this); |
| |
| if (!EmitStatements(func->body->statements)) { |
| return false; |
| } |
| |
| if (!Is<ast::ReturnStatement>(func->body->Last())) { |
| ast::ReturnStatement ret(ProgramID{}, Source{}); |
| if (!EmitStatement(&ret)) { |
| return false; |
| } |
| } |
| } |
| |
| line() << "}"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitIdentifier(std::ostream& out, |
| const ast::IdentifierExpression* expr) { |
| out << program_->Symbols().NameFor(expr->symbol); |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitLoop(const ast::LoopStatement* stmt) { |
| auto emit_continuing = [this, stmt]() { |
| if (stmt->continuing && !stmt->continuing->Empty()) { |
| if (!EmitBlock(stmt->continuing)) { |
| return false; |
| } |
| } |
| return true; |
| }; |
| |
| TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing); |
| line() << "while (true) {"; |
| { |
| ScopedIndent si(this); |
| if (!EmitStatements(stmt->body->statements)) { |
| return false; |
| } |
| if (!emit_continuing()) { |
| return false; |
| } |
| } |
| line() << "}"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitForLoop(const ast::ForLoopStatement* stmt) { |
| TextBuffer init_buf; |
| if (auto* init = stmt->initializer) { |
| TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf); |
| if (!EmitStatement(init)) { |
| return false; |
| } |
| } |
| |
| TextBuffer cond_pre; |
| std::stringstream cond_buf; |
| if (auto* cond = stmt->condition) { |
| TINT_SCOPED_ASSIGNMENT(current_buffer_, &cond_pre); |
| if (!EmitExpression(cond_buf, cond)) { |
| return false; |
| } |
| } |
| |
| TextBuffer cont_buf; |
| if (auto* cont = stmt->continuing) { |
| TINT_SCOPED_ASSIGNMENT(current_buffer_, &cont_buf); |
| if (!EmitStatement(cont)) { |
| return false; |
| } |
| } |
| |
| // If the for-loop has a multi-statement conditional and / or continuing, then |
| // we cannot emit this as a regular for-loop in MSL. Instead we need to |
| // generate a `while(true)` loop. |
| bool emit_as_loop = cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1; |
| |
| // If the for-loop has multi-statement initializer, or is going to be emitted |
| // as a `while(true)` loop, then declare the initializer statement(s) before |
| // the loop in a new block. |
| bool nest_in_block = |
| init_buf.lines.size() > 1 || (stmt->initializer && emit_as_loop); |
| if (nest_in_block) { |
| line() << "{"; |
| increment_indent(); |
| current_buffer_->Append(init_buf); |
| init_buf.lines.clear(); // Don't emit the initializer again in the 'for' |
| } |
| TINT_DEFER({ |
| if (nest_in_block) { |
| decrement_indent(); |
| line() << "}"; |
| } |
| }); |
| |
| if (emit_as_loop) { |
| auto emit_continuing = [&]() { |
| current_buffer_->Append(cont_buf); |
| return true; |
| }; |
| |
| TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing); |
| line() << "while (true) {"; |
| increment_indent(); |
| TINT_DEFER({ |
| decrement_indent(); |
| line() << "}"; |
| }); |
| |
| if (stmt->condition) { |
| current_buffer_->Append(cond_pre); |
| line() << "if (!(" << cond_buf.str() << ")) { break; }"; |
| } |
| |
| if (!EmitStatements(stmt->body->statements)) { |
| return false; |
| } |
| |
| if (!emit_continuing()) { |
| return false; |
| } |
| } else { |
| // For-loop can be generated. |
| { |
| auto out = line(); |
| out << "for"; |
| { |
| ScopedParen sp(out); |
| |
| if (!init_buf.lines.empty()) { |
| out << init_buf.lines[0].content << " "; |
| } else { |
| out << "; "; |
| } |
| |
| out << cond_buf.str() << "; "; |
| |
| if (!cont_buf.lines.empty()) { |
| out << TrimSuffix(cont_buf.lines[0].content, ";"); |
| } |
| } |
| out << " {"; |
| } |
| { |
| auto emit_continuing = [] { return true; }; |
| TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing); |
| if (!EmitStatementsWithIndent(stmt->body->statements)) { |
| return false; |
| } |
| } |
| line() << "}"; |
| } |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitDiscard(const ast::DiscardStatement*) { |
| // TODO(dsinclair): Verify this is correct when the discard semantics are |
| // defined for WGSL (https://github.com/gpuweb/gpuweb/issues/361) |
| line() << "discard_fragment();"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitIf(const ast::IfStatement* stmt) { |
| { |
| auto out = line(); |
| out << "if ("; |
| if (!EmitExpression(out, stmt->condition)) { |
| return false; |
| } |
| out << ") {"; |
| } |
| |
| if (!EmitStatementsWithIndent(stmt->body->statements)) { |
| return false; |
| } |
| |
| for (auto* e : stmt->else_statements) { |
| if (e->condition) { |
| line() << "} else {"; |
| increment_indent(); |
| |
| { |
| auto out = line(); |
| out << "if ("; |
| if (!EmitExpression(out, e->condition)) { |
| return false; |
| } |
| out << ") {"; |
| } |
| } else { |
| line() << "} else {"; |
| } |
| |
| if (!EmitStatementsWithIndent(e->body->statements)) { |
| return false; |
| } |
| } |
| |
| line() << "}"; |
| |
| for (auto* e : stmt->else_statements) { |
| if (e->condition) { |
| decrement_indent(); |
| line() << "}"; |
| } |
| } |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitMemberAccessor( |
| std::ostream& out, |
| const ast::MemberAccessorExpression* expr) { |
| auto write_lhs = [&] { |
| bool paren_lhs = !expr->structure->IsAnyOf< |
| ast::IndexAccessorExpression, ast::CallExpression, |
| ast::IdentifierExpression, ast::MemberAccessorExpression>(); |
| if (paren_lhs) { |
| out << "("; |
| } |
| if (!EmitExpression(out, expr->structure)) { |
| return false; |
| } |
| if (paren_lhs) { |
| out << ")"; |
| } |
| return true; |
| }; |
| |
| auto& sem = program_->Sem(); |
| |
| if (auto* swizzle = sem.Get(expr)->As<sem::Swizzle>()) { |
| // Metal 1.x does not support swizzling of packed vector types. |
| // For single element swizzles, we can use the index operator. |
| // For multi-element swizzles, we need to cast to a regular vector type |
| // first. Note that we do not currently allow assignments to swizzles, so |
| // the casting which will convert the l-value to r-value is fine. |
| if (swizzle->Indices().size() == 1) { |
| if (!write_lhs()) { |
| return false; |
| } |
| out << "[" << swizzle->Indices()[0] << "]"; |
| } else { |
| if (!EmitType(out, sem.Get(expr->structure)->Type()->UnwrapRef(), "")) { |
| return false; |
| } |
| out << "("; |
| if (!write_lhs()) { |
| return false; |
| } |
| out << ")." << program_->Symbols().NameFor(expr->member->symbol); |
| } |
| } else { |
| if (!write_lhs()) { |
| return false; |
| } |
| out << "."; |
| if (!EmitExpression(out, expr->member)) { |
| return false; |
| } |
| } |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitReturn(const ast::ReturnStatement* stmt) { |
| auto out = line(); |
| out << "return"; |
| if (stmt->value) { |
| out << " "; |
| if (!EmitExpression(out, stmt->value)) { |
| return false; |
| } |
| } |
| out << ";"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitBlock(const ast::BlockStatement* stmt) { |
| line() << "{"; |
| |
| if (!EmitStatementsWithIndent(stmt->statements)) { |
| return false; |
| } |
| |
| line() << "}"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitStatement(const ast::Statement* stmt) { |
| if (auto* a = stmt->As<ast::AssignmentStatement>()) { |
| return EmitAssign(a); |
| } |
| if (auto* b = stmt->As<ast::BlockStatement>()) { |
| return EmitBlock(b); |
| } |
| if (auto* b = stmt->As<ast::BreakStatement>()) { |
| return EmitBreak(b); |
| } |
| if (auto* c = stmt->As<ast::CallStatement>()) { |
| auto out = line(); |
| if (!EmitCall(out, c->expr)) { |
| return false; |
| } |
| out << ";"; |
| return true; |
| } |
| if (auto* c = stmt->As<ast::ContinueStatement>()) { |
| return EmitContinue(c); |
| } |
| if (auto* d = stmt->As<ast::DiscardStatement>()) { |
| return EmitDiscard(d); |
| } |
| if (stmt->As<ast::FallthroughStatement>()) { |
| line() << "/* fallthrough */"; |
| return true; |
| } |
| if (auto* i = stmt->As<ast::IfStatement>()) { |
| return EmitIf(i); |
| } |
| if (auto* l = stmt->As<ast::LoopStatement>()) { |
| return EmitLoop(l); |
| } |
| if (auto* l = stmt->As<ast::ForLoopStatement>()) { |
| return EmitForLoop(l); |
| } |
| if (auto* r = stmt->As<ast::ReturnStatement>()) { |
| return EmitReturn(r); |
| } |
| if (auto* s = stmt->As<ast::SwitchStatement>()) { |
| return EmitSwitch(s); |
| } |
| if (auto* v = stmt->As<ast::VariableDeclStatement>()) { |
| auto* var = program_->Sem().Get(v->variable); |
| return EmitVariable(var); |
| } |
| |
| diagnostics_.add_error( |
| diag::System::Writer, |
| "unknown statement type: " + std::string(stmt->TypeInfo().name)); |
| return false; |
| } |
| |
| bool GeneratorImpl::EmitStatements(const ast::StatementList& stmts) { |
| for (auto* s : stmts) { |
| if (!EmitStatement(s)) { |
| return false; |
| } |
| } |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitStatementsWithIndent(const ast::StatementList& stmts) { |
| ScopedIndent si(this); |
| return EmitStatements(stmts); |
| } |
| |
| bool GeneratorImpl::EmitSwitch(const ast::SwitchStatement* stmt) { |
| { |
| auto out = line(); |
| out << "switch("; |
| if (!EmitExpression(out, stmt->condition)) { |
| return false; |
| } |
| out << ") {"; |
| } |
| |
| { |
| ScopedIndent si(this); |
| for (auto* s : stmt->body) { |
| if (!EmitCase(s)) { |
| return false; |
| } |
| } |
| } |
| |
| line() << "}"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitType(std::ostream& out, |
| const sem::Type* type, |
| const std::string& name, |
| bool* name_printed /* = nullptr */) { |
| if (name_printed) { |
| *name_printed = false; |
| } |
| if (auto* atomic = type->As<sem::Atomic>()) { |
| if (atomic->Type()->Is<sem::I32>()) { |
| out << "atomic_int"; |
| return true; |
| } |
| if (atomic->Type()->Is<sem::U32>()) { |
| out << "atomic_uint"; |
| return true; |
| } |
| TINT_ICE(Writer, diagnostics_) |
| << "unhandled atomic type " << atomic->Type()->type_name(); |
| return false; |
| } |
| |
| if (auto* ary = type->As<sem::Array>()) { |
| const sem::Type* base_type = ary; |
| std::vector<uint32_t> sizes; |
| while (auto* arr = base_type->As<sem::Array>()) { |
| if (arr->IsRuntimeSized()) { |
| sizes.push_back(1); |
| } else { |
| sizes.push_back(arr->Count()); |
| } |
| base_type = arr->ElemType(); |
| } |
| if (!EmitType(out, base_type, "")) { |
| return false; |
| } |
| if (!name.empty()) { |
| out << " " << name; |
| if (name_printed) { |
| *name_printed = true; |
| } |
| } |
| for (uint32_t size : sizes) { |
| out << "[" << size << "]"; |
| } |
| return true; |
| } |
| |
| if (type->Is<sem::Bool>()) { |
| out << "bool"; |
| return true; |
| } |
| |
| if (type->Is<sem::F32>()) { |
| out << "float"; |
| return true; |
| } |
| |
| if (type->Is<sem::I32>()) { |
| out << "int"; |
| return true; |
| } |
| |
| if (auto* mat = type->As<sem::Matrix>()) { |
| if (!EmitType(out, mat->type(), "")) { |
| return false; |
| } |
| out << mat->columns() << "x" << mat->rows(); |
| return true; |
| } |
| |
| if (auto* ptr = type->As<sem::Pointer>()) { |
| if (ptr->Access() == ast::Access::kRead) { |
| out << "const "; |
| } |
| if (!EmitStorageClass(out, ptr->StorageClass())) { |
| return false; |
| } |
| out << " "; |
| if (ptr->StoreType()->Is<sem::Array>()) { |
| std::string inner = "(*" + name + ")"; |
| if (!EmitType(out, ptr->StoreType(), inner)) { |
| return false; |
| } |
| if (name_printed) { |
| *name_printed = true; |
| } |
| } else { |
| if (!EmitType(out, ptr->StoreType(), "")) { |
| return false; |
| } |
| out << "* " << name; |
| if (name_printed) { |
| *name_printed = true; |
| } |
| } |
| return true; |
| } |
| |
| if (type->Is<sem::Sampler>()) { |
| out << "sampler"; |
| return true; |
| } |
| |
| if (auto* str = type->As<sem::Struct>()) { |
| // The struct type emits as just the name. The declaration would be emitted |
| // as part of emitting the declared types. |
| out << StructName(str); |
| return true; |
| } |
| |
| if (auto* tex = type->As<sem::Texture>()) { |
| if (tex->IsAnyOf<sem::DepthTexture, sem::DepthMultisampledTexture>()) { |
| out << "depth"; |
| } else { |
| out << "texture"; |
| } |
| |
| switch (tex->dim()) { |
| case ast::TextureDimension::k1d: |
| out << "1d"; |
| break; |
| case ast::TextureDimension::k2d: |
| out << "2d"; |
| break; |
| case ast::TextureDimension::k2dArray: |
| out << "2d_array"; |
| break; |
| case ast::TextureDimension::k3d: |
| out << "3d"; |
| break; |
| case ast::TextureDimension::kCube: |
| out << "cube"; |
| break; |
| case ast::TextureDimension::kCubeArray: |
| out << "cube_array"; |
| break; |
| default: |
| diagnostics_.add_error(diag::System::Writer, |
| "Invalid texture dimensions"); |
| return false; |
| } |
| if (tex->IsAnyOf<sem::MultisampledTexture, |
| sem::DepthMultisampledTexture>()) { |
| out << "_ms"; |
| } |
| out << "<"; |
| if (tex->Is<sem::DepthTexture>()) { |
| out << "float, access::sample"; |
| } else if (tex->Is<sem::DepthMultisampledTexture>()) { |
| out << "float, access::read"; |
| } else if (auto* storage = tex->As<sem::StorageTexture>()) { |
| if (!EmitType(out, storage->type(), "")) { |
| return false; |
| } |
| |
| std::string access_str; |
| if (storage->access() == ast::Access::kRead) { |
| out << ", access::read"; |
| } else if (storage->access() == ast::Access::kWrite) { |
| out << ", access::write"; |
| } else { |
| diagnostics_.add_error(diag::System::Writer, |
| "Invalid access control for storage texture"); |
| return false; |
| } |
| } else if (auto* ms = tex->As<sem::MultisampledTexture>()) { |
| if (!EmitType(out, ms->type(), "")) { |
| return false; |
| } |
| out << ", access::read"; |
| } else if (auto* sampled = tex->As<sem::SampledTexture>()) { |
| if (!EmitType(out, sampled->type(), "")) { |
| return false; |
| } |
| out << ", access::sample"; |
| } else { |
| diagnostics_.add_error(diag::System::Writer, "invalid texture type"); |
| return false; |
| } |
| out << ">"; |
| return true; |
| } |
| |
| if (type->Is<sem::U32>()) { |
| out << "uint"; |
| return true; |
| } |
| |
| if (auto* vec = type->As<sem::Vector>()) { |
| if (!EmitType(out, vec->type(), "")) { |
| return false; |
| } |
| out << vec->Width(); |
| return true; |
| } |
| |
| if (type->Is<sem::Void>()) { |
| out << "void"; |
| return true; |
| } |
| |
| diagnostics_.add_error(diag::System::Writer, |
| "unknown type in EmitType: " + type->type_name()); |
| return false; |
| } |
| |
| bool GeneratorImpl::EmitTypeAndName(std::ostream& out, |
| const sem::Type* type, |
| const std::string& name) { |
| bool name_printed = false; |
| if (!EmitType(out, type, name, &name_printed)) { |
| return false; |
| } |
| if (!name_printed) { |
| out << " " << name; |
| } |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitStorageClass(std::ostream& out, ast::StorageClass sc) { |
| switch (sc) { |
| case ast::StorageClass::kFunction: |
| case ast::StorageClass::kPrivate: |
| case ast::StorageClass::kUniformConstant: |
| out << "thread"; |
| return true; |
| case ast::StorageClass::kWorkgroup: |
| out << "threadgroup"; |
| return true; |
| case ast::StorageClass::kStorage: |
| out << "device"; |
| return true; |
| case ast::StorageClass::kUniform: |
| out << "constant"; |
| return true; |
| default: |
| break; |
| } |
| TINT_ICE(Writer, diagnostics_) << "unhandled storage class: " << sc; |
| return false; |
| } |
| |
| bool GeneratorImpl::EmitPackedType(std::ostream& out, |
| const sem::Type* type, |
| const std::string& name) { |
| auto* vec = type->As<sem::Vector>(); |
| if (vec && vec->Width() == 3) { |
| out << "packed_"; |
| if (!EmitType(out, vec, "")) { |
| return false; |
| } |
| |
| if (vec->is_float_vector() && !matrix_packed_vector_overloads_) { |
| // Overload operators for matrix-vector arithmetic where the vector |
| // operand is packed, as these overloads to not exist in the metal |
| // namespace. |
| TextBuffer b; |
| TINT_DEFER(helpers_.Append(b)); |
| line(&b) << R"(template<typename T, int N, int M> |
| inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) { |
| return lhs * vec<T, N>(rhs); |
| } |
| |
| template<typename T, int N, int M> |
| inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) { |
| return vec<T, M>(lhs) * rhs; |
| } |
| )"; |
| matrix_packed_vector_overloads_ = true; |
| } |
| |
| return true; |
| } |
| |
| return EmitType(out, type, name); |
| } |
| |
| bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { |
| line(b) << "struct " << StructName(str) << " {"; |
| |
| bool is_host_shareable = str->IsHostShareable(); |
| |
| // Emits a `/* 0xnnnn */` byte offset comment for a struct member. |
| auto add_byte_offset_comment = [&](std::ostream& out, uint32_t offset) { |
| std::ios_base::fmtflags saved_flag_state(out.flags()); |
| out << "/* 0x" << std::hex << std::setfill('0') << std::setw(4) << offset |
| << " */ "; |
| out.flags(saved_flag_state); |
| }; |
| |
| auto add_padding = [&](uint32_t size, uint32_t msl_offset) { |
| std::string name; |
| do { |
| name = UniqueIdentifier("tint_pad"); |
| } while (str->FindMember(program_->Symbols().Get(name))); |
| |
| auto out = line(b); |
| add_byte_offset_comment(out, msl_offset); |
| out << "int8_t " << name << "[" << size << "];"; |
| }; |
| |
| b->IncrementIndent(); |
| |
| uint32_t msl_offset = 0; |
| for (auto* mem : str->Members()) { |
| auto out = line(b); |
| auto name = program_->Symbols().NameFor(mem->Name()); |
| auto wgsl_offset = mem->Offset(); |
| |
| if (is_host_shareable) { |
| if (wgsl_offset < msl_offset) { |
| // Unimplementable layout |
| TINT_ICE(Writer, diagnostics_) |
| << "Structure member WGSL offset (" << wgsl_offset |
| << ") is behind MSL offset (" << msl_offset << ")"; |
| return false; |
| } |
| |
| // Generate padding if required |
| if (auto padding = wgsl_offset - msl_offset) { |
| add_padding(padding, msl_offset); |
| msl_offset += padding; |
| } |
| |
| add_byte_offset_comment(out, msl_offset); |
| |
| if (!EmitPackedType(out, mem->Type(), name)) { |
| return false; |
| } |
| } else { |
| if (!EmitType(out, mem->Type(), name)) { |
| return false; |
| } |
| } |
| |
| auto* ty = mem->Type(); |
| |
| // Array member name will be output with the type |
| if (!ty->Is<sem::Array>()) { |
| out << " " << name; |
| } |
| |
| // Emit decorations |
| if (auto* decl = mem->Declaration()) { |
| for (auto* deco : decl->decorations) { |
| if (auto* builtin = deco->As<ast::BuiltinDecoration>()) { |
| auto attr = builtin_to_attribute(builtin->builtin); |
| if (attr.empty()) { |
| diagnostics_.add_error(diag::System::Writer, "unknown builtin"); |
| return false; |
| } |
| out << " [[" << attr << "]]"; |
| } else if (auto* loc = deco->As<ast::LocationDecoration>()) { |
| auto& pipeline_stage_uses = str->PipelineStageUses(); |
| if (pipeline_stage_uses.size() != 1) { |
| TINT_ICE(Writer, diagnostics_) |
| << "invalid entry point IO struct uses"; |
| } |
| |
| if (pipeline_stage_uses.count( |
| sem::PipelineStageUsage::kVertexInput)) { |
| out << " [[attribute(" + std::to_string(loc->value) + ")]]"; |
| } else if (pipeline_stage_uses.count( |
| sem::PipelineStageUsage::kVertexOutput)) { |
| out << " [[user(locn" + std::to_string(loc->value) + ")]]"; |
| } else if (pipeline_stage_uses.count( |
| sem::PipelineStageUsage::kFragmentInput)) { |
| out << " [[user(locn" + std::to_string(loc->value) + ")]]"; |
| } else if (pipeline_stage_uses.count( |
| sem::PipelineStageUsage::kFragmentOutput)) { |
| out << " [[color(" + std::to_string(loc->value) + ")]]"; |
| } else { |
| TINT_ICE(Writer, diagnostics_) |
| << "invalid use of location decoration"; |
| } |
| } else if (auto* interpolate = deco->As<ast::InterpolateDecoration>()) { |
| auto attr = interpolation_to_attribute(interpolate->type, |
| interpolate->sampling); |
| if (attr.empty()) { |
| diagnostics_.add_error(diag::System::Writer, |
| "unknown interpolation attribute"); |
| return false; |
| } |
| out << " [[" << attr << "]]"; |
| } else if (deco->Is<ast::InvariantDecoration>()) { |
| out << " [[invariant]]"; |
| has_invariant_ = true; |
| } else if (!deco->IsAnyOf<ast::StructMemberOffsetDecoration, |
| ast::StructMemberAlignDecoration, |
| ast::StructMemberSizeDecoration>()) { |
| TINT_ICE(Writer, diagnostics_) |
| << "unhandled struct member attribute: " << deco->Name(); |
| } |
| } |
| } |
| |
| out << ";"; |
| |
| if (is_host_shareable) { |
| // Calculate new MSL offset |
| auto size_align = MslPackedTypeSizeAndAlign(ty); |
| if (msl_offset % size_align.align) { |
| TINT_ICE(Writer, diagnostics_) |
| << "Misaligned MSL structure member " |
| << ty->FriendlyName(program_->Symbols()) << " " << name; |
| return false; |
| } |
| msl_offset += size_align.size; |
| } |
| } |
| |
| if (is_host_shareable && str->Size() != msl_offset) { |
| add_padding(str->Size() - msl_offset, msl_offset); |
| } |
| |
| b->DecrementIndent(); |
| |
| line(b) << "};"; |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitUnaryOp(std::ostream& out, |
| const ast::UnaryOpExpression* expr) { |
| // Handle `-e` when `e` is signed, so that we ensure that if `e` is the |
| // largest negative value, it returns `e`. |
| auto* expr_type = TypeOf(expr->expr)->UnwrapRef(); |
| if (expr->op == ast::UnaryOp::kNegation && |
| expr_type->is_signed_scalar_or_vector()) { |
| auto fn = |
| utils::GetOrCreate(unary_minus_funcs_, expr_type, [&]() -> std::string { |
| // e.g.: |
| // int tint_unary_minus(const int v) { |
| // return (v == -2147483648) ? v : -v; |
| // } |
| TextBuffer b; |
| TINT_DEFER(helpers_.Append(b)); |
| |
| auto fn_name = UniqueIdentifier("tint_unary_minus"); |
| { |
| auto decl = line(&b); |
| if (!EmitTypeAndName(decl, expr_type, fn_name)) { |
| return ""; |
| } |
| decl << "(const "; |
| if (!EmitType(decl, expr_type, "")) { |
| return ""; |
| } |
| decl << " v) {"; |
| } |
| |
| { |
| ScopedIndent si(&b); |
| const auto largest_negative_value = |
| std::to_string(std::numeric_limits<int32_t>::min()); |
| line(&b) << "return select(-v, v, v == " << largest_negative_value |
| << ");"; |
| } |
| line(&b) << "}"; |
| line(&b); |
| return fn_name; |
| }); |
| |
| out << fn << "("; |
| if (!EmitExpression(out, expr->expr)) { |
| return false; |
| } |
| out << ")"; |
| return true; |
| } |
| |
| switch (expr->op) { |
| case ast::UnaryOp::kAddressOf: |
| out << "&"; |
| break; |
| case ast::UnaryOp::kComplement: |
| out << "~"; |
| break; |
| case ast::UnaryOp::kIndirection: |
| out << "*"; |
| break; |
| case ast::UnaryOp::kNot: |
| out << "!"; |
| break; |
| case ast::UnaryOp::kNegation: |
| out << "-"; |
| break; |
| } |
| out << "("; |
| |
| if (!EmitExpression(out, expr->expr)) { |
| return false; |
| } |
| |
| out << ")"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitVariable(const sem::Variable* var) { |
| auto* decl = var->Declaration(); |
| |
| for (auto* deco : decl->decorations) { |
| if (!deco->Is<ast::InternalDecoration>()) { |
| TINT_ICE(Writer, diagnostics_) << "unexpected variable decoration"; |
| return false; |
| } |
| } |
| |
| auto out = line(); |
| |
| switch (var->StorageClass()) { |
| case ast::StorageClass::kFunction: |
| case ast::StorageClass::kUniformConstant: |
| case ast::StorageClass::kNone: |
| break; |
| case ast::StorageClass::kPrivate: |
| out << "thread "; |
| break; |
| case ast::StorageClass::kWorkgroup: |
| out << "threadgroup "; |
| break; |
| default: |
| TINT_ICE(Writer, diagnostics_) << "unhandled variable storage class"; |
| return false; |
| } |
| |
| auto* type = var->Type()->UnwrapRef(); |
| |
| std::string name = program_->Symbols().NameFor(decl->symbol); |
| if (decl->is_const) { |
| name = "const " + name; |
| } |
| if (!EmitType(out, type, name)) { |
| return false; |
| } |
| // Variable name is output as part of the type for arrays and pointers. |
| if (!type->Is<sem::Array>() && !type->Is<sem::Pointer>()) { |
| out << " " << name; |
| } |
| |
| if (decl->constructor != nullptr) { |
| out << " = "; |
| if (!EmitExpression(out, decl->constructor)) { |
| return false; |
| } |
| } else if (var->StorageClass() == ast::StorageClass::kPrivate || |
| var->StorageClass() == ast::StorageClass::kFunction || |
| var->StorageClass() == ast::StorageClass::kNone) { |
| out << " = "; |
| if (!EmitZeroValue(out, type)) { |
| return false; |
| } |
| } |
| out << ";"; |
| |
| return true; |
| } |
| |
| bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) { |
| for (auto* d : var->decorations) { |
| if (!d->Is<ast::OverrideDecoration>()) { |
| diagnostics_.add_error(diag::System::Writer, |
| "Decorated const values not valid"); |
| return false; |
| } |
| } |
| if (!var->is_const) { |
| diagnostics_.add_error(diag::System::Writer, "Expected a const value"); |
| return false; |
| } |
| |
| auto out = line(); |
| out << "constant "; |
| auto* type = program_->Sem().Get(var)->Type()->UnwrapRef(); |
| if (!EmitType(out, type, program_->Symbols().NameFor(var->symbol))) { |
| return false; |
| } |
| if (!type->Is<sem::Array>()) { |
| out << " " << program_->Symbols().NameFor(var->symbol); |
| } |
| |
| auto* global = program_->Sem().Get<sem::GlobalVariable>(var); |
| if (global && global->IsOverridable()) { |
| out << " [[function_constant(" << global->ConstantId() << ")]]"; |
| } else if (var->constructor != nullptr) { |
| out << " = "; |
| if (!EmitExpression(out, var->constructor)) { |
| return false; |
| } |
| } |
| out << ";"; |
| |
| return true; |
| } |
| |
| GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign( |
| const sem::Type* ty) { |
| if (ty->IsAnyOf<sem::U32, sem::I32, sem::F32>()) { |
| // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf |
| // 2.1 Scalar Data Types |
| return {4, 4}; |
| } |
| |
| if (auto* vec = ty->As<sem::Vector>()) { |
| auto num_els = vec->Width(); |
| auto* el_ty = vec->type(); |
| if (el_ty->IsAnyOf<sem::U32, sem::I32, sem::F32>()) { |
| // Use a packed_vec type for 3-element vectors only. |
| if (num_els == 3) { |
| // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf |
| // 2.2.3 Packed Vector Types |
| return SizeAndAlign{num_els * 4, 4}; |
| } else { |
| // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf |
| // 2.2 Vector Data Types |
| return SizeAndAlign{num_els * 4, num_els * 4}; |
| } |
| } |
| } |
| |
| if (auto* mat = ty->As<sem::Matrix>()) { |
| // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf |
| // 2.3 Matrix Data Types |
| auto cols = mat->columns(); |
| auto rows = mat->rows(); |
| auto* el_ty = mat->type(); |
| if (el_ty->IsAnyOf<sem::U32, sem::I32, sem::F32>()) { |
| static constexpr SizeAndAlign table[] = { |
| /* float2x2 */ {16, 8}, |
| /* float2x3 */ {32, 16}, |
| /* float2x4 */ {32, 16}, |
| /* float3x2 */ {24, 8}, |
| /* float3x3 */ {48, 16}, |
| /* float3x4 */ {48, 16}, |
| /* float4x2 */ {32, 8}, |
| /* float4x3 */ {64, 16}, |
| /* float4x4 */ {64, 16}, |
| }; |
| if (cols >= 2 && cols <= 4 && rows >= 2 && rows <= 4) { |
| return table[(3 * (cols - 2)) + (rows - 2)]; |
| } |
| } |
| } |
| |
| if (auto* arr = ty->As<sem::Array>()) { |
| if (!arr->IsStrideImplicit()) { |
| TINT_ICE(Writer, diagnostics_) |
| << "arrays with explicit strides should have " |
| "removed with the PadArrayElements transform"; |
| return {}; |
| } |
| auto num_els = std::max<uint32_t>(arr->Count(), 1); |
| return SizeAndAlign{arr->Stride() * num_els, arr->Align()}; |
| } |
| |
| if (auto* str = ty->As<sem::Struct>()) { |
| // TODO(crbug.com/tint/650): There's an assumption here that MSL's default |
| // structure size and alignment matches WGSL's. We need to confirm this. |
| return SizeAndAlign{str->Size(), str->Align()}; |
| } |
| |
| if (auto* atomic = ty->As<sem::Atomic>()) { |
| return MslPackedTypeSizeAndAlign(atomic->Type()); |
| } |
| |
| TINT_UNREACHABLE(Writer, diagnostics_) |
| << "Unhandled type " << ty->TypeInfo().name; |
| return {}; |
| } |
| |
| template <typename F> |
| bool GeneratorImpl::CallIntrinsicHelper(std::ostream& out, |
| const ast::CallExpression* call, |
| const sem::Intrinsic* intrinsic, |
| F&& build) { |
| // Generate the helper function if it hasn't been created already |
| auto fn = utils::GetOrCreate(intrinsics_, intrinsic, [&]() -> std::string { |
| TextBuffer b; |
| TINT_DEFER(helpers_.Append(b)); |
| |
| auto fn_name = |
| UniqueIdentifier(std::string("tint_") + sem::str(intrinsic->Type())); |
| std::vector<std::string> parameter_names; |
| { |
| auto decl = line(&b); |
| if (!EmitTypeAndName(decl, intrinsic->ReturnType(), fn_name)) { |
| return ""; |
| } |
| { |
| ScopedParen sp(decl); |
| for (auto* param : intrinsic->Parameters()) { |
| if (!parameter_names.empty()) { |
| decl << ", "; |
| } |
| auto param_name = "param_" + std::to_string(parameter_names.size()); |
| if (!EmitTypeAndName(decl, param->Type(), param_name)) { |
| return ""; |
| } |
| parameter_names.emplace_back(std::move(param_name)); |
| } |
| } |
| decl << " {"; |
| } |
| { |
| ScopedIndent si(&b); |
| if (!build(&b, parameter_names)) { |
| return ""; |
| } |
| } |
| line(&b) << "}"; |
| line(&b); |
| return fn_name; |
| }); |
| |
| if (fn.empty()) { |
| return false; |
| } |
| |
| // Call the helper |
| out << fn; |
| { |
| ScopedParen sp(out); |
| bool first = true; |
| for (auto* arg : call->args) { |
| if (!first) { |
| out << ", "; |
| } |
| first = false; |
| if (!EmitExpression(out, arg)) { |
| return false; |
| } |
| } |
| } |
| return true; |
| } |
| |
| } // namespace msl |
| } // namespace writer |
| } // namespace tint |
