// 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/tint/writer/spirv/builder.h"

#include <algorithm>
#include <limits>
#include <utility>

#include "spirv/unified1/GLSL.std.450.h"
#include "src/tint/ast/call_statement.h"
#include "src/tint/ast/fallthrough_statement.h"
#include "src/tint/ast/id_attribute.h"
#include "src/tint/ast/internal_attribute.h"
#include "src/tint/ast/traverse_expressions.h"
#include "src/tint/sem/array.h"
#include "src/tint/sem/atomic_type.h"
#include "src/tint/sem/builtin.h"
#include "src/tint/sem/call.h"
#include "src/tint/sem/depth_multisampled_texture_type.h"
#include "src/tint/sem/depth_texture_type.h"
#include "src/tint/sem/function.h"
#include "src/tint/sem/member_accessor_expression.h"
#include "src/tint/sem/module.h"
#include "src/tint/sem/multisampled_texture_type.h"
#include "src/tint/sem/reference_type.h"
#include "src/tint/sem/sampled_texture_type.h"
#include "src/tint/sem/statement.h"
#include "src/tint/sem/struct.h"
#include "src/tint/sem/type_constructor.h"
#include "src/tint/sem/type_conversion.h"
#include "src/tint/sem/variable.h"
#include "src/tint/sem/vector_type.h"
#include "src/tint/transform/add_spirv_block_attribute.h"
#include "src/tint/utils/defer.h"
#include "src/tint/utils/map.h"
#include "src/tint/writer/append_vector.h"

