// 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/inspector/inspector.h"

#include <limits>
#include <utility>

#include "src/ast/bool_literal.h"
#include "src/ast/call_expression.h"
#include "src/ast/float_literal.h"
#include "src/ast/interpolate_decoration.h"
#include "src/ast/location_decoration.h"
#include "src/ast/module.h"
#include "src/ast/override_decoration.h"
#include "src/ast/sint_literal.h"
#include "src/ast/uint_literal.h"
#include "src/sem/array.h"
#include "src/sem/call.h"
#include "src/sem/depth_multisampled_texture_type.h"
#include "src/sem/f32_type.h"
#include "src/sem/function.h"
#include "src/sem/i32_type.h"
#include "src/sem/matrix_type.h"
#include "src/sem/multisampled_texture_type.h"
#include "src/sem/sampled_texture_type.h"
#include "src/sem/statement.h"
#include "src/sem/storage_texture_type.h"
#include "src/sem/struct.h"
#include "src/sem/u32_type.h"
#include "src/sem/variable.h"
#include "src/sem/vector_type.h"
#include "src/sem/void_type.h"
#include "src/utils/math.h"
#include "src/utils/unique_vector.h"

namespace tint {
namespace inspector {

namespace {

void AppendResourceBindings(std::vector<ResourceBinding>* dest,
                            const std::vector<ResourceBinding>& orig) {
  TINT_ASSERT(Inspector, dest);
  if (!dest) {
    return;
  }

  dest->reserve(dest->size() + orig.size());
  dest->insert(dest->end(), orig.begin(), orig.end());
}

std::tuple<ComponentType, CompositionType> CalculateComponentAndComposition(
    const sem::Type* type) {
  if (type->is_float_scalar()) {
    return {ComponentType::kFloat, CompositionType::kScalar};
  } else if (type->is_float_vector()) {
    auto* vec = type->As<sem::Vector>();
    if (vec->Width() == 2) {
      return {ComponentType::kFloat, CompositionType::kVec2};
    } else if (vec->Width() == 3) {
      return {ComponentType::kFloat, CompositionType::kVec3};
    } else if (vec->Width() == 4) {
      return {ComponentType::kFloat, CompositionType::kVec4};
    }
  } else if (type->is_unsigned_integer_scalar()) {
    return {ComponentType::kUInt, CompositionType::kScalar};
  } else if (type->is_unsigned_integer_vector()) {
    auto* vec = type->As<sem::Vector>();
    if (vec->Width() == 2) {
      return {ComponentType::kUInt, CompositionType::kVec2};
    } else if (vec->Width() == 3) {
      return {ComponentType::kUInt, CompositionType::kVec3};
    } else if (vec->Width() == 4) {
      return {ComponentType::kUInt, CompositionType::kVec4};
    }
  } else if (type->is_signed_integer_scalar()) {
    return {ComponentType::kSInt, CompositionType::kScalar};
  } else if (type->is_signed_integer_vector()) {
    auto* vec = type->As<sem::Vector>();
    if (vec->Width() == 2) {
      return {ComponentType::kSInt, CompositionType::kVec2};
    } else if (vec->Width() == 3) {
      return {ComponentType::kSInt, CompositionType::kVec3};
    } else if (vec->Width() == 4) {
      return {ComponentType::kSInt, CompositionType::kVec4};
    }
  }
  return {ComponentType::kUnknown, CompositionType::kUnknown};
}

std::tuple<InterpolationType, InterpolationSampling> CalculateInterpolationData(
    const sem::Type* type,
    const ast::DecorationList& decorations) {
  auto* interpolation_decoration =
      ast::GetDecoration<ast::InterpolateDecoration>(decorations);
  if (type->is_integer_scalar_or_vector()) {
    return {InterpolationType::kFlat, InterpolationSampling::kNone};
  }

  if (!interpolation_decoration) {
    return {InterpolationType::kPerspective, InterpolationSampling::kCenter};
  }

  auto interpolation_type = interpolation_decoration->type;
  auto sampling = interpolation_decoration->sampling;
  if (interpolation_type != ast::InterpolationType::kFlat &&
      sampling == ast::InterpolationSampling::kNone) {
    sampling = ast::InterpolationSampling::kCenter;
  }
  return {ASTToInspectorInterpolationType(interpolation_type),
          ASTToInspectorInterpolationSampling(sampling)};
}

}  // namespace

Inspector::Inspector(const Program* program) : program_(program) {}

Inspector::~Inspector() = default;

std::vector<EntryPoint> Inspector::GetEntryPoints() {
  std::vector<EntryPoint> result;

  for (auto* func : program_->AST().Functions()) {
    if (!func->IsEntryPoint()) {
      continue;
    }

    auto* sem = program_->Sem().Get(func);

    EntryPoint entry_point;
    entry_point.name = program_->Symbols().NameFor(func->symbol);
    entry_point.remapped_name = program_->Symbols().NameFor(func->symbol);
    entry_point.stage = func->PipelineStage();

    auto wgsize = sem->WorkgroupSize();
    entry_point.workgroup_size_x = wgsize[0].value;
    entry_point.workgroup_size_y = wgsize[1].value;
    entry_point.workgroup_size_z = wgsize[2].value;
    if (wgsize[0].overridable_const || wgsize[1].overridable_const ||
        wgsize[2].overridable_const) {
      // TODO(crbug.com/tint/713): Handle overridable constants.
      TINT_ASSERT(Inspector, false);
    }

    for (auto* param : sem->Parameters()) {
      AddEntryPointInOutVariables(
          program_->Symbols().NameFor(param->Declaration()->symbol),
          param->Type(), param->Declaration()->decorations,
          entry_point.input_variables);

      entry_point.input_position_used |=
          ContainsBuiltin(ast::Builtin::kPosition, param->Type(),
                          param->Declaration()->decorations);
      entry_point.front_facing_used |=
          ContainsBuiltin(ast::Builtin::kFrontFacing, param->Type(),
                          param->Declaration()->decorations);
      entry_point.sample_index_used |=
          ContainsBuiltin(ast::Builtin::kSampleIndex, param->Type(),
                          param->Declaration()->decorations);
      entry_point.input_sample_mask_used |=
          ContainsBuiltin(ast::Builtin::kSampleMask, param->Type(),
                          param->Declaration()->decorations);
      entry_point.num_workgroups_used |=
          ContainsBuiltin(ast::Builtin::kNumWorkgroups, param->Type(),
                          param->Declaration()->decorations);
    }

    if (!sem->ReturnType()->Is<sem::Void>()) {
      AddEntryPointInOutVariables("<retval>", sem->ReturnType(),
                                  func->return_type_decorations,
                                  entry_point.output_variables);

      entry_point.output_sample_mask_used =
          ContainsBuiltin(ast::Builtin::kSampleMask, sem->ReturnType(),
                          func->return_type_decorations);
    }

    for (auto* var : sem->TransitivelyReferencedGlobals()) {
      auto* decl = var->Declaration();

      auto name = program_->Symbols().NameFor(decl->symbol);

      auto* global = var->As<sem::GlobalVariable>();
      if (global && global->IsPipelineConstant()) {
        OverridableConstant overridable_constant;
        overridable_constant.name = name;
        overridable_constant.numeric_id = global->ConstantId();
        auto* type = var->Type();
        TINT_ASSERT(Inspector, type->is_scalar());
        if (type->is_bool_scalar_or_vector()) {
          overridable_constant.type = OverridableConstant::Type::kBool;
        } else if (type->is_float_scalar()) {
          overridable_constant.type = OverridableConstant::Type::kFloat32;
        } else if (type->is_signed_integer_scalar()) {
          overridable_constant.type = OverridableConstant::Type::kInt32;
        } else if (type->is_unsigned_integer_scalar()) {
          overridable_constant.type = OverridableConstant::Type::kUint32;
        } else {
          TINT_UNREACHABLE(Inspector, diagnostics_);
        }

        overridable_constant.is_initialized =
            global->Declaration()->constructor;
        auto* override_deco = ast::GetDecoration<ast::OverrideDecoration>(
            global->Declaration()->decorations);
        overridable_constant.is_numeric_id_specified =
            override_deco ? override_deco->has_value : false;

        entry_point.overridable_constants.push_back(overridable_constant);
      }
    }

    result.push_back(std::move(entry_point));
  }

  return result;
}

std::string Inspector::GetRemappedNameForEntryPoint(
    const std::string& entry_point) {
  // TODO(rharrison): Reenable once all of the backends are using the renamed
  //                  entry points.

  //  auto* func = FindEntryPointByName(entry_point);
  //  if (!func) {
  //    return {};
  //  }
  //  return func->name();
  return entry_point;
}

std::map<uint32_t, Scalar> Inspector::GetConstantIDs() {
  std::map<uint32_t, Scalar> result;
  for (auto* var : program_->AST().GlobalVariables()) {
    auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
    if (!global || !global->IsPipelineConstant()) {
      continue;
    }

    // If there are conflicting defintions for a constant id, that is invalid
    // WGSL, so the resolver should catch it. Thus here the inspector just
    // assumes all definitions of the constant id are the same, so only needs
    // to find the first reference to constant id.
    uint32_t constant_id = global->ConstantId();
    if (result.find(constant_id) != result.end()) {
      continue;
    }

    if (!var->constructor) {
      result[constant_id] = Scalar();
      continue;
    }

    auto* literal = var->constructor->As<ast::Literal>();
    if (!literal) {
      // This is invalid WGSL, but handling gracefully.
      result[constant_id] = Scalar();
      continue;
    }

    if (auto* l = literal->As<ast::BoolLiteral>()) {
      result[constant_id] = Scalar(l->value);
      continue;
    }

    if (auto* l = literal->As<ast::UintLiteral>()) {
      result[constant_id] = Scalar(l->value);
      continue;
    }

    if (auto* l = literal->As<ast::SintLiteral>()) {
      result[constant_id] = Scalar(l->value);
      continue;
    }

    if (auto* l = literal->As<ast::FloatLiteral>()) {
      result[constant_id] = Scalar(l->value);
      continue;
    }

    result[constant_id] = Scalar();
  }

  return result;
}

std::map<std::string, uint32_t> Inspector::GetConstantNameToIdMap() {
  std::map<std::string, uint32_t> result;
  for (auto* var : program_->AST().GlobalVariables()) {
    auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
    if (global && global->IsPipelineConstant()) {
      auto name = program_->Symbols().NameFor(var->symbol);
      result[name] = global->ConstantId();
    }
  }
  return result;
}

uint32_t Inspector::GetStorageSize(const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return 0;
  }

  size_t size = 0;
  auto* func_sem = program_->Sem().Get(func);
  for (auto& ruv : func_sem->TransitivelyReferencedUniformVariables()) {
    const sem::Struct* s = ruv.first->Type()->UnwrapRef()->As<sem::Struct>();
    if (s && s->IsBlockDecorated()) {
      size += s->Size();
    }
  }
  for (auto& rsv : func_sem->TransitivelyReferencedStorageBufferVariables()) {
    const sem::Struct* s = rsv.first->Type()->UnwrapRef()->As<sem::Struct>();
    if (s) {
      size += s->Size();
    }
  }

  if (static_cast<uint64_t>(size) >
      static_cast<uint64_t>(std::numeric_limits<uint32_t>::max())) {
    return std::numeric_limits<uint32_t>::max();
  }
  return static_cast<uint32_t>(size);
}

std::vector<ResourceBinding> Inspector::GetResourceBindings(
    const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  std::vector<ResourceBinding> result;
  for (auto fn : {
           &Inspector::GetUniformBufferResourceBindings,
           &Inspector::GetStorageBufferResourceBindings,
           &Inspector::GetReadOnlyStorageBufferResourceBindings,
           &Inspector::GetSamplerResourceBindings,
           &Inspector::GetComparisonSamplerResourceBindings,
           &Inspector::GetSampledTextureResourceBindings,
           &Inspector::GetMultisampledTextureResourceBindings,
           &Inspector::GetWriteOnlyStorageTextureResourceBindings,
           &Inspector::GetDepthTextureResourceBindings,
           &Inspector::GetDepthMultisampledTextureResourceBindings,
           &Inspector::GetExternalTextureResourceBindings,
       }) {
    AppendResourceBindings(&result, (this->*fn)(entry_point));
  }
  return result;
}

std::vector<ResourceBinding> Inspector::GetUniformBufferResourceBindings(
    const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  std::vector<ResourceBinding> result;

  auto* func_sem = program_->Sem().Get(func);
  for (auto& ruv : func_sem->TransitivelyReferencedUniformVariables()) {
    auto* var = ruv.first;
    auto binding_info = ruv.second;

    auto* unwrapped_type = var->Type()->UnwrapRef();
    auto* str = unwrapped_type->As<sem::Struct>();
    if (str == nullptr) {
      continue;
    }

    if (!str->IsBlockDecorated()) {
      continue;
    }

    ResourceBinding entry;
    entry.resource_type = ResourceBinding::ResourceType::kUniformBuffer;
    entry.bind_group = binding_info.group->value;
    entry.binding = binding_info.binding->value;
    entry.size = str->Size();
    entry.size_no_padding = str->SizeNoPadding();

    result.push_back(entry);
  }

  return result;
}

std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindings(
    const std::string& entry_point) {
  return GetStorageBufferResourceBindingsImpl(entry_point, false);
}

std::vector<ResourceBinding>
Inspector::GetReadOnlyStorageBufferResourceBindings(
    const std::string& entry_point) {
  return GetStorageBufferResourceBindingsImpl(entry_point, true);
}

std::vector<ResourceBinding> Inspector::GetSamplerResourceBindings(
    const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  std::vector<ResourceBinding> result;

  auto* func_sem = program_->Sem().Get(func);
  for (auto& rs : func_sem->TransitivelyReferencedSamplerVariables()) {
    auto binding_info = rs.second;

    ResourceBinding entry;
    entry.resource_type = ResourceBinding::ResourceType::kSampler;
    entry.bind_group = binding_info.group->value;
    entry.binding = binding_info.binding->value;

    result.push_back(entry);
  }

  return result;
}

std::vector<ResourceBinding> Inspector::GetComparisonSamplerResourceBindings(
    const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  std::vector<ResourceBinding> result;

  auto* func_sem = program_->Sem().Get(func);
  for (auto& rcs :
       func_sem->TransitivelyReferencedComparisonSamplerVariables()) {
    auto binding_info = rcs.second;

    ResourceBinding entry;
    entry.resource_type = ResourceBinding::ResourceType::kComparisonSampler;
    entry.bind_group = binding_info.group->value;
    entry.binding = binding_info.binding->value;

    result.push_back(entry);
  }

  return result;
}

std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindings(
    const std::string& entry_point) {
  return GetSampledTextureResourceBindingsImpl(entry_point, false);
}

std::vector<ResourceBinding> Inspector::GetMultisampledTextureResourceBindings(
    const std::string& entry_point) {
  return GetSampledTextureResourceBindingsImpl(entry_point, true);
}

std::vector<ResourceBinding>
Inspector::GetWriteOnlyStorageTextureResourceBindings(
    const std::string& entry_point) {
  return GetStorageTextureResourceBindingsImpl(entry_point);
}

std::vector<ResourceBinding> Inspector::GetTextureResourceBindings(
    const std::string& entry_point,
    const tint::TypeInfo& texture_type,
    ResourceBinding::ResourceType resource_type) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  std::vector<ResourceBinding> result;
  auto* func_sem = program_->Sem().Get(func);
  for (auto& ref :
       func_sem->TransitivelyReferencedVariablesOfType(texture_type)) {
    auto* var = ref.first;
    auto binding_info = ref.second;

    ResourceBinding entry;
    entry.resource_type = resource_type;
    entry.bind_group = binding_info.group->value;
    entry.binding = binding_info.binding->value;

    auto* tex = var->Type()->UnwrapRef()->As<sem::Texture>();
    entry.dim =
        TypeTextureDimensionToResourceBindingTextureDimension(tex->dim());

    result.push_back(entry);
  }