namespace tint::writer::spirv {
namespace {

using BuiltinType = sem::BuiltinType;

const char kGLSLstd450[] = "GLSL.std.450";

uint32_t size_of(const InstructionList& instructions) {
  uint32_t size = 0;
  for (const auto& inst : instructions)
    size += inst.word_length();

  return size;
}

uint32_t pipeline_stage_to_execution_model(ast::PipelineStage stage) {
  SpvExecutionModel model = SpvExecutionModelVertex;

  switch (stage) {
    case ast::PipelineStage::kFragment:
      model = SpvExecutionModelFragment;
      break;
    case ast::PipelineStage::kVertex:
      model = SpvExecutionModelVertex;
      break;
    case ast::PipelineStage::kCompute:
      model = SpvExecutionModelGLCompute;
      break;
    case ast::PipelineStage::kNone:
      model = SpvExecutionModelMax;
      break;
  }
  return model;
}

bool LastIsFallthrough(const ast::BlockStatement* stmts) {
  return !stmts->Empty() && stmts->Last()->Is<ast::FallthroughStatement>();
}

/// Returns the matrix type that is `type` or that is wrapped by
/// one or more levels of an arrays inside of `type`.
/// @param type the given type, which must not be null
/// @returns the nested matrix type, or nullptr if none
const sem::Matrix* GetNestedMatrixType(const sem::Type* type) {
  while (auto* arr = type->As<sem::Array>()) {
    type = arr->ElemType();
  }
  return type->As<sem::Matrix>();
}

uint32_t builtin_to_glsl_method(const sem::Builtin* builtin) {
  switch (builtin->Type()) {
    case BuiltinType::kAcos:
      return GLSLstd450Acos;
    case BuiltinType::kAsin:
      return GLSLstd450Asin;
    case BuiltinType::kAtan:
      return GLSLstd450Atan;
    case BuiltinType::kAtan2:
      return GLSLstd450Atan2;
    case BuiltinType::kCeil:
      return GLSLstd450Ceil;
    case BuiltinType::kClamp:
      if (builtin->ReturnType()->is_float_scalar_or_vector()) {
        return GLSLstd450NClamp;
      } else if (builtin->ReturnType()->is_unsigned_scalar_or_vector()) {
        return GLSLstd450UClamp;
      } else {
        return GLSLstd450SClamp;
      }
    case BuiltinType::kCos:
      return GLSLstd450Cos;
    case BuiltinType::kCosh:
      return GLSLstd450Cosh;
    case BuiltinType::kCross:
      return GLSLstd450Cross;
    case BuiltinType::kDegrees:
      return GLSLstd450Degrees;
    case BuiltinType::kDeterminant:
      return GLSLstd450Determinant;
    case BuiltinType::kDistance:
      return GLSLstd450Distance;
    case BuiltinType::kExp:
      return GLSLstd450Exp;
    case BuiltinType::kExp2:
      return GLSLstd450Exp2;
    case BuiltinType::kFaceForward:
      return GLSLstd450FaceForward;
    case BuiltinType::kFloor:
      return GLSLstd450Floor;
    case BuiltinType::kFma:
      return GLSLstd450Fma;
    case BuiltinType::kFract:
      return GLSLstd450Fract;
    case BuiltinType::kFrexp:
      return GLSLstd450FrexpStruct;
    case BuiltinType::kInverseSqrt:
      return GLSLstd450InverseSqrt;
    case BuiltinType::kLdexp:
      return GLSLstd450Ldexp;
    case BuiltinType::kLength:
      return GLSLstd450Length;
    case BuiltinType::kLog:
      return GLSLstd450Log;
    case BuiltinType::kLog2:
      return GLSLstd450Log2;
    case BuiltinType::kMax:
      if (builtin->ReturnType()->is_float_scalar_or_vector()) {
        return GLSLstd450NMax;
      } else if (builtin->ReturnType()->is_unsigned_scalar_or_vector()) {
        return GLSLstd450UMax;
      } else {
        return GLSLstd450SMax;
      }
    case BuiltinType::kMin:
      if (builtin->ReturnType()->is_float_scalar_or_vector()) {
        return GLSLstd450NMin;
      } else if (builtin->ReturnType()->is_unsigned_scalar_or_vector()) {
        return GLSLstd450UMin;
      } else {
        return GLSLstd450SMin;
      }
    case BuiltinType::kMix:
      return GLSLstd450FMix;
    case BuiltinType::kModf:
      return GLSLstd450ModfStruct;
    case BuiltinType::kNormalize:
      return GLSLstd450Normalize;
    case BuiltinType::kPack4x8snorm:
      return GLSLstd450PackSnorm4x8;
    case BuiltinType::kPack4x8unorm:
      return GLSLstd450PackUnorm4x8;
    case BuiltinType::kPack2x16snorm:
      return GLSLstd450PackSnorm2x16;
    case BuiltinType::kPack2x16unorm:
      return GLSLstd450PackUnorm2x16;
    case BuiltinType::kPack2x16float:
      return GLSLstd450PackHalf2x16;
    case BuiltinType::kPow:
      return GLSLstd450Pow;
    case BuiltinType::kRadians:
      return GLSLstd450Radians;
    case BuiltinType::kReflect:
      return GLSLstd450Reflect;
    case BuiltinType::kRefract:
      return GLSLstd450Refract;
    case BuiltinType::kRound:
      return GLSLstd450RoundEven;
    case BuiltinType::kSign:
      return GLSLstd450FSign;
    case BuiltinType::kSin:
      return GLSLstd450Sin;
    case BuiltinType::kSinh:
      return GLSLstd450Sinh;
    case BuiltinType::kSmoothstep:
    case BuiltinType::kSmoothStep:
      return GLSLstd450SmoothStep;
    case BuiltinType::kSqrt:
      return GLSLstd450Sqrt;
    case BuiltinType::kStep:
      return GLSLstd450Step;
    case BuiltinType::kTan:
      return GLSLstd450Tan;
    case BuiltinType::kTanh:
      return GLSLstd450Tanh;
    case BuiltinType::kTrunc:
      return GLSLstd450Trunc;
    case BuiltinType::kUnpack4x8snorm:
      return GLSLstd450UnpackSnorm4x8;
    case BuiltinType::kUnpack4x8unorm:
      return GLSLstd450UnpackUnorm4x8;
    case BuiltinType::kUnpack2x16snorm:
      return GLSLstd450UnpackSnorm2x16;
    case BuiltinType::kUnpack2x16unorm:
      return GLSLstd450UnpackUnorm2x16;
    case BuiltinType::kUnpack2x16float:
      return GLSLstd450UnpackHalf2x16;
    default:
      break;
  }
  return 0;
}

/// @return the vector element type if ty is a vector, otherwise return ty.
const sem::Type* ElementTypeOf(const sem::Type* ty) {
  if (auto* v = ty->As<sem::Vector>()) {
    return v->type();
  }
  return ty;
}

}  // namespace

Builder::AccessorInfo::AccessorInfo() : source_id(0), source_type(nullptr) {}

Builder::AccessorInfo::~AccessorInfo() {}

Builder::Builder(const Program* program, bool zero_initialize_workgroup_memory)
    : builder_(ProgramBuilder::Wrap(program)),
      scope_stack_({}),
      zero_initialize_workgroup_memory_(zero_initialize_workgroup_memory) {}

Builder::~Builder() = default;

bool Builder::Build() {
  push_capability(SpvCapabilityShader);

  push_memory_model(spv::Op::OpMemoryModel,
                    {Operand::Int(SpvAddressingModelLogical),
                     Operand::Int(SpvMemoryModelGLSL450)});

  for (auto* var : builder_.AST().GlobalVariables()) {
    if (!GenerateGlobalVariable(var)) {
      return false;
    }
  }

  auto* mod = builder_.Sem().Module();
  for (auto* decl : mod->DependencyOrderedDeclarations()) {
    if (auto* func = decl->As<ast::Function>()) {
      if (!GenerateFunction(func)) {
        return false;
      }
    }
  }

  return true;
}

Operand Builder::result_op() {
  return Operand::Int(next_id());
}

uint32_t Builder::total_size() const {
  // The 5 covers the magic, version, generator, id bound and reserved.
  uint32_t size = 5;

  size += size_of(capabilities_);
  size += size_of(extensions_);
  size += size_of(ext_imports_);
  size += size_of(memory_model_);
  size += size_of(entry_points_);
  size += size_of(execution_modes_);
  size += size_of(debug_);
  size += size_of(annotations_);
  size += size_of(types_);
  for (const auto& func : functions_) {
    size += func.word_length();
  }

  return size;
}

void Builder::iterate(std::function<void(const Instruction&)> cb) const {
  for (const auto& inst : capabilities_) {
    cb(inst);
  }
  for (const auto& inst : extensions_) {
    cb(inst);
  }
  for (const auto& inst : ext_imports_) {
    cb(inst);
  }
  for (const auto& inst : memory_model_) {
    cb(inst);
  }
  for (const auto& inst : entry_points_) {
    cb(inst);
  }
  for (const auto& inst : execution_modes_) {
    cb(inst);
  }
  for (const auto& inst : debug_) {
    cb(inst);
  }
  for (const auto& inst : annotations_) {
    cb(inst);
  }
  for (const auto& inst : types_) {
    cb(inst);
  }
  for (const auto& func : functions_) {
    func.iterate(cb);
  }
}

void Builder::push_capability(uint32_t cap) {
  if (capability_set_.count(cap) == 0) {
    capability_set_.insert(cap);
    capabilities_.push_back(
        Instruction{spv::Op::OpCapability, {Operand::Int(cap)}});
  }
}

bool Builder::GenerateLabel(uint32_t id) {
  if (!push_function_inst(spv::Op::OpLabel, {Operand::Int(id)})) {
    return false;
  }
  current_label_id_ = id;
  return true;
}

bool Builder::GenerateAssignStatement(const ast::AssignmentStatement* assign) {
  if (assign->lhs->Is<ast::PhonyExpression>()) {
    auto rhs_id = GenerateExpression(assign->rhs);
    if (rhs_id == 0) {
      return false;
    }
    return true;
  } else {
    auto lhs_id = GenerateExpression(assign->lhs);
    if (lhs_id == 0) {
      return false;
    }
    auto rhs_id = GenerateExpressionWithLoadIfNeeded(assign->rhs);
    if (rhs_id == 0) {
      return false;
    }
    return GenerateStore(lhs_id, rhs_id);
  }
}

bool Builder::GenerateBreakStatement(const ast::BreakStatement*) {
  if (merge_stack_.empty()) {
    error_ = "Attempted to break without a merge block";
    return false;
  }
  if (!push_function_inst(spv::Op::OpBranch,
                          {Operand::Int(merge_stack_.back())})) {
    return false;
  }
  return true;
}

bool Builder::GenerateContinueStatement(const ast::ContinueStatement*) {
  if (continue_stack_.empty()) {
    error_ = "Attempted to continue without a continue block";
    return false;
  }
  if (!push_function_inst(spv::Op::OpBranch,
                          {Operand::Int(continue_stack_.back())})) {
    return false;
  }
  return true;
}

// TODO(dsinclair): This is generating an OpKill but the semantics of kill
// haven't been defined for WGSL yet. So, this may need to change.
// https://github.com/gpuweb/gpuweb/issues/676
bool Builder::GenerateDiscardStatement(const ast::DiscardStatement*) {
  if (!push_function_inst(spv::Op::OpKill, {})) {
    return false;
  }
  return true;
}

bool Builder::GenerateEntryPoint(const ast::Function* func, uint32_t id) {
  auto stage = pipeline_stage_to_execution_model(func->PipelineStage());
  if (stage == SpvExecutionModelMax) {
    error_ = "Unknown pipeline stage provided";
    return false;
  }

  OperandList operands = {
      Operand::Int(stage), Operand::Int(id),
      Operand::String(builder_.Symbols().NameFor(func->symbol))};

  auto* func_sem = builder_.Sem().Get(func);
  for (const auto* var : func_sem->TransitivelyReferencedGlobals()) {
    // For SPIR-V 1.3 we only output Input/output variables. If we update to
    // SPIR-V 1.4 or later this should be all variables.
    if (var->StorageClass() != ast::StorageClass::kInput &&
        var->StorageClass() != ast::StorageClass::kOutput) {
      continue;
    }

    uint32_t var_id = scope_stack_.Get(var->Declaration()->symbol);
    if (var_id == 0) {
      error_ = "unable to find ID for global variable: " +
               builder_.Symbols().NameFor(var->Declaration()->symbol);
      return false;
    }

    operands.push_back(Operand::Int(var_id));
  }
  push_entry_point(spv::Op::OpEntryPoint, operands);

  return true;
}

bool Builder::GenerateExecutionModes(const ast::Function* func, uint32_t id) {
  auto* func_sem = builder_.Sem().Get(func);

  // WGSL fragment shader origin is upper left
  if (func->PipelineStage() == ast::PipelineStage::kFragment) {
    push_execution_mode(
        spv::Op::OpExecutionMode,
        {Operand::Int(id), Operand::Int(SpvExecutionModeOriginUpperLeft)});
  } else if (func->PipelineStage() == ast::PipelineStage::kCompute) {
    auto& wgsize = func_sem->WorkgroupSize();

    // Check if the workgroup_size uses pipeline-overridable constants.
    if (wgsize[0].overridable_const || wgsize[1].overridable_const ||
        wgsize[2].overridable_const) {
      if (has_overridable_workgroup_size_) {
        // Only one stage can have a pipeline-overridable workgroup size.
        // TODO(crbug.com/tint/810): Use LocalSizeId to handle this scenario.
        TINT_ICE(Writer, builder_.Diagnostics())
            << "multiple stages using pipeline-overridable workgroup sizes";
      }
      has_overridable_workgroup_size_ = true;

      auto* vec3_u32 =
          builder_.create<sem::Vector>(builder_.create<sem::U32>(), 3u);
      uint32_t vec3_u32_type_id = GenerateTypeIfNeeded(vec3_u32);
      if (vec3_u32_type_id == 0) {
        return 0;
      }

      OperandList wgsize_ops;
      auto wgsize_result = result_op();
      wgsize_ops.push_back(Operand::Int(vec3_u32_type_id));
      wgsize_ops.push_back(wgsize_result);

      // Generate OpConstant instructions for each dimension.
      for (int i = 0; i < 3; i++) {
        auto constant = ScalarConstant::U32(wgsize[i].value);
        if (wgsize[i].overridable_const) {
          // Make the constant specializable.
          auto* sem_const = builder_.Sem().Get<sem::GlobalVariable>(
              wgsize[i].overridable_const);
          if (!sem_const->IsOverridable()) {
            TINT_ICE(Writer, builder_.Diagnostics())
                << "expected a pipeline-overridable constant";
          }
          constant.is_spec_op = true;
          constant.constant_id = sem_const->ConstantId();
        }

        auto result = GenerateConstantIfNeeded(constant);
        wgsize_ops.push_back(Operand::Int(result));
      }

      // Generate the WorkgroupSize builtin.
      push_type(spv::Op::OpSpecConstantComposite, wgsize_ops);
      push_annot(spv::Op::OpDecorate,
                 {wgsize_result, Operand::Int(SpvDecorationBuiltIn),
                  Operand::Int(SpvBuiltInWorkgroupSize)});
    } else {
      // Not overridable, so just use OpExecutionMode LocalSize.
      uint32_t x = wgsize[0].value;
      uint32_t y = wgsize[1].value;
      uint32_t z = wgsize[2].value;
      push_execution_mode(
          spv::Op::OpExecutionMode,
          {Operand::Int(id), Operand::Int(SpvExecutionModeLocalSize),
           Operand::Int(x), Operand::Int(y), Operand::Int(z)});
    }
  }

  for (auto builtin : func_sem->TransitivelyReferencedBuiltinVariables()) {
    if (builtin.second->builtin == ast::Builtin::kFragDepth) {
      push_execution_mode(
          spv::Op::OpExecutionMode,
          {Operand::Int(id), Operand::Int(SpvExecutionModeDepthReplacing)});
    }
  }

  return true;
}

uint32_t Builder::GenerateExpression(const ast::Expression* expr) {
  return Switch(
      expr,
      [&](const ast::IndexAccessorExpression* a) {
        return GenerateAccessorExpression(a);
      },
      [&](const ast::BinaryExpression* b) {
        return GenerateBinaryExpression(b);
      },
      [&](const ast::BitcastExpression* b) {
        return GenerateBitcastExpression(b);
      },
      [&](const ast::CallExpression* c) { return GenerateCallExpression(c); },
      [&](const ast::IdentifierExpression* i) {
        return GenerateIdentifierExpression(i);
      },
      [&](const ast::LiteralExpression* l) {
        return GenerateLiteralIfNeeded(nullptr, l);
      },
      [&](const ast::MemberAccessorExpression* m) {
        return GenerateAccessorExpression(m);
      },
      [&](const ast::UnaryOpExpression* u) {
        return GenerateUnaryOpExpression(u);
      },
      [&](Default) {
        error_ =
            "unknown expression type: " + std::string(expr->TypeInfo().name);
        return 0;
      });
}

bool Builder::GenerateFunction(const ast::Function* func_ast) {
  auto* func = builder_.Sem().Get(func_ast);

  uint32_t func_type_id = GenerateFunctionTypeIfNeeded(func);
  if (func_type_id == 0) {
    return false;
  }

  auto func_op = result_op();
  auto func_id = func_op.to_i();

  push_debug(spv::Op::OpName,
             {Operand::Int(func_id),
              Operand::String(builder_.Symbols().NameFor(func_ast->symbol))});

  auto ret_id = GenerateTypeIfNeeded(func->ReturnType());
  if (ret_id == 0) {
    return false;
  }

  scope_stack_.Push();
  TINT_DEFER(scope_stack_.Pop());

  auto definition_inst = Instruction{
      spv::Op::OpFunction,
      {Operand::Int(ret_id), func_op, Operand::Int(SpvFunctionControlMaskNone),
       Operand::Int(func_type_id)}};

  InstructionList params;
  for (auto* param : func->Parameters()) {
    auto param_op = result_op();
    auto param_id = param_op.to_i();

    auto param_type_id = GenerateTypeIfNeeded(param->Type());
    if (param_type_id == 0) {
      return false;
    }

    push_debug(spv::Op::OpName, {Operand::Int(param_id),
                                 Operand::String(builder_.Symbols().NameFor(
                                     param->Declaration()->symbol))});
    params.push_back(Instruction{spv::Op::OpFunctionParameter,
                                 {Operand::Int(param_type_id), param_op}});

    scope_stack_.Set(param->Declaration()->symbol, param_id);
  }

  push_function(Function{definition_inst, result_op(), std::move(params)});

  for (auto* stmt : func_ast->body->statements) {
    if (!GenerateStatement(stmt)) {
      return false;
    }
  }

  if (InsideBasicBlock()) {
    if (func->ReturnType()->Is<sem::Void>()) {
      push_function_inst(spv::Op::OpReturn, {});
    } else {
      auto zero = GenerateConstantNullIfNeeded(func->ReturnType());
      push_function_inst(spv::Op::OpReturnValue, {Operand::Int(zero)});
    }
  }

  if (func_ast->IsEntryPoint()) {
    if (!GenerateEntryPoint(func_ast, func_id)) {
      return false;
    }
    if (!GenerateExecutionModes(func_ast, func_id)) {
      return false;
    }
  }

  func_symbol_to_id_[func_ast->symbol] = func_id;

  return true;
}

uint32_t Builder::GenerateFunctionTypeIfNeeded(const sem::Function* func) {
  return utils::GetOrCreate(
      func_sig_to_id_, func->Signature(), [&]() -> uint32_t {
        auto func_op = result_op();
        auto func_type_id = func_op.to_i();

        auto ret_id = GenerateTypeIfNeeded(func->ReturnType());
        if (ret_id == 0) {
          return 0;
        }

        OperandList ops = {func_op, Operand::Int(ret_id)};
        for (auto* param : func->Parameters()) {
          auto param_type_id = GenerateTypeIfNeeded(param->Type());
          if (param_type_id == 0) {
            return 0;
          }
          ops.push_back(Operand::Int(param_type_id));
        }

        push_type(spv::Op::OpTypeFunction, std::move(ops));
        return func_type_id;
      });
}

bool Builder::GenerateFunctionVariable(const ast::Variable* var) {
  uint32_t init_id = 0;
  if (var->constructor) {
    init_id = GenerateExpressionWithLoadIfNeeded(var->constructor);
    if (init_id == 0) {
      return false;
    }
  }

  if (var->is_const) {
    if (!var->constructor) {
      error_ = "missing constructor for constant";
      return false;
    }
    scope_stack_.Set(var->symbol, init_id);
    spirv_id_to_variable_[init_id] = var;
    return true;
  }

  auto result = result_op();
  auto var_id = result.to_i();
  auto sc = ast::StorageClass::kFunction;
  auto* type = builder_.Sem().Get(var)->Type();
  auto type_id = GenerateTypeIfNeeded(type);
  if (type_id == 0) {
    return false;
  }

  push_debug(spv::Op::OpName,
             {Operand::Int(var_id),
              Operand::String(builder_.Symbols().NameFor(var->symbol))});

  // TODO(dsinclair) We could detect if the constructor is fully const and emit
  // an initializer value for the variable instead of doing the OpLoad.
  auto null_id = GenerateConstantNullIfNeeded(type->UnwrapRef());
  if (null_id == 0) {
    return 0;
  }
  push_function_var({Operand::Int(type_id), result,
                     Operand::Int(ConvertStorageClass(sc)),
                     Operand::Int(null_id)});

  if (var->constructor) {
    if (!GenerateStore(var_id, init_id)) {
      return false;
    }
  }

  scope_stack_.Set(var->symbol, var_id);
  spirv_id_to_variable_[var_id] = var;

  return true;
}

bool Builder::GenerateStore(uint32_t to, uint32_t from) {
  return push_function_inst(spv::Op::OpStore,
                            {Operand::Int(to), Operand::Int(from)});
}

bool Builder::GenerateGlobalVariable(const ast::Variable* var) {
  auto* sem = builder_.Sem().Get(var);
  auto* type = sem->Type()->UnwrapRef();

  uint32_t init_id = 0;
  if (var->constructor) {
    init_id = GenerateConstructorExpression(var, var->constructor);
    if (init_id == 0) {
      return false;
    }
  }

  if (var->is_const) {
    if (!var->constructor) {
      // Constants must have an initializer unless they are overridable.
      if (!var->is_overridable) {
        error_ = "missing constructor for constant";
        return false;
      }

      // SPIR-V requires specialization constants to have initializers.
      if (type->Is<sem::F32>()) {
        ast::FloatLiteralExpression l(ProgramID(), Source{}, 0.0f);
        init_id = GenerateLiteralIfNeeded(var, &l);
      } else if (type->Is<sem::U32>()) {
        ast::UintLiteralExpression l(ProgramID(), Source{}, 0);
        init_id = GenerateLiteralIfNeeded(var, &l);
      } else if (type->Is<sem::I32>()) {
        ast::SintLiteralExpression l(ProgramID(), Source{}, 0);
        init_id = GenerateLiteralIfNeeded(var, &l);
      } else if (type->Is<sem::Bool>()) {
        ast::BoolLiteralExpression l(ProgramID(), Source{}, false);
        init_id = GenerateLiteralIfNeeded(var, &l);
      } else {
        error_ = "invalid type for pipeline constant ID, must be scalar";
        return false;
      }
      if (init_id == 0) {
        return 0;
      }
    }
    push_debug(spv::Op::OpName,
               {Operand::Int(init_id),
                Operand::String(builder_.Symbols().NameFor(var->symbol))});

    scope_stack_.Set(var->symbol, init_id);
    spirv_id_to_variable_[init_id] = var;
    return true;
  }

  auto result = result_op();
  auto var_id = result.to_i();

  auto sc = sem->StorageClass() == ast::StorageClass::kNone
                ? ast::StorageClass::kPrivate
                : sem->StorageClass();

  auto type_id = GenerateTypeIfNeeded(sem->Type());
  if (type_id == 0) {
    return false;
  }

  push_debug(spv::Op::OpName,
             {Operand::Int(var_id),
              Operand::String(builder_.Symbols().NameFor(var->symbol))});

  OperandList ops = {Operand::Int(type_id), result,
                     Operand::Int(ConvertStorageClass(sc))};

  if (var->constructor) {
    ops.push_back(Operand::Int(init_id));
  } else {
    auto* st = type->As<sem::StorageTexture>();
    if (st || type->Is<sem::Struct>()) {
      // type is a sem::Struct or a sem::StorageTexture
      auto access = st ? st->access() : sem->Access();
      switch (access) {
        case ast::Access::kWrite:
          push_annot(
              spv::Op::OpDecorate,
              {Operand::Int(var_id), Operand::Int(SpvDecorationNonReadable)});
          break;
        case ast::Access::kRead:
          push_annot(
              spv::Op::OpDecorate,
              {Operand::Int(var_id), Operand::Int(SpvDecorationNonWritable)});
          break;
        case ast::Access::kUndefined:
        case ast::Access::kReadWrite:
          break;
      }
    }
    if (!type->Is<sem::Sampler>()) {
      // If we don't have a constructor and we're an Output or Private
      // variable, then WGSL requires that we zero-initialize.
      // If we're a Workgroup variable, and the
      // VK_KHR_zero_initialize_workgroup_memory extension is enabled, we should
      // also zero-initialize.
      if (sem->StorageClass() == ast::StorageClass::kPrivate ||
          sem->StorageClass() == ast::StorageClass::kOutput ||
          (zero_initialize_workgroup_memory_ &&
           sem->StorageClass() == ast::StorageClass::kWorkgroup)) {
        init_id = GenerateConstantNullIfNeeded(type);
        if (init_id == 0) {
          return 0;
        }
        ops.push_back(Operand::Int(init_id));
      }
    }
  }

  push_type(spv::Op::OpVariable, std::move(ops));

  for (auto* attr : var->attributes) {
    bool ok = Switch(
        attr,
        [&](const ast::BuiltinAttribute* builtin) {
          push_annot(spv::Op::OpDecorate,
                     {Operand::Int(var_id), Operand::Int(SpvDecorationBuiltIn),
                      Operand::Int(ConvertBuiltin(builtin->builtin,
                                                  sem->StorageClass()))});
          return true;
        },
        [&](const ast::LocationAttribute* location) {
          push_annot(spv::Op::OpDecorate,
                     {Operand::Int(var_id), Operand::Int(SpvDecorationLocation),
                      Operand::Int(location->value)});
          return true;
        },
        [&](const ast::InterpolateAttribute* interpolate) {
          AddInterpolationDecorations(var_id, interpolate->type,
                                      interpolate->sampling);
          return true;
        },
        [&](const ast::InvariantAttribute*) {
          push_annot(
              spv::Op::OpDecorate,
              {Operand::Int(var_id), Operand::Int(SpvDecorationInvariant)});
          return true;
        },
        [&](const ast::BindingAttribute* binding) {
          push_annot(spv::Op::OpDecorate,
                     {Operand::Int(var_id), Operand::Int(SpvDecorationBinding),
                      Operand::Int(binding->value)});
          return true;
        },
        [&](const ast::GroupAttribute* group) {
          push_annot(
              spv::Op::OpDecorate,
              {Operand::Int(var_id), Operand::Int(SpvDecorationDescriptorSet),
               Operand::Int(group->value)});
          return true;
        },
        [&](const ast::IdAttribute*) {
          return true;  // Spec constants are handled elsewhere
        },
        [&](const ast::InternalAttribute*) {
          return true;  // ignored
        },
        [&](Default) {
          error_ = "unknown attribute";
          return false;
        });
    if (!ok) {
      return false;
    }
  }

  scope_stack_.Set(var->symbol, var_id);
  spirv_id_to_variable_[var_id] = var;
  return true;
}

bool Builder::GenerateIndexAccessor(const ast::IndexAccessorExpression* expr,
                                    AccessorInfo* info) {
  auto idx_id = GenerateExpressionWithLoadIfNeeded(expr->index);
  if (idx_id == 0) {
    return 0;
  }

  // If the source is a reference, we access chain into it.
  // In the future, pointers may support access-chaining.
  // See https://github.com/gpuweb/gpuweb/pull/1580
  if (info->source_type->Is<sem::Reference>()) {
    info->access_chain_indices.push_back(idx_id);
    info->source_type = TypeOf(expr);
    return true;
  }

  auto result_type_id = GenerateTypeIfNeeded(TypeOf(expr));
  if (result_type_id == 0) {
    return false;
  }

  // We don't have a pointer, so we can just directly extract the value.
  auto extract = result_op();
  auto extract_id = extract.to_i();

  // If the index is compile-time constant, we use OpCompositeExtract.
  auto* idx = builder_.Sem().Get(expr->index);
  if (auto idx_constval = idx->ConstantValue()) {
    if (!push_function_inst(
            spv::Op::OpCompositeExtract,
            {
                Operand::Int(result_type_id),
                extract,
                Operand::Int(info->source_id),
                Operand::Int(idx_constval.ElementAs<uint32_t>(0)),
            })) {
      return false;
    }

    info->source_id = extract_id;
    info->source_type = TypeOf(expr);

    return true;
  }

  // If the source is a vector, we use OpVectorExtractDynamic.
  if (info->source_type->Is<sem::Vector>()) {
    if (!push_function_inst(
            spv::Op::OpVectorExtractDynamic,
            {Operand::Int(result_type_id), extract,
             Operand::Int(info->source_id), Operand::Int(idx_id)})) {
      return false;
    }

    info->source_id = extract_id;
    info->source_type = TypeOf(expr);

    return true;
  }

  TINT_ICE(Writer, builder_.Diagnostics())
      << "unsupported index accessor expression";
  return false;
}

bool Builder::GenerateMemberAccessor(const ast::MemberAccessorExpression* expr,
                                     AccessorInfo* info) {
  auto* expr_sem = builder_.Sem().Get(expr);
  auto* expr_type = expr_sem->Type();

  if (auto* access = expr_sem->As<sem::StructMemberAccess>()) {
    uint32_t idx = access->Member()->Index();

    if (info->source_type->Is<sem::Reference>()) {
      auto idx_id = GenerateConstantIfNeeded(ScalarConstant::U32(idx));
      if (idx_id == 0) {
        return 0;
      }
      info->access_chain_indices.push_back(idx_id);
      info->source_type = expr_type;
    } else {
      auto result_type_id = GenerateTypeIfNeeded(expr_type);
      if (result_type_id == 0) {
        return false;
      }

      auto extract = result_op();
      auto extract_id = extract.to_i();
      if (!push_function_inst(
              spv::Op::OpCompositeExtract,
              {Operand::Int(result_type_id), extract,
               Operand::Int(info->source_id), Operand::Int(idx)})) {
        return false;
      }

      info->source_id = extract_id;
      info->source_type = expr_type;
    }

    return true;
  }

  if (auto* swizzle = expr_sem->As<sem::Swizzle>()) {
    // Single element swizzle is either an access chain or a composite extract
    auto& indices = swizzle->Indices();
    if (indices.size() == 1) {
      if (info->source_type->Is<sem::Reference>()) {
        auto idx_id = GenerateConstantIfNeeded(ScalarConstant::U32(indices[0]));
        if (idx_id == 0) {
          return 0;
        }
        info->access_chain_indices.push_back(idx_id);
      } else {
        auto result_type_id = GenerateTypeIfNeeded(expr_type);
        if (result_type_id == 0) {
          return 0;
        }

        auto extract = result_op();
        auto extract_id = extract.to_i();
        if (!push_function_inst(
                spv::Op::OpCompositeExtract,
                {Operand::Int(result_type_id), extract,
                 Operand::Int(info->source_id), Operand::Int(indices[0])})) {
          return false;
        }

        info->source_id = extract_id;
        info->source_type = expr_type;
      }
      return true;
    }

    // Store the type away as it may change if we run the access chain
    auto* incoming_type = info->source_type;

    // Multi-item extract is a VectorShuffle. We have to emit any existing
    // access chain data, then load the access chain and shuffle that.
    if (!info->access_chain_indices.empty()) {
      auto result_type_id = GenerateTypeIfNeeded(info->source_type);
      if (result_type_id == 0) {
        return 0;
      }
      auto extract = result_op();
      auto extract_id = extract.to_i();

      OperandList ops = {Operand::Int(result_type_id), extract,
                         Operand::Int(info->source_id)};
      for (auto id : info->access_chain_indices) {
        ops.push_back(Operand::Int(id));
      }

      if (!push_function_inst(spv::Op::OpAccessChain, ops)) {
        return false;
      }

      info->source_id = GenerateLoadIfNeeded(expr_type, extract_id);
      info->source_type = expr_type->UnwrapRef();
      info->access_chain_indices.clear();
    }

    auto result_type_id = GenerateTypeIfNeeded(expr_type);
    if (result_type_id == 0) {
      return false;
    }

    auto vec_id = GenerateLoadIfNeeded(incoming_type, info->source_id);

    auto result = result_op();
    auto result_id = result.to_i();

    OperandList ops = {Operand::Int(result_type_id), result,
                       Operand::Int(vec_id), Operand::Int(vec_id)};

    for (auto idx : indices) {
      ops.push_back(Operand::Int(idx));
    }

    if (!push_function_inst(spv::Op::OpVectorShuffle, ops)) {
      return false;
    }
    info->source_id = result_id;
    info->source_type = expr_type;
    return true;
  }

  TINT_ICE(Writer, builder_.Diagnostics())
      << "unhandled member index type: " << expr_sem->TypeInfo().name;
  return false;
}

uint32_t Builder::GenerateAccessorExpression(const ast::Expression* expr) {
  if (!expr->IsAnyOf<ast::IndexAccessorExpression,
                     ast::MemberAccessorExpression>()) {
    TINT_ICE(Writer, builder_.Diagnostics()) << "expression is not an accessor";
    return 0;
  }

  // Gather a list of all the member and index accessors that are in this chain.
  // The list is built in reverse order as that's the order we need to access
  // the chain.
  std::vector<const ast::Expression*> accessors;
  const ast::Expression* source = expr;
  while (true) {
    if (auto* array = source->As<ast::IndexAccessorExpression>()) {
      accessors.insert(accessors.begin(), source);
      source = array->object;
    } else if (auto* member = source->As<ast::MemberAccessorExpression>()) {
      accessors.insert(accessors.begin(), source);
      source = member->structure;
    } else {
      break;
    }
  }

  AccessorInfo info;
  info.source_id = GenerateExpression(source);
  if (info.source_id == 0) {
    return 0;
  }
  info.source_type = TypeOf(source);

  // Note: Dynamic index on array and matrix values (lets) should have been
  // promoted to storage with the VarForDynamicIndex transform.

  for (auto* accessor : accessors) {
    bool ok = Switch(
        accessor,
        [&](const ast::IndexAccessorExpression* array) {
          return GenerateIndexAccessor(array, &info);
        },
        [&](const ast::MemberAccessorExpression* member) {
          return GenerateMemberAccessor(member, &info);
        },
        [&](Default) {
          error_ = "invalid accessor in list: " +
                   std::string(accessor->TypeInfo().name);
          return false;
        });
    if (!ok) {
      return false;
    }
  }

  if (!info.access_chain_indices.empty()) {
    auto* type = TypeOf(expr);
    auto result_type_id = GenerateTypeIfNeeded(type);
    if (result_type_id == 0) {
      return 0;
    }

    auto result = result_op();
    auto result_id = result.to_i();

    OperandList ops = {Operand::Int(result_type_id), result,
                       Operand::Int(info.source_id)};
    for (auto id : info.access_chain_indices) {
      ops.push_back(Operand::Int(id));
    }

    if (!push_function_inst(spv::Op::OpAccessChain, ops)) {
      return false;
    }
    info.source_id = result_id;
  }

  return info.source_id;
}

uint32_t Builder::GenerateIdentifierExpression(
    const ast::IdentifierExpression* expr) {
  uint32_t val = scope_stack_.Get(expr->symbol);
  if (val == 0) {
    error_ = "unable to find variable with identifier: " +
             builder_.Symbols().NameFor(expr->symbol);
  }
  return val;
}

uint32_t Builder::GenerateExpressionWithLoadIfNeeded(
    const sem::Expression* expr) {
  // The semantic node directly knows both the AST node and the resolved type.
  if (const auto id = GenerateExpression(expr->Declaration())) {
    return GenerateLoadIfNeeded(expr->Type(), id);
  }
  return 0;
}

uint32_t Builder::GenerateExpressionWithLoadIfNeeded(
    const ast::Expression* expr) {
  if (const auto id = GenerateExpression(expr)) {
    // Perform a lookup to get the resolved type.
    return GenerateLoadIfNeeded(TypeOf(expr), id);
  }
  return 0;
}

uint32_t Builder::GenerateLoadIfNeeded(const sem::Type* type, uint32_t id) {
  if (auto* ref = type->As<sem::Reference>()) {
    type = ref->StoreType();
  } else {
    return id;
  }

  auto type_id = GenerateTypeIfNeeded(type);
  auto result = result_op();
  auto result_id = result.to_i();
  if (!push_function_inst(spv::Op::OpLoad,
                          {Operand::Int(type_id), result, Operand::Int(id)})) {
    return 0;
  }
  return result_id;
}

uint32_t Builder::GenerateUnaryOpExpression(
    const ast::UnaryOpExpression* expr) {
  auto result = result_op();
  auto result_id = result.to_i();

  spv::Op op = spv::Op::OpNop;
  switch (expr->op) {
    case ast::UnaryOp::kComplement:
      op = spv::Op::OpNot;
      break;
    case ast::UnaryOp::kNegation:
      if (TypeOf(expr)->is_float_scalar_or_vector()) {
        op = spv::Op::OpFNegate;
      } else {
        op = spv::Op::OpSNegate;
      }
      break;
    case ast::UnaryOp::kNot:
      op = spv::Op::OpLogicalNot;
      break;
    case ast::UnaryOp::kAddressOf:
    case ast::UnaryOp::kIndirection:
      // Address-of converts a reference to a pointer, and dereference converts
      // a pointer to a reference. These are the same thing in SPIR-V, so this
      // is a no-op.
      return GenerateExpression(expr->expr);
  }

  auto val_id = GenerateExpressionWithLoadIfNeeded(expr->expr);
  if (val_id == 0) {
    return 0;
  }

  auto type_id = GenerateTypeIfNeeded(TypeOf(expr));
  if (type_id == 0) {
    return 0;
  }

  if (!push_function_inst(
          op, {Operand::Int(type_id), result, Operand::Int(val_id)})) {
    return false;
  }

  return result_id;
}

uint32_t Builder::GetGLSLstd450Import() {
  auto where = import_name_to_id_.find(kGLSLstd450);
  if (where != import_name_to_id_.end()) {
    return where->second;
  }

  // It doesn't exist yet. Generate it.
  auto result = result_op();
  auto id = result.to_i();

  push_ext_import(spv::Op::OpExtInstImport,
                  {result, Operand::String(kGLSLstd450)});

  // Remember it for later.
  import_name_to_id_[kGLSLstd450] = id;
  return id;
}

uint32_t Builder::GenerateConstructorExpression(const ast::Variable* var,
                                                const ast::Expression* expr) {
  if (auto* literal = expr->As<ast::LiteralExpression>()) {
    return GenerateLiteralIfNeeded(var, literal);
  }
  if (auto* call = builder_.Sem().Get<sem::Call>(expr)) {
    if (call->Target()->IsAnyOf<sem::TypeConstructor, sem::TypeConversion>()) {
      return GenerateTypeConstructorOrConversion(call, var);
    }
  }
  error_ = "unknown constructor expression";
  return 0;
}

bool Builder::IsConstructorConst(const ast::Expression* expr) {
  bool is_const = true;
  ast::TraverseExpressions(expr, builder_.Diagnostics(),
                           [&](const ast::Expression* e) {
                             if (e->Is<ast::LiteralExpression>()) {
                               return ast::TraverseAction::Descend;
                             }
                             if (auto* ce = e->As<ast::CallExpression>()) {
                               auto* call = builder_.Sem().Get(ce);
                               if (call->Target()->Is<sem::TypeConstructor>()) {
                                 return ast::TraverseAction::Descend;
                               }
                             }

                             is_const = false;
                             return ast::TraverseAction::Stop;
                           });
  return is_const;
}

uint32_t Builder::GenerateTypeConstructorOrConversion(
    const sem::Call* call,
    const ast::Variable* var) {
  auto& args = call->Arguments();
  auto* global_var = builder_.Sem().Get<sem::GlobalVariable>(var);
  auto* result_type = call->Type();

  // Generate the zero initializer if there are no values provided.
  if (args.empty()) {
    if (global_var && global_var->IsOverridable()) {
      auto constant_id = global_var->ConstantId();
      if (result_type->Is<sem::I32>()) {
        return GenerateConstantIfNeeded(
            ScalarConstant::I32(0).AsSpecOp(constant_id));
      }
      if (result_type->Is<sem::U32>()) {
        return GenerateConstantIfNeeded(
            ScalarConstant::U32(0).AsSpecOp(constant_id));
      }
      if (result_type->Is<sem::F32>()) {
        return GenerateConstantIfNeeded(
            ScalarConstant::F32(0).AsSpecOp(constant_id));
      }
      if (result_type->Is<sem::Bool>()) {
        return GenerateConstantIfNeeded(
            ScalarConstant::Bool(false).AsSpecOp(constant_id));
      }
    }
    return GenerateConstantNullIfNeeded(result_type->UnwrapRef());
  }

  std::ostringstream out;
  out << "__const_" << result_type->FriendlyName(builder_.Symbols()) << "_";

  result_type = result_type->UnwrapRef();
  bool constructor_is_const = IsConstructorConst(call->Declaration());
  if (has_error()) {
    return 0;
  }

  bool can_cast_or_copy = result_type->is_scalar();

  if (auto* res_vec = result_type->As<sem::Vector>()) {
    if (res_vec->type()->is_scalar()) {
      auto* value_type = args[0]->Type()->UnwrapRef();
      if (auto* val_vec = value_type->As<sem::Vector>()) {
        if (val_vec->type()->is_scalar()) {
          can_cast_or_copy = res_vec->Width() == val_vec->Width();
        }
      }
    }
  }

  if (can_cast_or_copy) {
    return GenerateCastOrCopyOrPassthrough(result_type, args[0]->Declaration(),
                                           global_var);
  }

  auto type_id = GenerateTypeIfNeeded(result_type);
  if (type_id == 0) {
    return 0;
  }

  bool result_is_constant_composite = constructor_is_const;
  bool result_is_spec_composite = false;

  if (auto* vec = result_type->As<sem::Vector>()) {
    result_type = vec->type();
  }

  OperandList ops;
  for (auto* e : args) {
    uint32_t id = 0;
    id = GenerateExpressionWithLoadIfNeeded(e);
    if (id == 0) {
      return 0;
    }

    auto* value_type = e->Type()->UnwrapRef();
    // If the result and value types are the same we can just use the object.
    // If the result is not a vector then we should have validated that the
    // value type is a correctly sized vector so we can just use it directly.
    if (result_type == value_type || result_type->Is<sem::Matrix>() ||
        result_type->Is<sem::Array>() || result_type->Is<sem::Struct>()) {
      out << "_" << id;

      ops.push_back(Operand::Int(id));
      continue;
    }

    // Both scalars, but not the same type so we need to generate a conversion
    // of the value.
    if (value_type->is_scalar() && result_type->is_scalar()) {
      id = GenerateCastOrCopyOrPassthrough(result_type, args[0]->Declaration(),
                                           global_var);
      out << "_" << id;
      ops.push_back(Operand::Int(id));
      continue;
    }

    // When handling vectors as the values there a few cases to take into
    // consideration:
    //  1. Module scoped vec3<f32>(vec2<f32>(1, 2), 3)  -> OpSpecConstantOp
    //  2. Function scoped vec3<f32>(vec2<f32>(1, 2), 3) ->  OpCompositeExtract
    //  3. Either array<vec3<f32>, 1>(vec3<f32>(1, 2, 3))  -> use the ID.
    //       -> handled above
    //
    // For cases 1 and 2, if the type is different we also may need to insert
    // a type cast.
    if (auto* vec = value_type->As<sem::Vector>()) {
      auto* vec_type = vec->type();

      auto value_type_id = GenerateTypeIfNeeded(vec_type);
      if (value_type_id == 0) {
        return 0;
      }

      for (uint32_t i = 0; i < vec->Width(); ++i) {
        auto extract = result_op();
        auto extract_id = extract.to_i();

        if (!global_var) {
          // A non-global initializer. Case 2.
          if (!push_function_inst(spv::Op::OpCompositeExtract,
                                  {Operand::Int(value_type_id), extract,
                                   Operand::Int(id), Operand::Int(i)})) {
            return false;
          }

          // We no longer have a constant composite, but have to do a
          // composite construction as these calls are inside a function.
          result_is_constant_composite = false;
        } else {
          // A global initializer, must use OpSpecConstantOp. Case 1.
          auto idx_id = GenerateConstantIfNeeded(ScalarConstant::U32(i));
          if (idx_id == 0) {
            return 0;
          }
          push_type(spv::Op::OpSpecConstantOp,
                    {Operand::Int(value_type_id), extract,
                     Operand::Int(SpvOpCompositeExtract), Operand::Int(id),
                     Operand::Int(idx_id)});

          result_is_spec_composite = true;
        }

        out << "_" << extract_id;
        ops.push_back(Operand::Int(extract_id));
      }
    } else {
      error_ = "Unhandled type cast value type";
      return 0;
    }
  }

  // For a single-value vector initializer, splat the initializer value.
  auto* const init_result_type = call->Type()->UnwrapRef();
  if (args.size() == 1 && init_result_type->is_scalar_vector() &&
      args[0]->Type()->UnwrapRef()->is_scalar()) {
    size_t vec_size = init_result_type->As<sem::Vector>()->Width();
    for (size_t i = 0; i < (vec_size - 1); ++i) {
      ops.push_back(ops[0]);
    }
  }

  auto str = out.str();
  auto val = type_constructor_to_id_.find(str);
  if (val != type_constructor_to_id_.end()) {
    return val->second;
  }

  auto result = result_op();
  ops.insert(ops.begin(), result);
  ops.insert(ops.begin(), Operand::Int(type_id));

  type_constructor_to_id_[str] = result.to_i();

  if (result_is_spec_composite) {
    push_type(spv::Op::OpSpecConstantComposite, ops);
  } else if (result_is_constant_composite) {
    push_type(spv::Op::OpConstantComposite, ops);
  } else {
    if (!push_function_inst(spv::Op::OpCompositeConstruct, ops)) {
      return 0;
    }
  }

  return result.to_i();
}

uint32_t Builder::GenerateCastOrCopyOrPassthrough(
    const sem::Type* to_type,
    const ast::Expression* from_expr,
    bool is_global_init) {
  // This should not happen as we rely on constant folding to obviate
  // casts/conversions for module-scope variables
  if (is_global_init) {
    TINT_ICE(Writer, builder_.Diagnostics())
        << "Module-level conversions are not supported. Conversions should "
           "have already been constant-folded by the FoldConstants transform.";
    return 0;
  }

  auto elem_type_of = [](const sem::Type* t) -> const sem::Type* {
    if (t->is_scalar()) {
      return t;
    }
    if (auto* v = t->As<sem::Vector>()) {
      return v->type();
    }
    return nullptr;
  };

  auto result = result_op();
  auto result_id = result.to_i();

  auto result_type_id = GenerateTypeIfNeeded(to_type);
  if (result_type_id == 0) {
    return 0;
  }

  auto val_id = GenerateExpressionWithLoadIfNeeded(from_expr);
  if (val_id == 0) {
    return 0;
  }

  auto* from_type = TypeOf(from_expr)->UnwrapRef();

  spv::Op op = spv::Op::OpNop;
  if ((from_type->Is<sem::I32>() && to_type->Is<sem::F32>()) ||
      (from_type->is_signed_integer_vector() && to_type->is_float_vector())) {
    op = spv::Op::OpConvertSToF;
  } else if ((from_type->Is<sem::U32>() && to_type->Is<sem::F32>()) ||
             (from_type->is_unsigned_integer_vector() &&
              to_type->is_float_vector())) {
    op = spv::Op::OpConvertUToF;
  } else if ((from_type->Is<sem::F32>() && to_type->Is<sem::I32>()) ||
             (from_type->is_float_vector() &&
              to_type->is_signed_integer_vector())) {
    op = spv::Op::OpConvertFToS;
  } else if ((from_type->Is<sem::F32>() && to_type->Is<sem::U32>()) ||
             (from_type->is_float_vector() &&
              to_type->is_unsigned_integer_vector())) {
    op = spv::Op::OpConvertFToU;
  } else if ((from_type->Is<sem::Bool>() && to_type->Is<sem::Bool>()) ||
             (from_type->Is<sem::U32>() && to_type->Is<sem::U32>()) ||
             (from_type->Is<sem::I32>() && to_type->Is<sem::I32>()) ||
             (from_type->Is<sem::F32>() && to_type->Is<sem::F32>()) ||
             (from_type->Is<sem::Vector>() && (from_type == to_type))) {
    return val_id;
  } else if ((from_type->Is<sem::I32>() && to_type->Is<sem::U32>()) ||
             (from_type->Is<sem::U32>() && to_type->Is<sem::I32>()) ||
             (from_type->is_signed_integer_vector() &&
              to_type->is_unsigned_integer_vector()) ||
             (from_type->is_unsigned_integer_vector() &&
              to_type->is_integer_scalar_or_vector())) {
    op = spv::Op::OpBitcast;
  } else if ((from_type->is_numeric_scalar() && to_type->Is<sem::Bool>()) ||
             (from_type->is_numeric_vector() && to_type->is_bool_vector())) {
    // Convert scalar (vector) to bool (vector)

    // Return the result of comparing from_expr with zero
    uint32_t zero = GenerateConstantNullIfNeeded(from_type);
    const auto* from_elem_type = elem_type_of(from_type);
    op = from_elem_type->is_integer_scalar() ? spv::Op::OpINotEqual
                                             : spv::Op::OpFUnordNotEqual;
    if (!push_function_inst(
            op, {Operand::Int(result_type_id), Operand::Int(result_id),
                 Operand::Int(val_id), Operand::Int(zero)})) {
      return 0;
    }

    return result_id;
  } else if (from_type->is_bool_scalar_or_vector() &&
             to_type->is_numeric_scalar_or_vector()) {
    // Convert bool scalar/vector to numeric scalar/vector.
    // Use the bool to select between 1 (if true) and 0 (if false).

    const auto* to_elem_type = elem_type_of(to_type);
    uint32_t one_id;
    uint32_t zero_id;
    if (to_elem_type->Is<sem::F32>()) {
      ast::FloatLiteralExpression one(ProgramID(), Source{}, 1.0f);
      ast::FloatLiteralExpression zero(ProgramID(), Source{}, 0.0f);
      one_id = GenerateLiteralIfNeeded(nullptr, &one);
      zero_id = GenerateLiteralIfNeeded(nullptr, &zero);
    } else if (to_elem_type->Is<sem::U32>()) {
      ast::UintLiteralExpression one(ProgramID(), Source{}, 1);
      ast::UintLiteralExpression zero(ProgramID(), Source{}, 0);
      one_id = GenerateLiteralIfNeeded(nullptr, &one);
      zero_id = GenerateLiteralIfNeeded(nullptr, &zero);
    } else if (to_elem_type->Is<sem::I32>()) {
      ast::SintLiteralExpression one(ProgramID(), Source{}, 1);
      ast::SintLiteralExpression zero(ProgramID(), Source{}, 0);
      one_id = GenerateLiteralIfNeeded(nullptr, &one);
      zero_id = GenerateLiteralIfNeeded(nullptr, &zero);
    } else {
      error_ = "invalid destination type for bool conversion";
      return false;
    }
    if (auto* to_vec = to_type->As<sem::Vector>()) {
      // Splat the scalars into vectors.
      one_id = GenerateConstantVectorSplatIfNeeded(to_vec, one_id);
      zero_id = GenerateConstantVectorSplatIfNeeded(to_vec, zero_id);
    }
    if (!one_id || !zero_id) {
      return false;
    }

    op = spv::Op::OpSelect;
    if (!push_function_inst(
            op, {Operand::Int(result_type_id), Operand::Int(result_id),
                 Operand::Int(val_id), Operand::Int(one_id),
                 Operand::Int(zero_id)})) {
      return 0;
    }

    return result_id;
  } else {
    TINT_ICE(Writer, builder_.Diagnostics()) << "Invalid from_type";
  }

  if (op == spv::Op::OpNop) {
    error_ = "unable to determine conversion type for cast, from: " +
             from_type->FriendlyName(builder_.Symbols()) +
             " to: " + to_type->FriendlyName(builder_.Symbols());
    return 0;
  }

  if (!push_function_inst(
          op, {Operand::Int(result_type_id), result, Operand::Int(val_id)})) {
    return 0;
  }

  return result_id;
}

uint32_t Builder::GenerateLiteralIfNeeded(const ast::Variable* var,
                                          const ast::LiteralExpression* lit) {
  ScalarConstant constant;

  auto* global = builder_.Sem().Get<sem::GlobalVariable>(var);
  if (global && global->IsOverridable()) {
    constant.is_spec_op = true;
    constant.constant_id = global->ConstantId();
  }

  Switch(
      lit,
      [&](const ast::BoolLiteralExpression* l) {
        constant.kind = ScalarConstant::Kind::kBool;
        constant.value.b = l->value;
      },
      [&](const ast::SintLiteralExpression* sl) {
        constant.kind = ScalarConstant::Kind::kI32;
        constant.value.i32 = sl->value;
      },
      [&](const ast::UintLiteralExpression* ul) {
        constant.kind = ScalarConstant::Kind::kU32;
        constant.value.u32 = ul->value;
      },
      [&](const ast::FloatLiteralExpression* fl) {
        constant.kind = ScalarConstant::Kind::kF32;
        constant.value.f32 = fl->value;
      },
      [&](Default) { error_ = "unknown literal type"; });

  if (!error_.empty()) {
    return false;
  }

  return GenerateConstantIfNeeded(constant);
}

uint32_t Builder::GenerateConstantIfNeeded(const ScalarConstant& constant) {
  auto it = const_to_id_.find(constant);
  if (it != const_to_id_.end()) {
    return it->second;
  }

  uint32_t type_id = 0;

  switch (constant.kind) {
    case ScalarConstant::Kind::kU32: {
      type_id = GenerateTypeIfNeeded(builder_.create<sem::U32>());
      break;
    }
    case ScalarConstant::Kind::kI32: {
      type_id = GenerateTypeIfNeeded(builder_.create<sem::I32>());
      break;
    }
    case ScalarConstant::Kind::kF32: {
      type_id = GenerateTypeIfNeeded(builder_.create<sem::F32>());
      break;
    }
    case ScalarConstant::Kind::kBool: {
      type_id = GenerateTypeIfNeeded(builder_.create<sem::Bool>());
      break;
    }
  }

  if (type_id == 0) {
    return 0;
  }

  auto result = result_op();
  auto result_id = result.to_i();

  if (constant.is_spec_op) {
    push_annot(spv::Op::OpDecorate,
               {Operand::Int(result_id), Operand::Int(SpvDecorationSpecId),
                Operand::Int(constant.constant_id)});
  }

  switch (constant.kind) {
    case ScalarConstant::Kind::kU32: {
      push_type(
          constant.is_spec_op ? spv::Op::OpSpecConstant : spv::Op::OpConstant,
          {Operand::Int(type_id), result, Operand::Int(constant.value.u32)});
      break;
    }
    case ScalarConstant::Kind::kI32: {
      push_type(
          constant.is_spec_op ? spv::Op::OpSpecConstant : spv::Op::OpConstant,
          {Operand::Int(type_id), result, Operand::Int(constant.value.i32)});
      break;
    }
    case ScalarConstant::Kind::kF32: {
      push_type(
          constant.is_spec_op ? spv::Op::OpSpecConstant : spv::Op::OpConstant,
          {Operand::Int(type_id), result, Operand::Float(constant.value.f32)});
      break;
    }
    case ScalarConstant::Kind::kBool: {
      if (constant.value.b) {
        push_type(constant.is_spec_op ? spv::Op::OpSpecConstantTrue
                                      : spv::Op::OpConstantTrue,
                  {Operand::Int(type_id), result});
      } else {
        push_type(constant.is_spec_op ? spv::Op::OpSpecConstantFalse
                                      : spv::Op::OpConstantFalse,
                  {Operand::Int(type_id), result});
      }
      break;
    }
  }

  const_to_id_[constant] = result_id;
  return result_id;
}

uint32_t Builder::GenerateConstantNullIfNeeded(const sem::Type* type) {
  auto type_id = GenerateTypeIfNeeded(type);
  if (type_id == 0) {
    return 0;
  }

  return utils::GetOrCreate(const_null_to_id_, type, [&] {
    auto result = result_op();

    push_type(spv::Op::OpConstantNull, {Operand::Int(type_id), result});

    return result.to_i();
  });
}

uint32_t Builder::GenerateConstantVectorSplatIfNeeded(const sem::Vector* type,
                                                      uint32_t value_id) {
  auto type_id = GenerateTypeIfNeeded(type);
  if (type_id == 0 || value_id == 0) {
    return 0;
  }

  uint64_t key = (static_cast<uint64_t>(type->Width()) << 32) + value_id;
  return utils::GetOrCreate(const_splat_to_id_, key, [&] {
    auto result = result_op();
    auto result_id = result.to_i();

    OperandList ops;
    ops.push_back(Operand::Int(type_id));
    ops.push_back(result);
    for (uint32_t i = 0; i < type->Width(); i++) {
      ops.push_back(Operand::Int(value_id));
    }
    push_type(spv::Op::OpConstantComposite, ops);

    const_splat_to_id_[key] = result_id;
    return result_id;
  });
}

uint32_t Builder::GenerateShortCircuitBinaryExpression(
    const ast::BinaryExpression* expr) {
  auto lhs_id = GenerateExpressionWithLoadIfNeeded(expr->lhs);
  if (lhs_id == 0) {
    return false;
  }

  // Get the ID of the basic block where control flow will diverge. It's the
  // last basic block generated for the left-hand-side of the operator.
  auto original_label_id = current_label_id_;

  auto type_id = GenerateTypeIfNeeded(TypeOf(expr));
  if (type_id == 0) {
    return 0;
  }

  auto merge_block = result_op();
  auto merge_block_id = merge_block.to_i();

  auto block = result_op();
  auto block_id = block.to_i();

  auto true_block_id = block_id;
  auto false_block_id = merge_block_id;

  // For a logical or we want to only check the RHS if the LHS is failed.
  if (expr->IsLogicalOr()) {
    std::swap(true_block_id, false_block_id);
  }

  if (!push_function_inst(spv::Op::OpSelectionMerge,
                          {Operand::Int(merge_block_id),
                           Operand::Int(SpvSelectionControlMaskNone)})) {
    return 0;
  }
  if (!push_function_inst(spv::Op::OpBranchConditional,
                          {Operand::Int(lhs_id), Operand::Int(true_block_id),
                           Operand::Int(false_block_id)})) {
    return 0;
  }

  // Output block to check the RHS
  if (!GenerateLabel(block_id)) {
    return 0;
  }
  auto rhs_id = GenerateExpressionWithLoadIfNeeded(expr->rhs);
  if (rhs_id == 0) {
    return 0;
  }

  // Get the block ID of the last basic block generated for the right-hand-side
  // expression. That block will be an immediate predecessor to the merge block.
  auto rhs_block_id = current_label_id_;
  if (!push_function_inst(spv::Op::OpBranch, {Operand::Int(merge_block_id)})) {
    return 0;
  }

  // Output the merge block
  if (!GenerateLabel(merge_block_id)) {
    return 0;
  }

  auto result = result_op();
  auto result_id = result.to_i();

  if (!push_function_inst(spv::Op::OpPhi,
                          {Operand::Int(type_id), result, Operand::Int(lhs_id),
                           Operand::Int(original_label_id),
                           Operand::Int(rhs_id), Operand::Int(rhs_block_id)})) {
    return 0;
  }

  return result_id;
}

uint32_t Builder::GenerateSplat(uint32_t scalar_id, const sem::Type* vec_type) {
  // Create a new vector to splat scalar into
  auto splat_vector = result_op();
  auto* splat_vector_type = builder_.create<sem::Pointer>(
      vec_type, ast::StorageClass::kFunction, ast::Access::kReadWrite);
  push_function_var(
      {Operand::Int(GenerateTypeIfNeeded(splat_vector_type)), splat_vector,
       Operand::Int(ConvertStorageClass(ast::StorageClass::kFunction)),
       Operand::Int(GenerateConstantNullIfNeeded(vec_type))});

  // Splat scalar into vector
  auto splat_result = result_op();
  OperandList ops;
  ops.push_back(Operand::Int(GenerateTypeIfNeeded(vec_type)));
  ops.push_back(splat_result);
  for (size_t i = 0; i < vec_type->As<sem::Vector>()->Width(); ++i) {
    ops.push_back(Operand::Int(scalar_id));
  }
  if (!push_function_inst(spv::Op::OpCompositeConstruct, ops)) {
    return 0;
  }

  return splat_result.to_i();
}

uint32_t Builder::GenerateMatrixAddOrSub(uint32_t lhs_id,
                                         uint32_t rhs_id,
                                         const sem::Matrix* type,
                                         spv::Op op) {
  // Example addition of two matrices:
  // %31 = OpLoad %mat3v4float %m34
  // %32 = OpLoad %mat3v4float %m34
  // %33 = OpCompositeExtract %v4float %31 0
  // %34 = OpCompositeExtract %v4float %32 0
  // %35 = OpFAdd %v4float %33 %34
  // %36 = OpCompositeExtract %v4float %31 1
  // %37 = OpCompositeExtract %v4float %32 1
  // %38 = OpFAdd %v4float %36 %37
  // %39 = OpCompositeExtract %v4float %31 2
  // %40 = OpCompositeExtract %v4float %32 2
  // %41 = OpFAdd %v4float %39 %40
  // %42 = OpCompositeConstruct %mat3v4float %35 %38 %41

  auto* column_type = builder_.create<sem::Vector>(type->type(), type->rows());
  auto column_type_id = GenerateTypeIfNeeded(column_type);

  OperandList ops;

  for (uint32_t i = 0; i < type->columns(); ++i) {
    // Extract column `i` from lhs mat
    auto lhs_column_id = result_op();
    if (!push_function_inst(spv::Op::OpCompositeExtract,
                            {Operand::Int(column_type_id), lhs_column_id,
                             Operand::Int(lhs_id), Operand::Int(i)})) {
      return 0;
    }

    // Extract column `i` from rhs mat
    auto rhs_column_id = result_op();
    if (!push_function_inst(spv::Op::OpCompositeExtract,
                            {Operand::Int(column_type_id), rhs_column_id,
                             Operand::Int(rhs_id), Operand::Int(i)})) {
      return 0;
    }

    // Add or subtract the two columns
    auto result = result_op();
    if (!push_function_inst(op, {Operand::Int(column_type_id), result,
                                 lhs_column_id, rhs_column_id})) {
      return 0;
    }

    ops.push_back(result);
  }

  // Create the result matrix from the added/subtracted column vectors
  auto result_mat_id = result_op();
  ops.insert(ops.begin(), result_mat_id);
  ops.insert(ops.begin(), Operand::Int(GenerateTypeIfNeeded(type)));
  if (!push_function_inst(spv::Op::OpCompositeConstruct, ops)) {
    return 0;
  }

  return result_mat_id.to_i();
}

uint32_t Builder::GenerateBinaryExpression(const ast::BinaryExpression* expr) {
  // There is special logic for short circuiting operators.
  if (expr->IsLogicalAnd() || expr->IsLogicalOr()) {
    return GenerateShortCircuitBinaryExpression(expr);
  }

  auto lhs_id = GenerateExpressionWithLoadIfNeeded(expr->lhs);
  if (lhs_id == 0) {
    return 0;
  }

  auto rhs_id = GenerateExpressionWithLoadIfNeeded(expr->rhs);
  if (rhs_id == 0) {
    return 0;
  }

  auto result = result_op();
  auto result_id = result.to_i();

  auto type_id = GenerateTypeIfNeeded(TypeOf(expr));
  if (type_id == 0) {
    return 0;
  }

  // Handle int and float and the vectors of those types. Other types
  // should have been rejected by validation.
  auto* lhs_type = TypeOf(expr->lhs)->UnwrapRef();
  auto* rhs_type = TypeOf(expr->rhs)->UnwrapRef();

  // Handle matrix-matrix addition and subtraction
  if ((expr->IsAdd() || expr->IsSubtract()) && lhs_type->is_float_matrix() &&
      rhs_type->is_float_matrix()) {
    auto* lhs_mat = lhs_type->As<sem::Matrix>();
    auto* rhs_mat = rhs_type->As<sem::Matrix>();

    // This should already have been validated by resolver
    if (lhs_mat->rows() != rhs_mat->rows() ||
        lhs_mat->columns() != rhs_mat->columns()) {
      error_ = "matrices must have same dimensionality for add or subtract";
      return 0;
    }

    return GenerateMatrixAddOrSub(
        lhs_id, rhs_id, lhs_mat,
        expr->IsAdd() ? spv::Op::OpFAdd : spv::Op::OpFSub);
  }

  // For vector-scalar arithmetic operations, splat scalar into a vector. We
  // skip this for multiply as we can use OpVectorTimesScalar.
  const bool is_float_scalar_vector_multiply =
      expr->IsMultiply() &&
      ((lhs_type->is_float_scalar() && rhs_type->is_float_vector()) ||
       (lhs_type->is_float_vector() && rhs_type->is_float_scalar()));

  if (expr->IsArithmetic() && !is_float_scalar_vector_multiply) {
    if (lhs_type->Is<sem::Vector>() && rhs_type->is_numeric_scalar()) {
      uint32_t splat_vector_id = GenerateSplat(rhs_id, lhs_type);
      if (splat_vector_id == 0) {
        return 0;
      }
      rhs_id = splat_vector_id;
      rhs_type = lhs_type;

    } else if (lhs_type->is_numeric_scalar() && rhs_type->Is<sem::Vector>()) {
      uint32_t splat_vector_id = GenerateSplat(lhs_id, rhs_type);
      if (splat_vector_id == 0) {
        return 0;
      }
      lhs_id = splat_vector_id;
      lhs_type = rhs_type;
    }
  }

  bool lhs_is_float_or_vec = lhs_type->is_float_scalar_or_vector();
  bool lhs_is_bool_or_vec = lhs_type->is_bool_scalar_or_vector();
  bool lhs_is_integer_or_vec = lhs_type->is_integer_scalar_or_vector();
  bool lhs_is_unsigned = lhs_type->is_unsigned_scalar_or_vector();

  spv::Op op = spv::Op::OpNop;
  if (expr->IsAnd()) {
    if (lhs_is_integer_or_vec) {
      op = spv::Op::OpBitwiseAnd;
    } else if (lhs_is_bool_or_vec) {
      op = spv::Op::OpLogicalAnd;
    } else {
      error_ = "invalid and expression";
      return 0;
    }
  } else if (expr->IsAdd()) {
    op = lhs_is_float_or_vec ? spv::Op::OpFAdd : spv::Op::OpIAdd;
  } else if (expr->IsDivide()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFDiv;
    } else if (lhs_is_unsigned) {
      op = spv::Op::OpUDiv;
    } else {
      op = spv::Op::OpSDiv;
    }
  } else if (expr->IsEqual()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFOrdEqual;
    } else if (lhs_is_bool_or_vec) {
      op = spv::Op::OpLogicalEqual;
    } else if (lhs_is_integer_or_vec) {
      op = spv::Op::OpIEqual;
    } else {
      error_ = "invalid equal expression";
      return 0;
    }
  } else if (expr->IsGreaterThan()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFOrdGreaterThan;
    } else if (lhs_is_unsigned) {
      op = spv::Op::OpUGreaterThan;
    } else {
      op = spv::Op::OpSGreaterThan;
    }
  } else if (expr->IsGreaterThanEqual()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFOrdGreaterThanEqual;
    } else if (lhs_is_unsigned) {
      op = spv::Op::OpUGreaterThanEqual;
    } else {
      op = spv::Op::OpSGreaterThanEqual;
    }
  } else if (expr->IsLessThan()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFOrdLessThan;
    } else if (lhs_is_unsigned) {
      op = spv::Op::OpULessThan;
    } else {
      op = spv::Op::OpSLessThan;
    }
  } else if (expr->IsLessThanEqual()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFOrdLessThanEqual;
    } else if (lhs_is_unsigned) {
      op = spv::Op::OpULessThanEqual;
    } else {
      op = spv::Op::OpSLessThanEqual;
    }
  } else if (expr->IsModulo()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFRem;
    } else if (lhs_is_unsigned) {
      op = spv::Op::OpUMod;
    } else {
      op = spv::Op::OpSMod;
    }
  } else if (expr->IsMultiply()) {
    if (lhs_type->is_integer_scalar_or_vector()) {
      // If the left hand side is an integer then this _has_ to be OpIMul as
      // there there is no other integer multiplication.
      op = spv::Op::OpIMul;
    } else if (lhs_type->is_float_scalar() && rhs_type->is_float_scalar()) {
      // Float scalars multiply with OpFMul
      op = spv::Op::OpFMul;
    } else if (lhs_type->is_float_vector() && rhs_type->is_float_vector()) {
      // Float vectors must be validated to be the same size and then use OpFMul
      op = spv::Op::OpFMul;
    } else if (lhs_type->is_float_scalar() && rhs_type->is_float_vector()) {
      // Scalar * Vector we need to flip lhs and rhs types
      // because OpVectorTimesScalar expects <vector>, <scalar>
      std::swap(lhs_id, rhs_id);
      op = spv::Op::OpVectorTimesScalar;
    } else if (lhs_type->is_float_vector() && rhs_type->is_float_scalar()) {
      // float vector * scalar
      op = spv::Op::OpVectorTimesScalar;
    } else if (lhs_type->is_float_scalar() && rhs_type->is_float_matrix()) {
      // Scalar * Matrix we need to flip lhs and rhs types because
      // OpMatrixTimesScalar expects <matrix>, <scalar>
      std::swap(lhs_id, rhs_id);
      op = spv::Op::OpMatrixTimesScalar;
    } else if (lhs_type->is_float_matrix() && rhs_type->is_float_scalar()) {
      // float matrix * scalar
      op = spv::Op::OpMatrixTimesScalar;
    } else if (lhs_type->is_float_vector() && rhs_type->is_float_matrix()) {
      // float vector * matrix
      op = spv::Op::OpVectorTimesMatrix;
    } else if (lhs_type->is_float_matrix() && rhs_type->is_float_vector()) {
      // float matrix * vector
      op = spv::Op::OpMatrixTimesVector;
    } else if (lhs_type->is_float_matrix() && rhs_type->is_float_matrix()) {
      // float matrix * matrix
      op = spv::Op::OpMatrixTimesMatrix;
    } else {
      error_ = "invalid multiply expression";
      return 0;
    }
  } else if (expr->IsNotEqual()) {
    if (lhs_is_float_or_vec) {
      op = spv::Op::OpFOrdNotEqual;
    } else if (lhs_is_bool_or_vec) {
      op = spv::Op::OpLogicalNotEqual;
    } else if (lhs_is_integer_or_vec) {
      op = spv::Op::OpINotEqual;
    } else {
      error_ = "invalid not-equal expression";
      return 0;
    }
  } else if (expr->IsOr()) {
    if (lhs_is_integer_or_vec) {
      op = spv::Op::OpBitwiseOr;
    } else if (lhs_is_bool_or_vec) {
      op = spv::Op::OpLogicalOr;
    } else {
      error_ = "invalid and expression";
      return 0;
    }
  } else if (expr->IsShiftLeft()) {
    op = spv::Op::OpShiftLeftLogical;
  } else if (expr->IsShiftRight() && lhs_type->is_signed_scalar_or_vector()) {
    // A shift right with a signed LHS is an arithmetic shift.
    op = spv::Op::OpShiftRightArithmetic;
  } else if (expr->IsShiftRight()) {
    op = spv::Op::OpShiftRightLogical;
  } else if (expr->IsSubtract()) {
    op = lhs_is_float_or_vec ? spv::Op::OpFSub : spv::Op::OpISub;
  } else if (expr->IsXor()) {
    op = spv::Op::OpBitwiseXor;
  } else {
    error_ = "unknown binary expression";
    return 0;
  }

  if (!push_function_inst(op, {Operand::Int(type_id), result,
                               Operand::Int(lhs_id), Operand::Int(rhs_id)})) {
    return 0;
  }
  return result_id;
}

bool Builder::GenerateBlockStatement(const ast::BlockStatement* stmt) {
  scope_stack_.Push();
  TINT_DEFER(scope_stack_.Pop());
  return GenerateBlockStatementWithoutScoping(stmt);
}

bool Builder::GenerateBlockStatementWithoutScoping(
    const ast::BlockStatement* stmt) {
  for (auto* block_stmt : stmt->statements) {
    if (!GenerateStatement(block_stmt)) {
      return false;
    }
  }
  return true;
}

uint32_t Builder::GenerateCallExpression(const ast::CallExpression* expr) {
  auto* call = builder_.Sem().Get(expr);
  auto* target = call->Target();
  return Switch(
      target,
      [&](const sem::Function* func) {
        return GenerateFunctionCall(call, func);
      },
      [&](const sem::Builtin* builtin) {
        return GenerateBuiltinCall(call, builtin);
      },
      [&](const sem::TypeConversion*) {
        return GenerateTypeConstructorOrConversion(call, nullptr);
      },
      [&](const sem::TypeConstructor*) {
        return GenerateTypeConstructorOrConversion(call, nullptr);
      },
      [&](Default) {
        TINT_ICE(Writer, builder_.Diagnostics())
            << "unhandled call target: " << target->TypeInfo().name;
        return 0;
      });
}