  return result;
}

std::vector<ResourceBinding> Inspector::GetDepthTextureResourceBindings(
    const std::string& entry_point) {
  return GetTextureResourceBindings(
      entry_point, TypeInfo::Of<sem::DepthTexture>(),
      ResourceBinding::ResourceType::kDepthTexture);
}

std::vector<ResourceBinding>
Inspector::GetDepthMultisampledTextureResourceBindings(
    const std::string& entry_point) {
  return GetTextureResourceBindings(
      entry_point, TypeInfo::Of<sem::DepthMultisampledTexture>(),
      ResourceBinding::ResourceType::kDepthMultisampledTexture);
}

std::vector<ResourceBinding> Inspector::GetExternalTextureResourceBindings(
    const std::string& entry_point) {
  return GetTextureResourceBindings(
      entry_point, TypeInfo::Of<sem::ExternalTexture>(),
      ResourceBinding::ResourceType::kExternalTexture);
}

std::vector<SamplerTexturePair> Inspector::GetSamplerTextureUses(
    const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  GenerateSamplerTargets();

  auto it = sampler_targets_->find(entry_point);
  if (it == sampler_targets_->end()) {
    return {};
  }
  return it->second;
}

uint32_t Inspector::GetWorkgroupStorageSize(const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return 0;
  }

  uint32_t total_size = 0;
  auto* func_sem = program_->Sem().Get(func);
  for (const sem::Variable* var : func_sem->TransitivelyReferencedGlobals()) {
    if (var->StorageClass() == ast::StorageClass::kWorkgroup) {
      auto* ty = var->Type()->UnwrapRef();
      uint32_t align = ty->Align();
      uint32_t size = ty->Size();

      // This essentially matches std430 layout rules from GLSL, which are in
      // turn specified as an upper bound for Vulkan layout sizing. Since D3D
      // and Metal are even less specific, we assume Vulkan behavior as a
      // good-enough approximation everywhere.
      total_size += utils::RoundUp(align, size);
    }
  }

  return total_size;
}