uint32_t Builder::GenerateFunctionCall(const sem::Call* call,
                                       const sem::Function*) {
  auto* expr = call->Declaration();
  auto* ident = expr->target.name;

  auto type_id = GenerateTypeIfNeeded(call->Type());
  if (type_id == 0) {
    return 0;
  }

  auto result = result_op();
  auto result_id = result.to_i();

  OperandList ops = {Operand::Int(type_id), result};

  auto func_id = func_symbol_to_id_[ident->symbol];
  if (func_id == 0) {
    error_ = "unable to find called function: " +
             builder_.Symbols().NameFor(ident->symbol);
    return 0;
  }
  ops.push_back(Operand::Int(func_id));

  size_t arg_idx = 0;
  for (auto* arg : expr->args) {
    auto id = GenerateExpressionWithLoadIfNeeded(arg);
    if (id == 0) {
      return 0;
    }
    ops.push_back(Operand::Int(id));
    arg_idx++;
  }

  if (!push_function_inst(spv::Op::OpFunctionCall, std::move(ops))) {
    return 0;
  }

  return result_id;
}

uint32_t Builder::GenerateBuiltinCall(const sem::Call* call,
                                      const sem::Builtin* builtin) {
  auto result = result_op();
  auto result_id = result.to_i();

  auto result_type_id = GenerateTypeIfNeeded(builtin->ReturnType());
  if (result_type_id == 0) {
    return 0;
  }

  if (builtin->IsFineDerivative() || builtin->IsCoarseDerivative()) {
    push_capability(SpvCapabilityDerivativeControl);
  }

  if (builtin->IsImageQuery()) {
    push_capability(SpvCapabilityImageQuery);
  }

  if (builtin->IsTexture()) {
    if (!GenerateTextureBuiltin(call, builtin, Operand::Int(result_type_id),
                                result)) {
      return 0;
    }
    return result_id;
  }

  if (builtin->IsBarrier()) {
    if (!GenerateControlBarrierBuiltin(builtin)) {
      return 0;
    }
    return result_id;
  }

  if (builtin->IsAtomic()) {
    if (!GenerateAtomicBuiltin(call, builtin, Operand::Int(result_type_id),
                               result)) {
      return 0;
    }
    return result_id;
  }

  // Generates the SPIR-V ID for the expression for the indexed call argument,
  // and loads it if necessary. Returns 0 on error.
  auto get_arg_as_value_id = [&](size_t i,
                                 bool generate_load = true) -> uint32_t {
    auto* arg = call->Arguments()[i];
    auto* param = builtin->Parameters()[i];
    auto val_id = GenerateExpression(arg->Declaration());
    if (val_id == 0) {
      return 0;
    }

    if (generate_load && !param->Type()->Is<sem::Pointer>()) {
      val_id = GenerateLoadIfNeeded(arg->Type(), val_id);
    }
    return val_id;
  };

  OperandList params = {Operand::Int(result_type_id), result};
  spv::Op op = spv::Op::OpNop;

  // Pushes the arguments for a GlslStd450 extended instruction, and sets op
  // to OpExtInst.
  auto glsl_std450 = [&](uint32_t inst_id) {
    auto set_id = GetGLSLstd450Import();
    params.push_back(Operand::Int(set_id));
    params.push_back(Operand::Int(inst_id));
    op = spv::Op::OpExtInst;
  };

  switch (builtin->Type()) {
    case BuiltinType::kAny:
      if (builtin->Parameters()[0]->Type()->Is<sem::Bool>()) {
        // any(v: bool) just resolves to v.
        return get_arg_as_value_id(0);
      }
      op = spv::Op::OpAny;
      break;
    case BuiltinType::kAll:
      if (builtin->Parameters()[0]->Type()->Is<sem::Bool>()) {
        // all(v: bool) just resolves to v.
        return get_arg_as_value_id(0);
      }
      op = spv::Op::OpAll;
      break;
    case BuiltinType::kArrayLength: {
      auto* address_of =
          call->Arguments()[0]->Declaration()->As<ast::UnaryOpExpression>();
      if (!address_of || address_of->op != ast::UnaryOp::kAddressOf) {
        error_ = "arrayLength() expected pointer to member access, got " +
                 std::string(address_of->TypeInfo().name);
        return 0;
      }
      auto* array_expr = address_of->expr;

      auto* accessor = array_expr->As<ast::MemberAccessorExpression>();
      if (!accessor) {
        error_ =
            "arrayLength() expected pointer to member access, got pointer to " +
            std::string(array_expr->TypeInfo().name);
        return 0;
      }

      auto struct_id = GenerateExpression(accessor->structure);
      if (struct_id == 0) {
        return 0;
      }
      params.push_back(Operand::Int(struct_id));

      auto* type = TypeOf(accessor->structure)->UnwrapRef();
      if (!type->Is<sem::Struct>()) {
        error_ = "invalid type (" + type->FriendlyName(builder_.Symbols()) +
                 ") for runtime array length";
        return 0;
      }
      // Runtime array must be the last member in the structure
      params.push_back(Operand::Int(uint32_t(
          type->As<sem::Struct>()->Declaration()->members.size() - 1)));

      if (!push_function_inst(spv::Op::OpArrayLength, params)) {
        return 0;
      }
      return result_id;
    }
    case BuiltinType::kCountOneBits:
      op = spv::Op::OpBitCount;
      break;
    case BuiltinType::kDot: {
      op = spv::Op::OpDot;
      auto* vec_ty = builtin->Parameters()[0]->Type()->As<sem::Vector>();
      if (vec_ty->type()->is_integer_scalar()) {
        // TODO(crbug.com/tint/1267): OpDot requires floating-point types, but
        // WGSL also supports integer types. SPV_KHR_integer_dot_product adds
        // support for integer vectors. Use it if it is available.
        auto el_ty = Operand::Int(GenerateTypeIfNeeded(vec_ty->type()));
        auto vec_a = Operand::Int(get_arg_as_value_id(0));
        auto vec_b = Operand::Int(get_arg_as_value_id(1));
        if (vec_a.to_i() == 0 || vec_b.to_i() == 0) {
          return 0;
        }

        auto sum = Operand::Int(0);
        for (uint32_t i = 0; i < vec_ty->Width(); i++) {
          auto a = result_op();
          auto b = result_op();
          auto mul = result_op();
          if (!push_function_inst(spv::Op::OpCompositeExtract,
                                  {el_ty, a, vec_a, Operand::Int(i)}) ||
              !push_function_inst(spv::Op::OpCompositeExtract,
                                  {el_ty, b, vec_b, Operand::Int(i)}) ||
              !push_function_inst(spv::Op::OpIMul, {el_ty, mul, a, b})) {
            return 0;
          }
          if (i == 0) {
            sum = mul;
          } else {
            auto prev_sum = sum;
            auto is_last_el = i == (vec_ty->Width() - 1);
            sum = is_last_el ? Operand::Int(result_id) : result_op();
            if (!push_function_inst(spv::Op::OpIAdd,
                                    {el_ty, sum, prev_sum, mul})) {
              return 0;
            }
          }
        }
        return result_id;
      }
      break;
    }
    case BuiltinType::kDpdx:
      op = spv::Op::OpDPdx;
      break;
    case BuiltinType::kDpdxCoarse:
      op = spv::Op::OpDPdxCoarse;
      break;
    case BuiltinType::kDpdxFine:
      op = spv::Op::OpDPdxFine;
      break;
    case BuiltinType::kDpdy:
      op = spv::Op::OpDPdy;
      break;
    case BuiltinType::kDpdyCoarse:
      op = spv::Op::OpDPdyCoarse;
      break;
    case BuiltinType::kDpdyFine:
      op = spv::Op::OpDPdyFine;
      break;
    case BuiltinType::kExtractBits:
      op = builtin->Parameters()[0]->Type()->is_unsigned_scalar_or_vector()
               ? spv::Op::OpBitFieldUExtract
               : spv::Op::OpBitFieldSExtract;
      break;
    case BuiltinType::kFwidth:
      op = spv::Op::OpFwidth;
      break;
    case BuiltinType::kFwidthCoarse:
      op = spv::Op::OpFwidthCoarse;
      break;
    case BuiltinType::kFwidthFine:
      op = spv::Op::OpFwidthFine;
      break;
    case BuiltinType::kInsertBits:
      op = spv::Op::OpBitFieldInsert;
      break;
    case BuiltinType::kMix: {
      auto std450 = Operand::Int(GetGLSLstd450Import());

      auto a_id = get_arg_as_value_id(0);
      auto b_id = get_arg_as_value_id(1);
      auto f_id = get_arg_as_value_id(2);
      if (!a_id || !b_id || !f_id) {
        return 0;
      }

      // If the interpolant is scalar but the objects are vectors, we need to
      // splat the interpolant into a vector of the same size.
      auto* result_vector_type = builtin->ReturnType()->As<sem::Vector>();
      if (result_vector_type && builtin->Parameters()[2]->Type()->is_scalar()) {
        f_id = GenerateSplat(f_id, builtin->Parameters()[0]->Type());
        if (f_id == 0) {
          return 0;
        }
      }

      if (!push_function_inst(spv::Op::OpExtInst,
                              {Operand::Int(result_type_id), result, std450,
                               Operand::Int(GLSLstd450FMix), Operand::Int(a_id),
                               Operand::Int(b_id), Operand::Int(f_id)})) {
        return 0;
      }
      return result_id;
    }
    case BuiltinType::kReverseBits:
      op = spv::Op::OpBitReverse;
      break;
    case BuiltinType::kSelect: {
      // Note: Argument order is different in WGSL and SPIR-V
      auto cond_id = get_arg_as_value_id(2);
      auto true_id = get_arg_as_value_id(1);
      auto false_id = get_arg_as_value_id(0);
      if (!cond_id || !true_id || !false_id) {
        return 0;
      }

      // If the condition is scalar but the objects are vectors, we need to
      // splat the condition into a vector of the same size.
      // TODO(jrprice): If we're targeting SPIR-V 1.4, we don't need to do this.
      auto* result_vector_type = builtin->ReturnType()->As<sem::Vector>();
      if (result_vector_type && builtin->Parameters()[2]->Type()->is_scalar()) {
        auto* bool_vec_ty = builder_.create<sem::Vector>(
            builder_.create<sem::Bool>(), result_vector_type->Width());
        if (!GenerateTypeIfNeeded(bool_vec_ty)) {
          return 0;
        }
        cond_id = GenerateSplat(cond_id, bool_vec_ty);
        if (cond_id == 0) {
          return 0;
        }
      }

      if (!push_function_inst(
              spv::Op::OpSelect,
              {Operand::Int(result_type_id), result, Operand::Int(cond_id),
               Operand::Int(true_id), Operand::Int(false_id)})) {
        return 0;
      }
      return result_id;
    }
    case BuiltinType::kTranspose:
      op = spv::Op::OpTranspose;
      break;
    case BuiltinType::kAbs:
      if (builtin->ReturnType()->is_unsigned_scalar_or_vector()) {
        // abs() only operates on *signed* integers.
        // This is a no-op for unsigned integers.
        return get_arg_as_value_id(0);
      }
      if (builtin->ReturnType()->is_float_scalar_or_vector()) {
        glsl_std450(GLSLstd450FAbs);
      } else {
        glsl_std450(GLSLstd450SAbs);
      }
      break;
    default: {
      auto inst_id = builtin_to_glsl_method(builtin);
      if (inst_id == 0) {
        error_ = "unknown method " + std::string(builtin->str());
        return 0;
      }
      glsl_std450(inst_id);
      break;
    }
  }

  if (op == spv::Op::OpNop) {
    error_ = "unable to determine operator for: " + std::string(builtin->str());
    return 0;
  }

  for (size_t i = 0; i < call->Arguments().size(); i++) {
    if (auto val_id = get_arg_as_value_id(i)) {
      params.emplace_back(Operand::Int(val_id));
    } else {
      return 0;
    }
  }

  if (!push_function_inst(op, params)) {
    return 0;
  }

  return result_id;
}

bool Builder::GenerateTextureBuiltin(const sem::Call* call,
                                     const sem::Builtin* builtin,
                                     Operand result_type,
                                     Operand result_id) {
  using Usage = sem::ParameterUsage;

  auto& signature = builtin->Signature();
  auto& arguments = call->Arguments();

  // Generates the given expression, returning the operand ID
  auto gen = [&](const sem::Expression* expr) {
    const auto val_id = GenerateExpressionWithLoadIfNeeded(expr);
    return Operand::Int(val_id);
  };

  // Returns the argument with the given usage
  auto arg = [&](Usage usage) {
    int idx = signature.IndexOf(usage);
    return (idx >= 0) ? arguments[idx] : nullptr;
  };

  // Generates the argument with the given usage, returning the operand ID
  auto gen_arg = [&](Usage usage) {
    auto* argument = arg(usage);
    if (!argument) {
      TINT_ICE(Writer, builder_.Diagnostics())
          << "missing argument " << static_cast<int>(usage);
    }
    return gen(argument);
  };

  auto* texture = arg(Usage::kTexture);
  if (!texture) {
    TINT_ICE(Writer, builder_.Diagnostics()) << "missing texture argument";
  }

  auto* texture_type = texture->Type()->UnwrapRef()->As<sem::Texture>();

  auto op = spv::Op::OpNop;

  // Custom function to call after the texture-builtin op has been generated.
  std::function<bool()> post_emission = [] { return true; };

  // Populate the spirv_params with common parameters
  OperandList spirv_params;
  spirv_params.reserve(8);  // Enough to fit most parameter lists

  // Extra image operands, appended to spirv_params.
  struct ImageOperand {
    SpvImageOperandsMask mask;
    Operand operand;
  };
  std::vector<ImageOperand> image_operands;
  image_operands.reserve(4);  // Enough to fit most parameter lists

  // Appends `result_type` and `result_id` to `spirv_params`
  auto append_result_type_and_id_to_spirv_params = [&]() {
    spirv_params.emplace_back(std::move(result_type));
    spirv_params.emplace_back(std::move(result_id));
  };

  // Appends a result type and id to `spirv_params`, possibly adding a
  // post_emission step.
  //
  // If the texture is a depth texture, then this function wraps the result of
  // the op with a OpCompositeExtract to evaluate to the first element of the
  // returned vector. This is done as the WGSL texture reading functions for
  // depths return a single float scalar instead of a vector.
  //
  // If the texture is not a depth texture, then this function simply delegates
  // to calling append_result_type_and_id_to_spirv_params().
  auto append_result_type_and_id_to_spirv_params_for_read = [&]() {
    if (texture_type
            ->IsAnyOf<sem::DepthTexture, sem::DepthMultisampledTexture>()) {
      auto* f32 = builder_.create<sem::F32>();
      auto* spirv_result_type = builder_.create<sem::Vector>(f32, 4u);
      auto spirv_result = result_op();
      post_emission = [=] {
        return push_function_inst(
            spv::Op::OpCompositeExtract,
            {result_type, result_id, spirv_result, Operand::Int(0)});
      };
      auto spirv_result_type_id = GenerateTypeIfNeeded(spirv_result_type);
      if (spirv_result_type_id == 0) {
        return false;
      }
      spirv_params.emplace_back(Operand::Int(spirv_result_type_id));
      spirv_params.emplace_back(spirv_result);
      return true;
    }

    append_result_type_and_id_to_spirv_params();
    return true;
  };

  // Appends a result type and id to `spirv_params`, by first swizzling the
  // result of the op with `swizzle`.
  auto append_result_type_and_id_to_spirv_params_swizzled =
      [&](uint32_t spirv_result_width, std::vector<uint32_t> swizzle) {
        if (swizzle.empty()) {
          append_result_type_and_id_to_spirv_params();
        } else {
          // Assign post_emission to swizzle the result of the call to
          // OpImageQuerySize[Lod].
          auto* element_type = ElementTypeOf(call->Type());
          auto spirv_result = result_op();
          auto* spirv_result_type =
              builder_.create<sem::Vector>(element_type, spirv_result_width);
          if (swizzle.size() > 1) {
            post_emission = [=] {
              OperandList operands{
                  result_type,
                  result_id,
                  spirv_result,
                  spirv_result,
              };
              for (auto idx : swizzle) {
                operands.emplace_back(Operand::Int(idx));
              }
              return push_function_inst(spv::Op::OpVectorShuffle, operands);
            };
          } else {
            post_emission = [=] {
              return push_function_inst(spv::Op::OpCompositeExtract,
                                        {result_type, result_id, spirv_result,
                                         Operand::Int(swizzle[0])});
            };
          }
          auto spirv_result_type_id = GenerateTypeIfNeeded(spirv_result_type);
          if (spirv_result_type_id == 0) {
            return false;
          }
          spirv_params.emplace_back(Operand::Int(spirv_result_type_id));
          spirv_params.emplace_back(spirv_result);
        }
        return true;
      };

  auto append_coords_to_spirv_params = [&]() -> bool {
    if (auto* array_index = arg(Usage::kArrayIndex)) {
      // Array index needs to be appended to the coordinates.
      auto* packed = AppendVector(&builder_, arg(Usage::kCoords)->Declaration(),
                                  array_index->Declaration());
      auto param = GenerateExpression(packed->Declaration());
      if (param == 0) {
        return false;
      }
      spirv_params.emplace_back(Operand::Int(param));
    } else {
      spirv_params.emplace_back(gen_arg(Usage::kCoords));  // coordinates
    }
    return true;
  };

  auto append_image_and_coords_to_spirv_params = [&]() -> bool {
    auto sampler_param = gen_arg(Usage::kSampler);
    auto texture_param = gen_arg(Usage::kTexture);
    auto sampled_image =
        GenerateSampledImage(texture_type, texture_param, sampler_param);

    // Populate the spirv_params with the common parameters
    spirv_params.emplace_back(Operand::Int(sampled_image));  // sampled image
    return append_coords_to_spirv_params();
  };

  switch (builtin->Type()) {
    case BuiltinType::kTextureDimensions: {
      // Number of returned elements from OpImageQuerySize[Lod] may not match
      // those of textureDimensions().
      // This might be due to an extra vector scalar describing the number of
      // array elements or textureDimensions() returning a vec3 for cubes
      // when only width / height is returned by OpImageQuerySize[Lod]
      // (see https://github.com/gpuweb/gpuweb/issues/1345).
      // Handle these mismatches by swizzling the returned vector.
      std::vector<uint32_t> swizzle;
      uint32_t spirv_dims = 0;
      switch (texture_type->dim()) {
        case ast::TextureDimension::kNone:
          error_ = "texture dimension is kNone";
          return false;
        case ast::TextureDimension::k1d:
        case ast::TextureDimension::k2d:
        case ast::TextureDimension::k3d:
        case ast::TextureDimension::kCube:
          break;  // No swizzle needed
        case ast::TextureDimension::kCubeArray:
        case ast::TextureDimension::k2dArray:
          swizzle = {0, 1};  // Strip array index
          spirv_dims = 3;    // [width, height, array_count]
          break;
      }

      if (!append_result_type_and_id_to_spirv_params_swizzled(spirv_dims,
                                                              swizzle)) {
        return false;
      }

      spirv_params.emplace_back(gen_arg(Usage::kTexture));
      if (texture_type->IsAnyOf<sem::MultisampledTexture,       //
                                sem::DepthMultisampledTexture,  //
                                sem::StorageTexture>()) {
        op = spv::Op::OpImageQuerySize;
      } else if (auto* level = arg(Usage::kLevel)) {
        op = spv::Op::OpImageQuerySizeLod;
        spirv_params.emplace_back(gen(level));
      } else {
        ast::SintLiteralExpression i32_0(ProgramID(), Source{}, 0);
        op = spv::Op::OpImageQuerySizeLod;
        spirv_params.emplace_back(
            Operand::Int(GenerateLiteralIfNeeded(nullptr, &i32_0)));
      }
      break;
    }
    case BuiltinType::kTextureNumLayers: {
      uint32_t spirv_dims = 0;
      switch (texture_type->dim()) {
        default:
          error_ = "texture is not arrayed";
          return false;
        case ast::TextureDimension::k2dArray:
        case ast::TextureDimension::kCubeArray:
          spirv_dims = 3;
          break;
      }

      // OpImageQuerySize[Lod] packs the array count as the last element of the
      // returned vector. Extract this.
      if (!append_result_type_and_id_to_spirv_params_swizzled(
              spirv_dims, {spirv_dims - 1})) {
        return false;
      }

      spirv_params.emplace_back(gen_arg(Usage::kTexture));

      if (texture_type->Is<sem::MultisampledTexture>() ||
          texture_type->Is<sem::StorageTexture>()) {
        op = spv::Op::OpImageQuerySize;
      } else {
        ast::SintLiteralExpression i32_0(ProgramID(), Source{}, 0);
        op = spv::Op::OpImageQuerySizeLod;
        spirv_params.emplace_back(
            Operand::Int(GenerateLiteralIfNeeded(nullptr, &i32_0)));
      }
      break;
    }
    case BuiltinType::kTextureNumLevels: {
      op = spv::Op::OpImageQueryLevels;
      append_result_type_and_id_to_spirv_params();
      spirv_params.emplace_back(gen_arg(Usage::kTexture));
      break;
    }
    case BuiltinType::kTextureNumSamples: {
      op = spv::Op::OpImageQuerySamples;
      append_result_type_and_id_to_spirv_params();
      spirv_params.emplace_back(gen_arg(Usage::kTexture));
      break;
    }
    case BuiltinType::kTextureLoad: {
      op = texture_type->Is<sem::StorageTexture>() ? spv::Op::OpImageRead
                                                   : spv::Op::OpImageFetch;
      append_result_type_and_id_to_spirv_params_for_read();
      spirv_params.emplace_back(gen_arg(Usage::kTexture));
      if (!append_coords_to_spirv_params()) {
        return false;
      }

      if (auto* level = arg(Usage::kLevel)) {
        image_operands.emplace_back(
            ImageOperand{SpvImageOperandsLodMask, gen(level)});
      }

      if (auto* sample_index = arg(Usage::kSampleIndex)) {
        image_operands.emplace_back(
            ImageOperand{SpvImageOperandsSampleMask, gen(sample_index)});
      }

      break;
    }
    case BuiltinType::kTextureStore: {
      op = spv::Op::OpImageWrite;
      spirv_params.emplace_back(gen_arg(Usage::kTexture));
      if (!append_coords_to_spirv_params()) {
        return false;
      }
      spirv_params.emplace_back(gen_arg(Usage::kValue));
      break;
    }
    case BuiltinType::kTextureGather: {
      op = spv::Op::OpImageGather;
      append_result_type_and_id_to_spirv_params();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      if (signature.IndexOf(Usage::kComponent) < 0) {
        spirv_params.emplace_back(
            Operand::Int(GenerateConstantIfNeeded(ScalarConstant::I32(0))));
      } else {
        spirv_params.emplace_back(gen_arg(Usage::kComponent));
      }
      break;
    }
    case BuiltinType::kTextureGatherCompare: {
      op = spv::Op::OpImageDrefGather;
      append_result_type_and_id_to_spirv_params();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      spirv_params.emplace_back(gen_arg(Usage::kDepthRef));
      break;
    }
    case BuiltinType::kTextureSample: {
      op = spv::Op::OpImageSampleImplicitLod;
      append_result_type_and_id_to_spirv_params_for_read();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      break;
    }
    case BuiltinType::kTextureSampleBias: {
      op = spv::Op::OpImageSampleImplicitLod;
      append_result_type_and_id_to_spirv_params_for_read();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      image_operands.emplace_back(
          ImageOperand{SpvImageOperandsBiasMask, gen_arg(Usage::kBias)});
      break;
    }
    case BuiltinType::kTextureSampleLevel: {
      op = spv::Op::OpImageSampleExplicitLod;
      append_result_type_and_id_to_spirv_params_for_read();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      auto level = Operand::Int(0);
      if (arg(Usage::kLevel)->Type()->UnwrapRef()->Is<sem::I32>()) {
        // Depth textures have i32 parameters for the level, but SPIR-V expects
        // F32. Cast.
        auto f32_type_id = GenerateTypeIfNeeded(builder_.create<sem::F32>());
        if (f32_type_id == 0) {
          return 0;
        }
        level = result_op();
        if (!push_function_inst(
                spv::Op::OpConvertSToF,
                {Operand::Int(f32_type_id), level, gen_arg(Usage::kLevel)})) {
          return 0;
        }
      } else {
        level = gen_arg(Usage::kLevel);
      }
      image_operands.emplace_back(ImageOperand{SpvImageOperandsLodMask, level});
      break;
    }
    case BuiltinType::kTextureSampleGrad: {
      op = spv::Op::OpImageSampleExplicitLod;
      append_result_type_and_id_to_spirv_params_for_read();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      image_operands.emplace_back(
          ImageOperand{SpvImageOperandsGradMask, gen_arg(Usage::kDdx)});
      image_operands.emplace_back(
          ImageOperand{SpvImageOperandsGradMask, gen_arg(Usage::kDdy)});
      break;
    }
    case BuiltinType::kTextureSampleCompare: {
      op = spv::Op::OpImageSampleDrefImplicitLod;
      append_result_type_and_id_to_spirv_params();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      spirv_params.emplace_back(gen_arg(Usage::kDepthRef));
      break;
    }
    case BuiltinType::kTextureSampleCompareLevel: {
      op = spv::Op::OpImageSampleDrefExplicitLod;
      append_result_type_and_id_to_spirv_params();
      if (!append_image_and_coords_to_spirv_params()) {
        return false;
      }
      spirv_params.emplace_back(gen_arg(Usage::kDepthRef));

      ast::FloatLiteralExpression float_0(ProgramID(), Source{}, 0.0);
      image_operands.emplace_back(ImageOperand{
          SpvImageOperandsLodMask,
          Operand::Int(GenerateLiteralIfNeeded(nullptr, &float_0))});
      break;
    }
    default:
      TINT_UNREACHABLE(Writer, builder_.Diagnostics());
      return false;
  }

  if (auto* offset = arg(Usage::kOffset)) {
    image_operands.emplace_back(
        ImageOperand{SpvImageOperandsConstOffsetMask, gen(offset)});
  }

  if (!image_operands.empty()) {
    std::sort(image_operands.begin(), image_operands.end(),
              [](auto& a, auto& b) { return a.mask < b.mask; });
    uint32_t mask = 0;
    for (auto& image_operand : image_operands) {
      mask |= image_operand.mask;
    }
    spirv_params.emplace_back(Operand::Int(mask));
    for (auto& image_operand : image_operands) {
      spirv_params.emplace_back(image_operand.operand);
    }
  }

  if (op == spv::Op::OpNop) {
    error_ = "unable to determine operator for: " + std::string(builtin->str());
    return false;
  }

  if (!push_function_inst(op, spirv_params)) {
    return false;
  }

  return post_emission();
}