const ast::Function* Inspector::FindEntryPointByName(const std::string& name) {
  auto* func = program_->AST().Functions().Find(program_->Symbols().Get(name));
  if (!func) {
    diagnostics_.add_error(diag::System::Inspector, name + " was not found!");
    return nullptr;
  }

  if (!func->IsEntryPoint()) {
    diagnostics_.add_error(diag::System::Inspector,
                           name + " is not an entry point!");
    return nullptr;
  }

  return func;
}

void Inspector::AddEntryPointInOutVariables(
    std::string name,
    const sem::Type* type,
    const ast::DecorationList& decorations,
    std::vector<StageVariable>& variables) const {
  // Skip builtins.
  if (ast::HasDecoration<ast::BuiltinDecoration>(decorations)) {
    return;
  }

  auto* unwrapped_type = type->UnwrapRef();

  if (auto* struct_ty = unwrapped_type->As<sem::Struct>()) {
    // Recurse into members.
    for (auto* member : struct_ty->Members()) {
      AddEntryPointInOutVariables(
          name + "." +
              program_->Symbols().NameFor(member->Declaration()->symbol),
          member->Type(), member->Declaration()->decorations, variables);
    }
    return;
  }

  // Base case: add the variable.

  StageVariable stage_variable;
  stage_variable.name = name;
  std::tie(stage_variable.component_type, stage_variable.composition_type) =
      CalculateComponentAndComposition(type);

  auto* location = ast::GetDecoration<ast::LocationDecoration>(decorations);
  TINT_ASSERT(Inspector, location != nullptr);
  stage_variable.has_location_decoration = true;
  stage_variable.location_decoration = location->value;

  std::tie(stage_variable.interpolation_type,
           stage_variable.interpolation_sampling) =
      CalculateInterpolationData(type, decorations);

  variables.push_back(stage_variable);
}

bool Inspector::ContainsBuiltin(ast::Builtin builtin,
                                const sem::Type* type,
                                const ast::DecorationList& decorations) const {
  auto* unwrapped_type = type->UnwrapRef();

  if (auto* struct_ty = unwrapped_type->As<sem::Struct>()) {
    // Recurse into members.
    for (auto* member : struct_ty->Members()) {
      if (ContainsBuiltin(builtin, member->Type(),
                          member->Declaration()->decorations)) {
        return true;
      }
    }
    return false;
  }

  // Base case: check for builtin
  auto* builtin_declaration =
      ast::GetDecoration<ast::BuiltinDecoration>(decorations);
  if (!builtin_declaration || builtin_declaration->builtin != builtin) {
    return false;
  }

  return true;
}

std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
    const std::string& entry_point,
    bool read_only) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  auto* func_sem = program_->Sem().Get(func);
  std::vector<ResourceBinding> result;
  for (auto& rsv : func_sem->TransitivelyReferencedStorageBufferVariables()) {
    auto* var = rsv.first;
    auto binding_info = rsv.second;

    if (read_only != (var->Access() == ast::Access::kRead)) {
      continue;
    }

    auto* str = var->Type()->UnwrapRef()->As<sem::Struct>();
    if (!str) {
      continue;
    }

    ResourceBinding entry;
    entry.resource_type =
        read_only ? ResourceBinding::ResourceType::kReadOnlyStorageBuffer
                  : ResourceBinding::ResourceType::kStorageBuffer;
    entry.bind_group = binding_info.group->value;
    entry.binding = binding_info.binding->value;
    entry.size = str->Size();
    entry.size_no_padding = str->SizeNoPadding();

    result.push_back(entry);
  }

  return result;
}