bool Builder::GenerateControlBarrierBuiltin(const sem::Builtin* builtin) {
  auto const op = spv::Op::OpControlBarrier;
  uint32_t execution = 0;
  uint32_t memory = 0;
  uint32_t semantics = 0;

  // TODO(crbug.com/tint/661): Combine sequential barriers to a single
  // instruction.
  if (builtin->Type() == sem::BuiltinType::kWorkgroupBarrier) {
    execution = static_cast<uint32_t>(spv::Scope::Workgroup);
    memory = static_cast<uint32_t>(spv::Scope::Workgroup);
    semantics =
        static_cast<uint32_t>(spv::MemorySemanticsMask::AcquireRelease) |
        static_cast<uint32_t>(spv::MemorySemanticsMask::WorkgroupMemory);
  } else if (builtin->Type() == sem::BuiltinType::kStorageBarrier) {
    execution = static_cast<uint32_t>(spv::Scope::Workgroup);
    memory = static_cast<uint32_t>(spv::Scope::Workgroup);
    semantics =
        static_cast<uint32_t>(spv::MemorySemanticsMask::AcquireRelease) |
        static_cast<uint32_t>(spv::MemorySemanticsMask::UniformMemory);
  } else {
    error_ = "unexpected barrier builtin type ";
    error_ += sem::str(builtin->Type());
    return false;
  }

  auto execution_id = GenerateConstantIfNeeded(ScalarConstant::U32(execution));
  auto memory_id = GenerateConstantIfNeeded(ScalarConstant::U32(memory));
  auto semantics_id = GenerateConstantIfNeeded(ScalarConstant::U32(semantics));
  if (execution_id == 0 || memory_id == 0 || semantics_id == 0) {
    return false;
  }

  return push_function_inst(op, {
                                    Operand::Int(execution_id),
                                    Operand::Int(memory_id),
                                    Operand::Int(semantics_id),
                                });
}

bool Builder::GenerateAtomicBuiltin(const sem::Call* call,
                                    const sem::Builtin* builtin,
                                    Operand result_type,
                                    Operand result_id) {
  auto is_value_signed = [&] {
    return builtin->Parameters()[1]->Type()->Is<sem::I32>();
  };

  auto storage_class =
      builtin->Parameters()[0]->Type()->As<sem::Pointer>()->StorageClass();

  uint32_t memory_id = 0;
  switch (
      builtin->Parameters()[0]->Type()->As<sem::Pointer>()->StorageClass()) {
    case ast::StorageClass::kWorkgroup:
      memory_id = GenerateConstantIfNeeded(
          ScalarConstant::U32(static_cast<uint32_t>(spv::Scope::Workgroup)));
      break;
    case ast::StorageClass::kStorage:
      memory_id = GenerateConstantIfNeeded(
          ScalarConstant::U32(static_cast<uint32_t>(spv::Scope::Device)));
      break;
    default:
      TINT_UNREACHABLE(Writer, builder_.Diagnostics())
          << "unhandled atomic storage class " << storage_class;
      return false;
  }
  if (memory_id == 0) {
    return false;
  }

  uint32_t semantics_id = GenerateConstantIfNeeded(ScalarConstant::U32(
      static_cast<uint32_t>(spv::MemorySemanticsMask::MaskNone)));
  if (semantics_id == 0) {
    return false;
  }

  uint32_t pointer_id = GenerateExpression(call->Arguments()[0]->Declaration());
  if (pointer_id == 0) {
    return false;
  }

  uint32_t value_id = 0;
  if (call->Arguments().size() > 1) {
    value_id = GenerateExpressionWithLoadIfNeeded(call->Arguments().back());
    if (value_id == 0) {
      return false;
    }
  }

  Operand pointer = Operand::Int(pointer_id);
  Operand value = Operand::Int(value_id);
  Operand memory = Operand::Int(memory_id);
  Operand semantics = Operand::Int(semantics_id);

  switch (builtin->Type()) {
    case sem::BuiltinType::kAtomicLoad:
      return push_function_inst(spv::Op::OpAtomicLoad, {
                                                           result_type,
                                                           result_id,
                                                           pointer,
                                                           memory,
                                                           semantics,
                                                       });
    case sem::BuiltinType::kAtomicStore:
      return push_function_inst(spv::Op::OpAtomicStore, {
                                                            pointer,
                                                            memory,
                                                            semantics,
                                                            value,
                                                        });
    case sem::BuiltinType::kAtomicAdd:
      return push_function_inst(spv::Op::OpAtomicIAdd, {
                                                           result_type,
                                                           result_id,
                                                           pointer,
                                                           memory,
                                                           semantics,
                                                           value,
                                                       });
    case sem::BuiltinType::kAtomicSub:
      return push_function_inst(spv::Op::OpAtomicISub, {
                                                           result_type,
                                                           result_id,
                                                           pointer,
                                                           memory,
                                                           semantics,
                                                           value,
                                                       });
    case sem::BuiltinType::kAtomicMax:
      return push_function_inst(
          is_value_signed() ? spv::Op::OpAtomicSMax : spv::Op::OpAtomicUMax,
          {
              result_type,
              result_id,
              pointer,
              memory,
              semantics,
              value,
          });
    case sem::BuiltinType::kAtomicMin:
      return push_function_inst(
          is_value_signed() ? spv::Op::OpAtomicSMin : spv::Op::OpAtomicUMin,
          {
              result_type,
              result_id,
              pointer,
              memory,
              semantics,
              value,
          });
    case sem::BuiltinType::kAtomicAnd:
      return push_function_inst(spv::Op::OpAtomicAnd, {
                                                          result_type,
                                                          result_id,
                                                          pointer,
                                                          memory,
                                                          semantics,
                                                          value,
                                                      });
    case sem::BuiltinType::kAtomicOr:
      return push_function_inst(spv::Op::OpAtomicOr, {
                                                         result_type,
                                                         result_id,
                                                         pointer,
                                                         memory,
                                                         semantics,
                                                         value,
                                                     });
    case sem::BuiltinType::kAtomicXor:
      return push_function_inst(spv::Op::OpAtomicXor, {
                                                          result_type,
                                                          result_id,
                                                          pointer,
                                                          memory,
                                                          semantics,
                                                          value,
                                                      });
    case sem::BuiltinType::kAtomicExchange:
      return push_function_inst(spv::Op::OpAtomicExchange, {
                                                               result_type,
                                                               result_id,
                                                               pointer,
                                                               memory,
                                                               semantics,
                                                               value,
                                                           });
    case sem::BuiltinType::kAtomicCompareExchangeWeak: {
      auto comparator = GenerateExpression(call->Arguments()[1]->Declaration());
      if (comparator == 0) {
        return false;
      }

      auto* value_sem_type = TypeOf(call->Arguments()[2]->Declaration());

      auto value_type = GenerateTypeIfNeeded(value_sem_type);
      if (value_type == 0) {
        return false;
      }

      auto* bool_sem_ty = builder_.create<sem::Bool>();
      auto bool_type = GenerateTypeIfNeeded(bool_sem_ty);
      if (bool_type == 0) {
        return false;
      }

      // original_value := OpAtomicCompareExchange(pointer, memory, semantics,
      //                                           semantics, value, comparator)
      auto original_value = result_op();
      if (!push_function_inst(spv::Op::OpAtomicCompareExchange,
                              {
                                  Operand::Int(value_type),
                                  original_value,
                                  pointer,
                                  memory,
                                  semantics,
                                  semantics,
                                  value,
                                  Operand::Int(comparator),
                              })) {
        return false;
      }

      // values_equal := original_value == value
      auto values_equal = result_op();
      if (!push_function_inst(spv::Op::OpIEqual, {
                                                     Operand::Int(bool_type),
                                                     values_equal,
                                                     original_value,
                                                     value,
                                                 })) {
        return false;
      }

      // zero := T(0)
      // one := T(1)
      uint32_t zero = 0;
      uint32_t one = 0;
      if (value_sem_type->Is<sem::I32>()) {
        zero = GenerateConstantIfNeeded(ScalarConstant::I32(0u));
        one = GenerateConstantIfNeeded(ScalarConstant::I32(1u));
      } else if (value_sem_type->Is<sem::U32>()) {
        zero = GenerateConstantIfNeeded(ScalarConstant::U32(0u));
        one = GenerateConstantIfNeeded(ScalarConstant::U32(1u));
      } else {
        TINT_UNREACHABLE(Writer, builder_.Diagnostics())
            << "unsupported atomic type " << value_sem_type->TypeInfo().name;
      }
      if (zero == 0 || one == 0) {
        return false;
      }

      // xchg_success := values_equal ? one : zero
      auto xchg_success = result_op();
      if (!push_function_inst(spv::Op::OpSelect, {
                                                     Operand::Int(value_type),
                                                     xchg_success,
                                                     values_equal,
                                                     Operand::Int(one),
                                                     Operand::Int(zero),
                                                 })) {
        return false;
      }

      // result := vec2<T>(original_value, xchg_success)
      return push_function_inst(spv::Op::OpCompositeConstruct,
                                {
                                    result_type,
                                    result_id,
                                    original_value,
                                    xchg_success,
                                });
    }
    default:
      TINT_UNREACHABLE(Writer, builder_.Diagnostics())
          << "unhandled atomic builtin " << builtin->Type();
      return false;
  }
}