std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindingsImpl(
    const std::string& entry_point,
    bool multisampled_only) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  std::vector<ResourceBinding> result;
  auto* func_sem = program_->Sem().Get(func);
  auto referenced_variables =
      multisampled_only
          ? func_sem->TransitivelyReferencedMultisampledTextureVariables()
          : func_sem->TransitivelyReferencedSampledTextureVariables();
  for (auto& ref : referenced_variables) {
    auto* var = ref.first;
    auto binding_info = ref.second;

    ResourceBinding entry;
    entry.resource_type =
        multisampled_only ? ResourceBinding::ResourceType::kMultisampledTexture
                          : ResourceBinding::ResourceType::kSampledTexture;
    entry.bind_group = binding_info.group->value;
    entry.binding = binding_info.binding->value;

    auto* texture_type = var->Type()->UnwrapRef()->As<sem::Texture>();
    entry.dim = TypeTextureDimensionToResourceBindingTextureDimension(
        texture_type->dim());

    const sem::Type* base_type = nullptr;
    if (multisampled_only) {
      base_type = texture_type->As<sem::MultisampledTexture>()->type();
    } else {
      base_type = texture_type->As<sem::SampledTexture>()->type();
    }
    entry.sampled_kind = BaseTypeToSampledKind(base_type);

    result.push_back(entry);
  }

  return result;
}

std::vector<ResourceBinding> Inspector::GetStorageTextureResourceBindingsImpl(
    const std::string& entry_point) {
  auto* func = FindEntryPointByName(entry_point);
  if (!func) {
    return {};
  }

  auto* func_sem = program_->Sem().Get(func);
  std::vector<ResourceBinding> result;
  for (auto& ref :
       func_sem->TransitivelyReferencedVariablesOfType<sem::StorageTexture>()) {
    auto* var = ref.first;
    auto binding_info = ref.second;

    auto* texture_type = var->Type()->UnwrapRef()->As<sem::StorageTexture>();

    ResourceBinding entry;
    entry.resource_type =
        ResourceBinding::ResourceType::kWriteOnlyStorageTexture;
    entry.bind_group = binding_info.group->value;
    entry.binding = binding_info.binding->value;

    entry.dim = TypeTextureDimensionToResourceBindingTextureDimension(
        texture_type->dim());

    auto* base_type = texture_type->type();
    entry.sampled_kind = BaseTypeToSampledKind(base_type);
    entry.image_format = TypeImageFormatToResourceBindingImageFormat(
        texture_type->image_format());

    result.push_back(entry);
  }

  return result;
}

void Inspector::GenerateSamplerTargets() {
  // Do not re-generate, since |program_| should not change during the lifetime
  // of the inspector.
  if (sampler_targets_ != nullptr) {
    return;
  }

  sampler_targets_ = std::make_unique<std::unordered_map<
      std::string, utils::UniqueVector<SamplerTexturePair>>>();

  auto& sem = program_->Sem();

  for (auto* node : program_->ASTNodes().Objects()) {
    auto* c = node->As<ast::CallExpression>();
    if (!c) {
      continue;
    }

    auto* call = sem.Get(c);
    if (!call) {
      continue;
    }

    auto* i = call->Target()->As<sem::Intrinsic>();
    if (!i) {
      continue;
    }

    const auto& signature = i->Signature();
    int sampler_index = signature.IndexOf(sem::ParameterUsage::kSampler);
    if (sampler_index == -1) {
      continue;
    }

    int texture_index = signature.IndexOf(sem::ParameterUsage::kTexture);
    if (texture_index == -1) {
      continue;
    }

    auto* call_func = call->Stmt()->Function();
    std::vector<const sem::Function*> entry_points;
    if (call_func->Declaration()->IsEntryPoint()) {
      entry_points = {call_func};
    } else {
      entry_points = call_func->AncestorEntryPoints();
    }

    if (entry_points.empty()) {
      continue;
    }

    auto* t = c->args[texture_index];
    auto* s = c->args[sampler_index];

    GetOriginatingResources(
        std::array<const ast::Expression*, 2>{t, s},
        [&](std::array<const sem::GlobalVariable*, 2> globals) {
          auto* texture = globals[0];
          sem::BindingPoint texture_binding_point = {
              texture->Declaration()->BindingPoint().group->value,
              texture->Declaration()->BindingPoint().binding->value};

          auto* sampler = globals[1];
          sem::BindingPoint sampler_binding_point = {
              sampler->Declaration()->BindingPoint().group->value,
              sampler->Declaration()->BindingPoint().binding->value};

          for (auto* entry_point : entry_points) {
            const auto& ep_name =
                program_->Symbols().NameFor(entry_point->Declaration()->symbol);
            (*sampler_targets_)[ep_name].add(
                {sampler_binding_point, texture_binding_point});
          }
        });
  }
}