uint32_t Builder::GenerateSampledImage(const sem::Type* texture_type,
                                       Operand texture_operand,
                                       Operand sampler_operand) {
  // DepthTexture is always declared as SampledTexture.
  // The Vulkan spec says: The "Depth" operand of OpTypeImage is ignored.
  // In SPIRV, 0 means not depth, 1 means depth, and 2 means unknown.
  // Using anything other than 0 is problematic on various Vulkan drivers.
  if (auto* depthTextureType = texture_type->As<sem::DepthTexture>()) {
    texture_type = builder_.create<sem::SampledTexture>(
        depthTextureType->dim(), builder_.create<sem::F32>());
  }

  uint32_t sampled_image_type_id = utils::GetOrCreate(
      texture_type_to_sampled_image_type_id_, texture_type, [&] {
        // We need to create the sampled image type and cache the result.
        auto sampled_image_type = result_op();
        auto texture_type_id = GenerateTypeIfNeeded(texture_type);
        push_type(spv::Op::OpTypeSampledImage,
                  {sampled_image_type, Operand::Int(texture_type_id)});
        return sampled_image_type.to_i();
      });

  auto sampled_image = result_op();
  if (!push_function_inst(spv::Op::OpSampledImage,
                          {Operand::Int(sampled_image_type_id), sampled_image,
                           texture_operand, sampler_operand})) {
    return 0;
  }

  return sampled_image.to_i();
}

uint32_t Builder::GenerateBitcastExpression(
    const ast::BitcastExpression* expr) {
  auto result = result_op();
  auto result_id = result.to_i();

  auto result_type_id = GenerateTypeIfNeeded(TypeOf(expr));
  if (result_type_id == 0) {
    return 0;
  }

  auto val_id = GenerateExpressionWithLoadIfNeeded(expr->expr);
  if (val_id == 0) {
    return 0;
  }

  // Bitcast does not allow same types, just emit a CopyObject
  auto* to_type = TypeOf(expr)->UnwrapRef();
  auto* from_type = TypeOf(expr->expr)->UnwrapRef();
  if (to_type == from_type) {
    if (!push_function_inst(
            spv::Op::OpCopyObject,
            {Operand::Int(result_type_id), result, Operand::Int(val_id)})) {
      return 0;
    }
    return result_id;
  }

  if (!push_function_inst(spv::Op::OpBitcast, {Operand::Int(result_type_id),
                                               result, Operand::Int(val_id)})) {
    return 0;
  }

  return result_id;
}

bool Builder::GenerateConditionalBlock(
    const ast::Expression* cond,
    const ast::BlockStatement* true_body,
    size_t cur_else_idx,
    const ast::ElseStatementList& else_stmts) {
  auto cond_id = GenerateExpressionWithLoadIfNeeded(cond);
  if (cond_id == 0) {
    return false;
  }

  auto merge_block = result_op();
  auto merge_block_id = merge_block.to_i();

  if (!push_function_inst(spv::Op::OpSelectionMerge,
                          {Operand::Int(merge_block_id),
                           Operand::Int(SpvSelectionControlMaskNone)})) {
    return false;
  }

  auto true_block = result_op();
  auto true_block_id = true_block.to_i();

  // if there are no more else statements we branch on false to the merge
  // block otherwise we branch to the false block
  auto false_block_id =
      cur_else_idx < else_stmts.size() ? next_id() : merge_block_id;

  if (!push_function_inst(spv::Op::OpBranchConditional,
                          {Operand::Int(cond_id), Operand::Int(true_block_id),
                           Operand::Int(false_block_id)})) {
    return false;
  }

  // Output true block
  if (!GenerateLabel(true_block_id)) {
    return false;
  }
  if (!GenerateBlockStatement(true_body)) {
    return false;
  }
  // We only branch if the last element of the body didn't already branch.
  if (InsideBasicBlock()) {
    if (!push_function_inst(spv::Op::OpBranch,
                            {Operand::Int(merge_block_id)})) {
      return false;
    }
  }

  // Start the false block if needed
  if (false_block_id != merge_block_id) {
    if (!GenerateLabel(false_block_id)) {
      return false;
    }

    auto* else_stmt = else_stmts[cur_else_idx];
    // Handle the else case by just outputting the statements.
    if (!else_stmt->condition) {
      if (!GenerateBlockStatement(else_stmt->body)) {
        return false;
      }
    } else {
      if (!GenerateConditionalBlock(else_stmt->condition, else_stmt->body,
                                    cur_else_idx + 1, else_stmts)) {
        return false;
      }
    }
    if (InsideBasicBlock()) {
      if (!push_function_inst(spv::Op::OpBranch,
                              {Operand::Int(merge_block_id)})) {
        return false;
      }
    }
  }

  // Output the merge block
  return GenerateLabel(merge_block_id);
}

bool Builder::GenerateIfStatement(const ast::IfStatement* stmt) {
  if (!continuing_stack_.empty() &&
      stmt == continuing_stack_.back().last_statement->As<ast::IfStatement>()) {
    const ContinuingInfo& ci = continuing_stack_.back();
    // Match one of two patterns: the break-if and break-unless patterns.
    //
    // The break-if pattern:
    //  continuing { ...
    //    if (cond) { break; }
    //  }
    //
    // The break-unless pattern:
    //  continuing { ...
    //    if (cond) {} else {break;}
    //  }
    auto is_just_a_break = [](const ast::BlockStatement* block) {
      return block && (block->statements.size() == 1) &&
             block->Last()->Is<ast::BreakStatement>();
    };
    if (is_just_a_break(stmt->body) && stmt->else_statements.empty()) {
      // It's a break-if.
      TINT_ASSERT(Writer, !backedge_stack_.empty());
      const auto cond_id = GenerateExpressionWithLoadIfNeeded(stmt->condition);
      if (!cond_id) {
        return false;
      }
      backedge_stack_.back() =
          Backedge(spv::Op::OpBranchConditional,
                   {Operand::Int(cond_id), Operand::Int(ci.break_target_id),
                    Operand::Int(ci.loop_header_id)});
      return true;
    } else if (stmt->body->Empty()) {
      const auto& es = stmt->else_statements;
      if (es.size() == 1 && !es.back()->condition &&
          is_just_a_break(es.back()->body)) {
        // It's a break-unless.
        TINT_ASSERT(Writer, !backedge_stack_.empty());
        const auto cond_id =
            GenerateExpressionWithLoadIfNeeded(stmt->condition);
        if (!cond_id) {
          return false;
        }
        backedge_stack_.back() =
            Backedge(spv::Op::OpBranchConditional,
                     {Operand::Int(cond_id), Operand::Int(ci.loop_header_id),
                      Operand::Int(ci.break_target_id)});
        return true;
      }
    }
  }

  if (!GenerateConditionalBlock(stmt->condition, stmt->body, 0,
                                stmt->else_statements)) {
    return false;
  }
  return true;
}

bool Builder::GenerateSwitchStatement(const ast::SwitchStatement* stmt) {
  auto merge_block = result_op();
  auto merge_block_id = merge_block.to_i();

  merge_stack_.push_back(merge_block_id);

  auto cond_id = GenerateExpressionWithLoadIfNeeded(stmt->condition);
  if (cond_id == 0) {
    return false;
  }

  auto default_block = result_op();
  auto default_block_id = default_block.to_i();

  OperandList params = {Operand::Int(cond_id), Operand::Int(default_block_id)};

  std::vector<uint32_t> case_ids;
  for (const auto* item : stmt->body) {
    if (item->IsDefault()) {
      case_ids.push_back(default_block_id);
      continue;
    }

    auto block = result_op();
    auto block_id = block.to_i();

    case_ids.push_back(block_id);
    for (auto* selector : item->selectors) {
      auto* int_literal = selector->As<ast::IntLiteralExpression>();
      if (!int_literal) {
        error_ = "expected integer literal for switch case label";
        return false;
      }

      params.push_back(Operand::Int(int_literal->ValueAsU32()));
      params.push_back(Operand::Int(block_id));
    }
  }

  if (!push_function_inst(spv::Op::OpSelectionMerge,
                          {Operand::Int(merge_block_id),
                           Operand::Int(SpvSelectionControlMaskNone)})) {
    return false;
  }
  if (!push_function_inst(spv::Op::OpSwitch, params)) {
    return false;
  }

  bool generated_default = false;
  auto& body = stmt->body;
  // We output the case statements in order they were entered in the original
  // source. Each fallthrough goes to the next case entry, so is a forward
  // branch, otherwise the branch is to the merge block which comes after
  // the switch statement.
  for (uint32_t i = 0; i < body.size(); i++) {
    auto* item = body[i];

    if (item->IsDefault()) {
      generated_default = true;
    }

    if (!GenerateLabel(case_ids[i])) {
      return false;
    }
    if (!GenerateBlockStatement(item->body)) {
      return false;
    }

    if (LastIsFallthrough(item->body)) {
      if (i == (body.size() - 1)) {
        // This case is caught by Resolver validation
        TINT_UNREACHABLE(Writer, builder_.Diagnostics());
        return false;
      }
      if (!push_function_inst(spv::Op::OpBranch,
                              {Operand::Int(case_ids[i + 1])})) {
        return false;
      }
    } else if (InsideBasicBlock()) {
      if (!push_function_inst(spv::Op::OpBranch,
                              {Operand::Int(merge_block_id)})) {
        return false;
      }
    }
  }

  if (!generated_default) {
    if (!GenerateLabel(default_block_id)) {
      return false;
    }
    if (!push_function_inst(spv::Op::OpBranch,
                            {Operand::Int(merge_block_id)})) {
      return false;
    }
  }

  merge_stack_.pop_back();

  return GenerateLabel(merge_block_id);
}

bool Builder::GenerateReturnStatement(const ast::ReturnStatement* stmt) {
  if (stmt->value) {
    auto val_id = GenerateExpressionWithLoadIfNeeded(stmt->value);
    if (val_id == 0) {
      return false;
    }
    if (!push_function_inst(spv::Op::OpReturnValue, {Operand::Int(val_id)})) {
      return false;
    }
  } else {
    if (!push_function_inst(spv::Op::OpReturn, {})) {
      return false;
    }
  }

  return true;
}

bool Builder::GenerateLoopStatement(const ast::LoopStatement* stmt) {
  auto loop_header = result_op();
  auto loop_header_id = loop_header.to_i();
  if (!push_function_inst(spv::Op::OpBranch, {Operand::Int(loop_header_id)})) {
    return false;
  }
  if (!GenerateLabel(loop_header_id)) {
    return false;
  }

  auto merge_block = result_op();
  auto merge_block_id = merge_block.to_i();
  auto continue_block = result_op();
  auto continue_block_id = continue_block.to_i();

  auto body_block = result_op();
  auto body_block_id = body_block.to_i();

  if (!push_function_inst(
          spv::Op::OpLoopMerge,
          {Operand::Int(merge_block_id), Operand::Int(continue_block_id),
           Operand::Int(SpvLoopControlMaskNone)})) {
    return false;
  }

  continue_stack_.push_back(continue_block_id);
  merge_stack_.push_back(merge_block_id);

  // Usually, the backedge is a simple branch.  This will be modified if the
  // backedge block in the continuing construct has an exiting edge.
  backedge_stack_.emplace_back(spv::Op::OpBranch,
                               OperandList{Operand::Int(loop_header_id)});

  if (!push_function_inst(spv::Op::OpBranch, {Operand::Int(body_block_id)})) {
    return false;
  }
  if (!GenerateLabel(body_block_id)) {
    return false;
  }

  // We need variables from the body to be visible in the continuing block, so
  // manage scope outside of GenerateBlockStatement.
  {
    scope_stack_.Push();
    TINT_DEFER(scope_stack_.Pop());

    if (!GenerateBlockStatementWithoutScoping(stmt->body)) {
      return false;
    }

    // We only branch if the last element of the body didn't already branch.
    if (InsideBasicBlock()) {
      if (!push_function_inst(spv::Op::OpBranch,
                              {Operand::Int(continue_block_id)})) {
        return false;
      }
    }

    if (!GenerateLabel(continue_block_id)) {
      return false;
    }
    if (stmt->continuing && !stmt->continuing->Empty()) {
      continuing_stack_.emplace_back(stmt->continuing->Last(), loop_header_id,
                                     merge_block_id);
      if (!GenerateBlockStatementWithoutScoping(stmt->continuing)) {
        return false;
      }
      continuing_stack_.pop_back();
    }
  }

  // Generate the backedge.
  TINT_ASSERT(Writer, !backedge_stack_.empty());
  const Backedge& backedge = backedge_stack_.back();
  if (!push_function_inst(backedge.opcode, backedge.operands)) {
    return false;
  }
  backedge_stack_.pop_back();

  merge_stack_.pop_back();
  continue_stack_.pop_back();

  return GenerateLabel(merge_block_id);
}

bool Builder::GenerateStatement(const ast::Statement* stmt) {
  return Switch(
      stmt,
      [&](const ast::AssignmentStatement* a) {
        return GenerateAssignStatement(a);
      },
      [&](const ast::BlockStatement* b) {  //
        return GenerateBlockStatement(b);
      },
      [&](const ast::BreakStatement* b) {  //
        return GenerateBreakStatement(b);
      },
      [&](const ast::CallStatement* c) {
        return GenerateCallExpression(c->expr) != 0;
      },
      [&](const ast::ContinueStatement* c) {
        return GenerateContinueStatement(c);
      },
      [&](const ast::DiscardStatement* d) {
        return GenerateDiscardStatement(d);
      },
      [&](const ast::FallthroughStatement*) {
        // Do nothing here, the fallthrough gets handled by the switch code.
        return true;
      },
      [&](const ast::IfStatement* i) {  //
        return GenerateIfStatement(i);
      },
      [&](const ast::LoopStatement* l) {  //
        return GenerateLoopStatement(l);
      },
      [&](const ast::ReturnStatement* r) {  //
        return GenerateReturnStatement(r);
      },
      [&](const ast::SwitchStatement* s) {  //
        return GenerateSwitchStatement(s);
      },
      [&](const ast::VariableDeclStatement* v) {
        return GenerateVariableDeclStatement(v);
      },
      [&](Default) {
        error_ = "Unknown statement: " + std::string(stmt->TypeInfo().name);
        return false;
      });
}

bool Builder::GenerateVariableDeclStatement(
    const ast::VariableDeclStatement* stmt) {
  return GenerateFunctionVariable(stmt->variable);
}

uint32_t Builder::GenerateTypeIfNeeded(const sem::Type* type) {
  if (type == nullptr) {
    error_ = "attempting to generate type from null type";
    return 0;
  }

  // Atomics are a type in WGSL, but aren't a distinct type in SPIR-V.
  // Just emit the type inside the atomic.
  if (auto* atomic = type->As<sem::Atomic>()) {
    return GenerateTypeIfNeeded(atomic->Type());
  }

  // DepthTexture is always declared as SampledTexture.
  // The Vulkan spec says: The "Depth" operand of OpTypeImage is ignored.
  // In SPIRV, 0 means not depth, 1 means depth, and 2 means unknown.
  // Using anything other than 0 is problematic on various Vulkan drivers.
  if (auto* depthTextureType = type->As<sem::DepthTexture>()) {
    type = builder_.create<sem::SampledTexture>(depthTextureType->dim(),
                                                builder_.create<sem::F32>());
  } else if (auto* multisampledDepthTextureType =
                 type->As<sem::DepthMultisampledTexture>()) {
    type = builder_.create<sem::MultisampledTexture>(
        multisampledDepthTextureType->dim(), builder_.create<sem::F32>());
  }

  // Pointers and references with differing accesses should not result in a
  // different SPIR-V types, so we explicitly ignore the access.
  // Pointers and References both map to a SPIR-V pointer type.
  // Transform a Reference to a Pointer to prevent these having duplicated
  // definitions in the generated SPIR-V. Note that nested pointers and
  // references are not legal in WGSL, so only considering the top-level type is
  // fine.
  if (auto* ptr = type->As<sem::Pointer>()) {
    type = builder_.create<sem::Pointer>(ptr->StoreType(), ptr->StorageClass(),
                                         ast::kReadWrite);
  } else if (auto* ref = type->As<sem::Reference>()) {
    type = builder_.create<sem::Pointer>(ref->StoreType(), ref->StorageClass(),
                                         ast::kReadWrite);
  }

  return utils::GetOrCreate(type_to_id_, type, [&]() -> uint32_t {
    auto result = result_op();
    auto id = result.to_i();
    bool ok = Switch(
        type,
        [&](const sem::Array* arr) {  //
          return GenerateArrayType(arr, result);
        },
        [&](const sem::Bool*) {
          push_type(spv::Op::OpTypeBool, {result});
          return true;
        },
        [&](const sem::F32*) {
          push_type(spv::Op::OpTypeFloat, {result, Operand::Int(32)});
          return true;
        },
        [&](const sem::I32*) {
          push_type(spv::Op::OpTypeInt,
                    {result, Operand::Int(32), Operand::Int(1)});
          return true;
        },
        [&](const sem::Matrix* mat) {  //
          return GenerateMatrixType(mat, result);
        },
        [&](const sem::Pointer* ptr) {  //
          return GeneratePointerType(ptr, result);
        },
        [&](const sem::Reference* ref) {  //
          return GenerateReferenceType(ref, result);
        },
        [&](const sem::Struct* str) {  //
          return GenerateStructType(str, result);
        },
        [&](const sem::U32*) {
          push_type(spv::Op::OpTypeInt,
                    {result, Operand::Int(32), Operand::Int(0)});
          return true;
        },
        [&](const sem::Vector* vec) {  //
          return GenerateVectorType(vec, result);
        },
        [&](const sem::Void*) {
          push_type(spv::Op::OpTypeVoid, {result});
          return true;
        },
        [&](const sem::StorageTexture* tex) {
          if (!GenerateTextureType(tex, result)) {
            return false;
          }

          // Register all three access types of StorageTexture names. In
          // SPIR-V, we must output a single type, while the variable is
          // annotated with the access type. Doing this ensures we de-dupe.
          type_to_id_[builder_.create<sem::StorageTexture>(
              tex->dim(), tex->texel_format(), ast::Access::kRead,
              tex->type())] = id;
          type_to_id_[builder_.create<sem::StorageTexture>(
              tex->dim(), tex->texel_format(), ast::Access::kWrite,
              tex->type())] = id;
          type_to_id_[builder_.create<sem::StorageTexture>(
              tex->dim(), tex->texel_format(), ast::Access::kReadWrite,
              tex->type())] = id;
          return true;
        },
        [&](const sem::Texture* tex) {
          return GenerateTextureType(tex, result);
        },
        [&](const sem::Sampler* s) {
          push_type(spv::Op::OpTypeSampler, {result});

          // Register both of the sampler type names. In SPIR-V they're the same
          // sampler type, so we need to match that when we do the dedup check.
          if (s->kind() == ast::SamplerKind::kSampler) {
            type_to_id_[builder_.create<sem::Sampler>(
                ast::SamplerKind::kComparisonSampler)] = id;
          } else {
            type_to_id_[builder_.create<sem::Sampler>(
                ast::SamplerKind::kSampler)] = id;
          }
          return true;
        },
        [&](Default) {
          error_ = "unable to convert type: " +
                   type->FriendlyName(builder_.Symbols());
          return false;
        });

    if (!ok) {
      return 0;
    }

    return id;
  });
}

bool Builder::GenerateTextureType(const sem::Texture* texture,
                                  const Operand& result) {
  if (texture->Is<sem::ExternalTexture>()) {
    TINT_ICE(Writer, builder_.Diagnostics())
        << "Multiplanar external texture transform was not run.";
    return false;
  }

  uint32_t array_literal = 0u;
  const auto dim = texture->dim();
  if (dim == ast::TextureDimension::k2dArray ||
      dim == ast::TextureDimension::kCubeArray) {
    array_literal = 1u;
  }

  uint32_t dim_literal = SpvDim2D;
  if (dim == ast::TextureDimension::k1d) {
    dim_literal = SpvDim1D;
    if (texture->Is<sem::SampledTexture>()) {
      push_capability(SpvCapabilitySampled1D);
    } else if (texture->Is<sem::StorageTexture>()) {
      push_capability(SpvCapabilityImage1D);
    }
  }
  if (dim == ast::TextureDimension::k3d) {
    dim_literal = SpvDim3D;
  }
  if (dim == ast::TextureDimension::kCube ||
      dim == ast::TextureDimension::kCubeArray) {
    dim_literal = SpvDimCube;
  }

  uint32_t ms_literal = 0u;
  if (texture->IsAnyOf<sem::MultisampledTexture,
                       sem::DepthMultisampledTexture>()) {
    ms_literal = 1u;
  }

  uint32_t depth_literal = 0u;
  // The Vulkan spec says: The "Depth" operand of OpTypeImage is ignored.
  // In SPIRV, 0 means not depth, 1 means depth, and 2 means unknown.
  // Using anything other than 0 is problematic on various Vulkan drivers.

  uint32_t sampled_literal = 2u;
  if (texture->IsAnyOf<sem::MultisampledTexture, sem::SampledTexture,
                       sem::DepthTexture, sem::DepthMultisampledTexture>()) {
    sampled_literal = 1u;
  }

  if (dim == ast::TextureDimension::kCubeArray) {
    if (texture->IsAnyOf<sem::SampledTexture, sem::DepthTexture>()) {
      push_capability(SpvCapabilitySampledCubeArray);
    }
  }

  uint32_t type_id = Switch(
      texture,
      [&](const sem::DepthTexture*) {
        return GenerateTypeIfNeeded(builder_.create<sem::F32>());
      },
      [&](const sem::DepthMultisampledTexture*) {
        return GenerateTypeIfNeeded(builder_.create<sem::F32>());
      },
      [&](const sem::SampledTexture* t) {
        return GenerateTypeIfNeeded(t->type());
      },
      [&](const sem::MultisampledTexture* t) {
        return GenerateTypeIfNeeded(t->type());
      },
      [&](const sem::StorageTexture* t) {
        return GenerateTypeIfNeeded(t->type());
      },
      [&](Default) { return 0u; });
  if (type_id == 0u) {
    return false;
  }

  uint32_t format_literal = SpvImageFormat_::SpvImageFormatUnknown;
  if (auto* t = texture->As<sem::StorageTexture>()) {
    format_literal = convert_texel_format_to_spv(t->texel_format());
  }

  push_type(spv::Op::OpTypeImage,
            {result, Operand::Int(type_id), Operand::Int(dim_literal),
             Operand::Int(depth_literal), Operand::Int(array_literal),
             Operand::Int(ms_literal), Operand::Int(sampled_literal),
             Operand::Int(format_literal)});

  return true;
}

bool Builder::GenerateArrayType(const sem::Array* ary, const Operand& result) {
  auto elem_type = GenerateTypeIfNeeded(ary->ElemType());
  if (elem_type == 0) {
    return false;
  }

  auto result_id = result.to_i();
  if (ary->IsRuntimeSized()) {
    push_type(spv::Op::OpTypeRuntimeArray, {result, Operand::Int(elem_type)});
  } else {
    auto len_id = GenerateConstantIfNeeded(ScalarConstant::U32(ary->Count()));
    if (len_id == 0) {
      return false;
    }

    push_type(spv::Op::OpTypeArray,
              {result, Operand::Int(elem_type), Operand::Int(len_id)});
  }

  push_annot(spv::Op::OpDecorate,
             {Operand::Int(result_id), Operand::Int(SpvDecorationArrayStride),
              Operand::Int(ary->Stride())});
  return true;
}

bool Builder::GenerateMatrixType(const sem::Matrix* mat,
                                 const Operand& result) {
  auto* col_type = builder_.create<sem::Vector>(mat->type(), mat->rows());
  auto col_type_id = GenerateTypeIfNeeded(col_type);
  if (has_error()) {
    return false;
  }

  push_type(spv::Op::OpTypeMatrix,
            {result, Operand::Int(col_type_id), Operand::Int(mat->columns())});
  return true;
}

bool Builder::GeneratePointerType(const sem::Pointer* ptr,
                                  const Operand& result) {
  auto subtype_id = GenerateTypeIfNeeded(ptr->StoreType());
  if (subtype_id == 0) {
    return false;
  }

  auto stg_class = ConvertStorageClass(ptr->StorageClass());
  if (stg_class == SpvStorageClassMax) {
    error_ = "invalid storage class for pointer";
    return false;
  }

  push_type(spv::Op::OpTypePointer,
            {result, Operand::Int(stg_class), Operand::Int(subtype_id)});

  return true;
}

bool Builder::GenerateReferenceType(const sem::Reference* ref,
                                    const Operand& result) {
  auto subtype_id = GenerateTypeIfNeeded(ref->StoreType());
  if (subtype_id == 0) {
    return false;
  }

  auto stg_class = ConvertStorageClass(ref->StorageClass());
  if (stg_class == SpvStorageClassMax) {
    error_ = "invalid storage class for reference";
    return false;
  }

  push_type(spv::Op::OpTypePointer,
            {result, Operand::Int(stg_class), Operand::Int(subtype_id)});

  return true;
}

bool Builder::GenerateStructType(const sem::Struct* struct_type,
                                 const Operand& result) {
  auto struct_id = result.to_i();

  if (struct_type->Name().IsValid()) {
    push_debug(
        spv::Op::OpName,
        {Operand::Int(struct_id),
         Operand::String(builder_.Symbols().NameFor(struct_type->Name()))});
  }

  OperandList ops;
  ops.push_back(result);

  auto* decl = struct_type->Declaration();
  if (decl &&
      ast::HasAttribute<transform::AddSpirvBlockAttribute::SpirvBlockAttribute>(
          decl->attributes)) {
    push_annot(spv::Op::OpDecorate,
               {Operand::Int(struct_id), Operand::Int(SpvDecorationBlock)});
  }

  for (uint32_t i = 0; i < struct_type->Members().size(); ++i) {
    auto mem_id = GenerateStructMember(struct_id, i, struct_type->Members()[i]);
    if (mem_id == 0) {
      return false;
    }

    ops.push_back(Operand::Int(mem_id));
  }

  push_type(spv::Op::OpTypeStruct, std::move(ops));
  return true;
}

uint32_t Builder::GenerateStructMember(uint32_t struct_id,
                                       uint32_t idx,
                                       const sem::StructMember* member) {
  push_debug(spv::Op::OpMemberName,
             {Operand::Int(struct_id), Operand::Int(idx),
              Operand::String(builder_.Symbols().NameFor(member->Name()))});

  // Note: This will generate layout annotations for *all* structs, whether or
  // not they are used in host-shareable variables. This is officially ok in
  // SPIR-V 1.0 through 1.3. If / when we migrate to using SPIR-V 1.4 we'll have
  // to only generate the layout info for structs used for certain storage
  // classes.

  push_annot(
      spv::Op::OpMemberDecorate,
      {Operand::Int(struct_id), Operand::Int(idx),
       Operand::Int(SpvDecorationOffset), Operand::Int(member->Offset())});

  // Infer and emit matrix layout.
  auto* matrix_type = GetNestedMatrixType(member->Type());
  if (matrix_type) {
    push_annot(spv::Op::OpMemberDecorate,
               {Operand::Int(struct_id), Operand::Int(idx),
                Operand::Int(SpvDecorationColMajor)});
    if (!matrix_type->type()->Is<sem::F32>()) {
      error_ = "matrix scalar element type must be f32";
      return 0;
    }
    const auto scalar_elem_size = 4;
    const auto effective_row_count = (matrix_type->rows() == 2) ? 2 : 4;
    push_annot(spv::Op::OpMemberDecorate,
               {Operand::Int(struct_id), Operand::Int(idx),
                Operand::Int(SpvDecorationMatrixStride),
                Operand::Int(effective_row_count * scalar_elem_size)});
  }

  return GenerateTypeIfNeeded(member->Type());
}

bool Builder::GenerateVectorType(const sem::Vector* vec,
                                 const Operand& result) {
  auto type_id = GenerateTypeIfNeeded(vec->type());
  if (has_error()) {
    return false;
  }

  push_type(spv::Op::OpTypeVector,
            {result, Operand::Int(type_id), Operand::Int(vec->Width())});
  return true;
}

SpvStorageClass Builder::ConvertStorageClass(ast::StorageClass klass) const {
  switch (klass) {
    case ast::StorageClass::kInvalid:
      return SpvStorageClassMax;
    case ast::StorageClass::kInput:
      return SpvStorageClassInput;
    case ast::StorageClass::kOutput:
      return SpvStorageClassOutput;
    case ast::StorageClass::kUniform:
      return SpvStorageClassUniform;
    case ast::StorageClass::kWorkgroup:
      return SpvStorageClassWorkgroup;
    case ast::StorageClass::kUniformConstant:
      return SpvStorageClassUniformConstant;
    case ast::StorageClass::kStorage:
      return SpvStorageClassStorageBuffer;
    case ast::StorageClass::kPrivate:
      return SpvStorageClassPrivate;
    case ast::StorageClass::kFunction:
      return SpvStorageClassFunction;
    case ast::StorageClass::kNone:
      break;
  }
  return SpvStorageClassMax;
}

SpvBuiltIn Builder::ConvertBuiltin(ast::Builtin builtin,
                                   ast::StorageClass storage) {
  switch (builtin) {
    case ast::Builtin::kPosition:
      if (storage == ast::StorageClass::kInput) {
        return SpvBuiltInFragCoord;
      } else if (storage == ast::StorageClass::kOutput) {
        return SpvBuiltInPosition;
      } else {
        TINT_ICE(Writer, builder_.Diagnostics())
            << "invalid storage class for builtin";
        break;
      }
    case ast::Builtin::kVertexIndex:
      return SpvBuiltInVertexIndex;
    case ast::Builtin::kInstanceIndex:
      return SpvBuiltInInstanceIndex;
    case ast::Builtin::kFrontFacing:
      return SpvBuiltInFrontFacing;
    case ast::Builtin::kFragDepth:
      return SpvBuiltInFragDepth;
    case ast::Builtin::kLocalInvocationId:
      return SpvBuiltInLocalInvocationId;
    case ast::Builtin::kLocalInvocationIndex:
      return SpvBuiltInLocalInvocationIndex;
    case ast::Builtin::kGlobalInvocationId:
      return SpvBuiltInGlobalInvocationId;
    case ast::Builtin::kPointSize:
      return SpvBuiltInPointSize;
    case ast::Builtin::kWorkgroupId:
      return SpvBuiltInWorkgroupId;
    case ast::Builtin::kNumWorkgroups:
      return SpvBuiltInNumWorkgroups;
    case ast::Builtin::kSampleIndex:
      push_capability(SpvCapabilitySampleRateShading);
      return SpvBuiltInSampleId;
    case ast::Builtin::kSampleMask:
      return SpvBuiltInSampleMask;
    case ast::Builtin::kNone:
      break;
  }
  return SpvBuiltInMax;
}

void Builder::AddInterpolationDecorations(uint32_t id,
                                          ast::InterpolationType type,
                                          ast::InterpolationSampling sampling) {
  switch (type) {
    case ast::InterpolationType::kLinear:
      push_annot(spv::Op::OpDecorate,
                 {Operand::Int(id), Operand::Int(SpvDecorationNoPerspective)});
      break;
    case ast::InterpolationType::kFlat:
      push_annot(spv::Op::OpDecorate,
                 {Operand::Int(id), Operand::Int(SpvDecorationFlat)});
      break;
    case ast::InterpolationType::kPerspective:
      break;
  }
  switch (sampling) {
    case ast::InterpolationSampling::kCentroid:
      push_annot(spv::Op::OpDecorate,
                 {Operand::Int(id), Operand::Int(SpvDecorationCentroid)});
      break;
    case ast::InterpolationSampling::kSample:
      push_capability(SpvCapabilitySampleRateShading);
      push_annot(spv::Op::OpDecorate,
                 {Operand::Int(id), Operand::Int(SpvDecorationSample)});
      break;
    case ast::InterpolationSampling::kCenter:
    case ast::InterpolationSampling::kNone:
      break;
  }
}

SpvImageFormat Builder::convert_texel_format_to_spv(
    const ast::TexelFormat format) {
  switch (format) {
    case ast::TexelFormat::kR32Uint:
      return SpvImageFormatR32ui;
    case ast::TexelFormat::kR32Sint:
      return SpvImageFormatR32i;
    case ast::TexelFormat::kR32Float:
      return SpvImageFormatR32f;
    case ast::TexelFormat::kRgba8Unorm:
      return SpvImageFormatRgba8;
    case ast::TexelFormat::kRgba8Snorm:
      return SpvImageFormatRgba8Snorm;
    case ast::TexelFormat::kRgba8Uint:
      return SpvImageFormatRgba8ui;
    case ast::TexelFormat::kRgba8Sint:
      return SpvImageFormatRgba8i;
    case ast::TexelFormat::kRg32Uint:
      push_capability(SpvCapabilityStorageImageExtendedFormats);
      return SpvImageFormatRg32ui;
    case ast::TexelFormat::kRg32Sint:
      push_capability(SpvCapabilityStorageImageExtendedFormats);
      return SpvImageFormatRg32i;
    case ast::TexelFormat::kRg32Float:
      push_capability(SpvCapabilityStorageImageExtendedFormats);
      return SpvImageFormatRg32f;
    case ast::TexelFormat::kRgba16Uint:
      return SpvImageFormatRgba16ui;
    case ast::TexelFormat::kRgba16Sint:
      return SpvImageFormatRgba16i;
    case ast::TexelFormat::kRgba16Float:
      return SpvImageFormatRgba16f;
    case ast::TexelFormat::kRgba32Uint:
      return SpvImageFormatRgba32ui;
    case ast::TexelFormat::kRgba32Sint:
      return SpvImageFormatRgba32i;
    case ast::TexelFormat::kRgba32Float:
      return SpvImageFormatRgba32f;
    case ast::TexelFormat::kNone:
      return SpvImageFormatUnknown;
  }
  return SpvImageFormatUnknown;
}

bool Builder::push_function_inst(spv::Op op, const OperandList& operands) {
  if (functions_.empty()) {
    std::ostringstream ss;
    ss << "Internal error: trying to add SPIR-V instruction " << int(op)
       << " outside a function";
    error_ = ss.str();
    return false;
  }
  functions_.back().push_inst(op, operands);
  return true;
}

bool Builder::InsideBasicBlock() const {
  if (functions_.empty()) {
    return false;
  }
  const auto& instructions = functions_.back().instructions();
  if (instructions.empty()) {
    // The Function object does not explicitly represent its entry block
    // label.  So return *true* because an empty list means the only
    // thing in the function is that entry block label.
    return true;
  }
  const auto& inst = instructions.back();
  switch (inst.opcode()) {
    case spv::Op::OpBranch:
    case spv::Op::OpBranchConditional:
    case spv::Op::OpSwitch:
    case spv::Op::OpReturn:
    case spv::Op::OpReturnValue:
    case spv::Op::OpUnreachable:
    case spv::Op::OpKill:
    case spv::Op::OpTerminateInvocation:
      return false;
    default:
      break;
  }
  return true;
}

Builder::ContinuingInfo::ContinuingInfo(
    const ast::Statement* the_last_statement,
    uint32_t loop_id,
    uint32_t break_id)
    : last_statement(the_last_statement),
      loop_header_id(loop_id),
      break_target_id(break_id) {
  TINT_ASSERT(Writer, last_statement != nullptr);
  TINT_ASSERT(Writer, loop_header_id != 0u);
  TINT_ASSERT(Writer, break_target_id != 0u);
}

Builder::Backedge::Backedge(spv::Op the_opcode, OperandList the_operands)
    : opcode(the_opcode), operands(the_operands) {}

Builder::Backedge::Backedge(const Builder::Backedge& other) = default;
Builder::Backedge& Builder::Backedge::operator=(
    const Builder::Backedge& other) = default;
Builder::Backedge::~Backedge() = default;

}  // namespace tint::writer::spirv