template <size_t N, typename F>
void Inspector::GetOriginatingResources(
    std::array<const ast::Expression*, N> exprs,
    F&& callback) {
  if (!program_->IsValid()) {
    TINT_ICE(Inspector, diagnostics_)
        << "attempting to get originating resources in invalid program";
    return;
  }

  auto& sem = program_->Sem();

  std::array<const sem::GlobalVariable*, N> globals{};
  std::array<const sem::Parameter*, N> parameters{};
  utils::UniqueVector<const ast::CallExpression*> callsites;

  for (size_t i = 0; i < N; i++) {
    auto*& expr = exprs[i];
    // Resolve each of the expressions
    while (true) {
      if (auto* user = sem.Get<sem::VariableUser>(expr)) {
        auto* var = user->Variable();

        if (auto* global = tint::As<sem::GlobalVariable>(var)) {
          // Found the global resource declaration.
          globals[i] = global;
          break;  // Done with this expression.
        }

        if (auto* local = tint::As<sem::LocalVariable>(var)) {
          // Chase the variable
          expr = local->Declaration()->constructor;
          if (!expr) {
            TINT_ICE(Inspector, diagnostics_)
                << "resource variable had no initializer";
            return;
          }
          continue;  // Continue chasing the expression in this function
        }

        if (auto* param = tint::As<sem::Parameter>(var)) {
          // Gather each of the callers of this function
          auto* func = tint::As<sem::Function>(param->Owner());
          if (func->CallSites().empty()) {
            // One or more of the expressions is a parameter, but this function
            // is not called. Ignore.
            return;
          }
          for (auto* call : func->CallSites()) {
            callsites.add(call->Declaration());
          }
          // Need to evaluate each function call with the group of
          // expressions, so move on to the next expression.
          parameters[i] = param;
          break;
        }

        TINT_ICE(Inspector, diagnostics_)
            << "unexpected variable type " << var->TypeInfo().name;
      }

      if (auto* unary = tint::As<ast::UnaryOpExpression>(expr)) {
        switch (unary->op) {
          case ast::UnaryOp::kAddressOf:
          case ast::UnaryOp::kIndirection:
            // `*` and `&` are the only valid unary ops for a resource type,
            // and must be balanced in order for the program to have passed
            // validation. Just skip past these.
            expr = unary->expr;
            continue;
          default: {
            TINT_ICE(Inspector, diagnostics_)
                << "unexpected unary op on resource: " << unary->op;
            return;
          }
        }
      }

      TINT_ICE(Inspector, diagnostics_)
          << "cannot resolve originating resource with expression type "
          << expr->TypeInfo().name;
      return;
    }
  }

  if (callsites.size()) {
    for (auto* call_expr : callsites) {
      // Make a copy of the expressions for this callsite
      std::array<const ast::Expression*, N> call_exprs = exprs;
      // Patch all the parameter expressions with their argument
      for (size_t i = 0; i < N; i++) {
        if (auto* param = parameters[i]) {
          call_exprs[i] = call_expr->args[param->Index()];
        }
      }
      // Now call GetOriginatingResources() with from the callsite
      GetOriginatingResources(call_exprs, callback);
    }
  } else {
    // All the expressions resolved to globals
    callback(globals);
  }
}

}  // namespace inspector
}  // namespace tint