src/tint/lang/msl/intrinsic/data.cc

// Copyright 2023 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
//    list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
//    this list of conditions and the following disclaimer in the documentation
//    and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
//    contributors may be used to endorse or promote products derived from
//    this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

////////////////////////////////////////////////////////////////////////////////
// File generated by 'tools/src/cmd/gen' using the template:
//   src/tint/lang/msl/intrinsic/data.cc.tmpl
//
// To regenerate run: './tools/run gen'
//
//                       Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////

#include <limits>
#include <string>

#include "src/tint/lang/core/intrinsic/type_matchers.h"
#include "src/tint/lang/msl/intrinsic/dialect.h"
#include "src/tint/utils/text/string_stream.h"

namespace tint::msl::intrinsic {

using namespace tint::core::intrinsic;  // NOLINT(build/namespaces)

namespace {

using ConstEvalFunctionIndex = tint::core::intrinsic::ConstEvalFunctionIndex;
using IntrinsicInfo = tint::core::intrinsic::IntrinsicInfo;
using MatcherIndicesIndex = tint::core::intrinsic::MatcherIndicesIndex;
using MatchState = tint::core::intrinsic::MatchState;
using Number = tint::core::intrinsic::Number;
using NumberMatcher = tint::core::intrinsic::NumberMatcher;
using NumberMatcherIndex = tint::core::intrinsic::NumberMatcherIndex;
using OverloadFlag = tint::core::intrinsic::OverloadFlag;
using OverloadFlags = tint::core::intrinsic::OverloadFlags;
using OverloadIndex = tint::core::intrinsic::OverloadIndex;
using OverloadInfo = tint::core::intrinsic::OverloadInfo;
using ParameterIndex = tint::core::intrinsic::ParameterIndex;
using ParameterInfo = tint::core::intrinsic::ParameterInfo;
using StringStream = tint::StringStream;
using TemplateIndex = tint::core::intrinsic::TemplateIndex;
using Type = tint::core::type::Type;
using TypeMatcher = tint::core::intrinsic::TypeMatcher;
using TypeMatcherIndex = tint::core::intrinsic::TypeMatcherIndex;

template <size_t N>
using TemplateNumberMatcher = tint::core::intrinsic::TemplateNumberMatcher<N>;

template <size_t N>
using TemplateTypeMatcher = tint::core::intrinsic::TemplateTypeMatcher<N>;

// clang-format off

/// TypeMatcher for 'type bool'
constexpr TypeMatcher kBoolMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchBool(state, ty)) {
      return nullptr;
    }
    return BuildBool(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("bool");
  }
};


/// TypeMatcher for 'type i32'
constexpr TypeMatcher kI32Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchI32(state, ty)) {
      return nullptr;
    }
    return BuildI32(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("i32");
  }
};


/// TypeMatcher for 'type u32'
constexpr TypeMatcher kU32Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchU32(state, ty)) {
      return nullptr;
    }
    return BuildU32(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("u32");
  }
};


/// TypeMatcher for 'type f32'
constexpr TypeMatcher kF32Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchF32(state, ty)) {
      return nullptr;
    }
    return BuildF32(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("f32");
  }
};


/// TypeMatcher for 'type vec2'
constexpr TypeMatcher kVec2Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchVec2(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildVec2(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("vec2", "<", T, ">");
  }
};


/// TypeMatcher for 'type vec3'
constexpr TypeMatcher kVec3Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchVec3(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildVec3(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("vec3", "<", T, ">");
  }
};


/// TypeMatcher for 'type vec4'
constexpr TypeMatcher kVec4Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchVec4(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildVec4(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("vec4", "<", T, ">");
  }
};


/// TypeMatcher for 'type atomic'
constexpr TypeMatcher kAtomicMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchAtomic(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildAtomic(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("atomic", "<", T, ">");
  }
};


/// TypeMatcher for 'type ptr'
constexpr TypeMatcher kPtrMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  Number S = Number::invalid;
  const Type* T = nullptr;
  Number A = Number::invalid;
    if (!MatchPtr(state, ty, S, T, A)) {
      return nullptr;
    }
    S = state.Num(S);
    if (!S.IsValid()) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    A = state.Num(A);
    if (!A.IsValid()) {
      return nullptr;
    }
    return BuildPtr(state, ty, S, T, A);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText S;
  state->PrintNum(S);StyledText T;
  state->PrintType(T);StyledText A;
  state->PrintNum(A);
    out << style::Type("ptr", "<", S, ", ", T, ", ", A, ">");
  }
};


/// TypeMatcher for 'type sampler'
constexpr TypeMatcher kSamplerMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchSampler(state, ty)) {
      return nullptr;
    }
    return BuildSampler(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("sampler");
  }
};


/// TypeMatcher for 'type texture_1d'
constexpr TypeMatcher kTexture1DMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchTexture1D(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildTexture1D(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("texture_1d", "<", T, ">");
  }
};


/// TypeMatcher for 'type texture_2d'
constexpr TypeMatcher kTexture2DMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchTexture2D(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildTexture2D(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("texture_2d", "<", T, ">");
  }
};


/// TypeMatcher for 'type texture_2d_array'
constexpr TypeMatcher kTexture2DArrayMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchTexture2DArray(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildTexture2DArray(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("texture_2d_array", "<", T, ">");
  }
};


/// TypeMatcher for 'type texture_3d'
constexpr TypeMatcher kTexture3DMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchTexture3D(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildTexture3D(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("texture_3d", "<", T, ">");
  }
};


/// TypeMatcher for 'type texture_cube'
constexpr TypeMatcher kTextureCubeMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchTextureCube(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildTextureCube(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("texture_cube", "<", T, ">");
  }
};


/// TypeMatcher for 'type texture_cube_array'
constexpr TypeMatcher kTextureCubeArrayMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
  const Type* T = nullptr;
    if (!MatchTextureCubeArray(state, ty, T)) {
      return nullptr;
    }
    T = state.Type(T);
    if (T == nullptr) {
      return nullptr;
    }
    return BuildTextureCubeArray(state, ty, T);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
  state->PrintType(T);
    out << style::Type("texture_cube_array", "<", T, ">");
  }
};


/// TypeMatcher for 'type texture_depth_2d'
constexpr TypeMatcher kTextureDepth2DMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchTextureDepth2D(state, ty)) {
      return nullptr;
    }
    return BuildTextureDepth2D(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("texture_depth_2d");
  }
};


/// TypeMatcher for 'type texture_depth_2d_array'
constexpr TypeMatcher kTextureDepth2DArrayMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchTextureDepth2DArray(state, ty)) {
      return nullptr;
    }
    return BuildTextureDepth2DArray(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("texture_depth_2d_array");
  }
};


/// TypeMatcher for 'type texture_depth_cube'
constexpr TypeMatcher kTextureDepthCubeMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchTextureDepthCube(state, ty)) {
      return nullptr;
    }
    return BuildTextureDepthCube(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("texture_depth_cube");
  }
};


/// TypeMatcher for 'type texture_depth_cube_array'
constexpr TypeMatcher kTextureDepthCubeArrayMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (!MatchTextureDepthCubeArray(state, ty)) {
      return nullptr;
    }
    return BuildTextureDepthCubeArray(state, ty);
  },
/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
    out << style::Type("texture_depth_cube_array");
  }
};


/// TypeMatcher for 'match iu32'
constexpr TypeMatcher kIu32Matcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
    if (MatchI32(state, ty)) {
      return BuildI32(state, ty);
    }
    if (MatchU32(state, ty)) {
      return BuildU32(state, ty);
    }
    return nullptr;
  },
/* print */ [](MatchState*, StyledText& out) {
    // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
    // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
 kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
};

/// EnumMatcher for 'match read_write'
constexpr NumberMatcher kReadWriteMatcher {
/* match */ [](MatchState&, Number number) -> Number {
    if (number.IsAny() || number.Value() == static_cast<uint32_t>(core::Access::kReadWrite)) {
      return Number(static_cast<uint32_t>(core::Access::kReadWrite));
    }
    return Number::invalid;
  },
/* print */ [](MatchState*, StyledText& out) {
  out<< style::Enum("read_write");
  }
};

/// EnumMatcher for 'match function'
constexpr NumberMatcher kFunctionMatcher {
/* match */ [](MatchState&, Number number) -> Number {
    if (number.IsAny() || number.Value() == static_cast<uint32_t>(core::AddressSpace::kFunction)) {
      return Number(static_cast<uint32_t>(core::AddressSpace::kFunction));
    }
    return Number::invalid;
  },
/* print */ [](MatchState*, StyledText& out) {
  out<< style::Enum("function");
  }
};

/// EnumMatcher for 'match workgroup_or_storage'
constexpr NumberMatcher kWorkgroupOrStorageMatcher {
/* match */ [](MatchState&, Number number) -> Number {
    switch (static_cast<core::AddressSpace>(number.Value())) {
      case core::AddressSpace::kWorkgroup:
      case core::AddressSpace::kStorage:
        return number;
      default:
        return Number::invalid;
    }
  },
/* print */ [](MatchState*, StyledText& out) {
  out<< style::Enum("workgroup")<< style::Plain(" or ") << style::Enum("storage");
  }
};

/// Type and number matchers

/// The template types, types, and type matchers
constexpr TypeMatcher kTypeMatchers[] = {
  /* [0] */ TemplateTypeMatcher<0>::matcher,
  /* [1] */ TemplateTypeMatcher<1>::matcher,
  /* [2] */ kBoolMatcher,
  /* [3] */ kI32Matcher,
  /* [4] */ kU32Matcher,
  /* [5] */ kF32Matcher,
  /* [6] */ kVec2Matcher,
  /* [7] */ kVec3Matcher,
  /* [8] */ kVec4Matcher,
  /* [9] */ kAtomicMatcher,
  /* [10] */ kPtrMatcher,
  /* [11] */ kSamplerMatcher,
  /* [12] */ kTexture1DMatcher,
  /* [13] */ kTexture2DMatcher,
  /* [14] */ kTexture2DArrayMatcher,
  /* [15] */ kTexture3DMatcher,
  /* [16] */ kTextureCubeMatcher,
  /* [17] */ kTextureCubeArrayMatcher,
  /* [18] */ kTextureDepth2DMatcher,
  /* [19] */ kTextureDepth2DArrayMatcher,
  /* [20] */ kTextureDepthCubeMatcher,
  /* [21] */ kTextureDepthCubeArrayMatcher,
  /* [22] */ kIu32Matcher,
};

/// The template numbers, and number matchers
constexpr NumberMatcher kNumberMatchers[] = {
  /* [0] */ TemplateNumberMatcher<0>::matcher,
  /* [1] */ TemplateNumberMatcher<1>::matcher,
  /* [2] */ kReadWriteMatcher,
  /* [3] */ kFunctionMatcher,
  /* [4] */ kWorkgroupOrStorageMatcher,
};

constexpr MatcherIndex kMatcherIndices[] = {
  /* [0] */ MatcherIndex(10),
  /* [1] */ MatcherIndex(1),
  /* [2] */ MatcherIndex(9),
  /* [3] */ MatcherIndex(0),
  /* [4] */ MatcherIndex(2),
  /* [5] */ MatcherIndex(10),
  /* [6] */ MatcherIndex(3),
  /* [7] */ MatcherIndex(0),
  /* [8] */ MatcherIndex(2),
  /* [9] */ MatcherIndex(8),
  /* [10] */ MatcherIndex(5),
  /* [11] */ MatcherIndex(12),
  /* [12] */ MatcherIndex(5),
  /* [13] */ MatcherIndex(13),
  /* [14] */ MatcherIndex(5),
  /* [15] */ MatcherIndex(6),
  /* [16] */ MatcherIndex(5),
  /* [17] */ MatcherIndex(6),
  /* [18] */ MatcherIndex(3),
  /* [19] */ MatcherIndex(14),
  /* [20] */ MatcherIndex(5),
  /* [21] */ MatcherIndex(15),
  /* [22] */ MatcherIndex(5),
  /* [23] */ MatcherIndex(7),
  /* [24] */ MatcherIndex(5),
  /* [25] */ MatcherIndex(7),
  /* [26] */ MatcherIndex(3),
  /* [27] */ MatcherIndex(16),
  /* [28] */ MatcherIndex(5),
  /* [29] */ MatcherIndex(17),
  /* [30] */ MatcherIndex(5),
  /* [31] */ MatcherIndex(22),
  /* [32] */ MatcherIndex(4),
  /* [33] */ MatcherIndex(11),
  /* [34] */ MatcherIndex(18),
  /* [35] */ MatcherIndex(19),
  /* [36] */ MatcherIndex(20),
  /* [37] */ MatcherIndex(21),
};

static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
              "MatcherIndicesIndex is not large enough to index kMatcherIndices");

constexpr ParameterInfo kParameters[] = {
  {
    /* [0] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(0),
  },
  {
    /* [1] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(5),
  },
  {
    /* [2] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(3),
  },
  {
    /* [3] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(32),
  },
  {
    /* [4] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(32),
  },
  {
    /* [5] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(19),
  },
  {
    /* [6] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [7] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(15),
  },
  {
    /* [8] */
    /* usage */ core::ParameterUsage::kArrayIndex,
    /* matcher_indices */ MatcherIndicesIndex(3),
  },
  {
    /* [9] */
    /* usage */ core::ParameterUsage::kOffset,
    /* matcher_indices */ MatcherIndicesIndex(17),
  },
  {
    /* [10] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(35),
  },
  {
    /* [11] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [12] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(15),
  },
  {
    /* [13] */
    /* usage */ core::ParameterUsage::kArrayIndex,
    /* matcher_indices */ MatcherIndicesIndex(3),
  },
  {
    /* [14] */
    /* usage */ core::ParameterUsage::kOffset,
    /* matcher_indices */ MatcherIndicesIndex(17),
  },
  {
    /* [15] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(13),
  },
  {
    /* [16] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [17] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(15),
  },
  {
    /* [18] */
    /* usage */ core::ParameterUsage::kOffset,
    /* matcher_indices */ MatcherIndicesIndex(17),
  },
  {
    /* [19] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(21),
  },
  {
    /* [20] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [21] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(23),
  },
  {
    /* [22] */
    /* usage */ core::ParameterUsage::kOffset,
    /* matcher_indices */ MatcherIndicesIndex(25),
  },
  {
    /* [23] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(29),
  },
  {
    /* [24] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [25] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(23),
  },
  {
    /* [26] */
    /* usage */ core::ParameterUsage::kArrayIndex,
    /* matcher_indices */ MatcherIndicesIndex(3),
  },
  {
    /* [27] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(34),
  },
  {
    /* [28] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [29] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(15),
  },
  {
    /* [30] */
    /* usage */ core::ParameterUsage::kOffset,
    /* matcher_indices */ MatcherIndicesIndex(17),
  },
  {
    /* [31] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(37),
  },
  {
    /* [32] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [33] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(23),
  },
  {
    /* [34] */
    /* usage */ core::ParameterUsage::kArrayIndex,
    /* matcher_indices */ MatcherIndicesIndex(3),
  },
  {
    /* [35] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(0),
  },
  {
    /* [36] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(3),
  },
  {
    /* [37] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(32),
  },
  {
    /* [38] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(11),
  },
  {
    /* [39] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [40] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(10),
  },
  {
    /* [41] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(27),
  },
  {
    /* [42] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [43] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(23),
  },
  {
    /* [44] */
    /* usage */ core::ParameterUsage::kTexture,
    /* matcher_indices */ MatcherIndicesIndex(36),
  },
  {
    /* [45] */
    /* usage */ core::ParameterUsage::kSampler,
    /* matcher_indices */ MatcherIndicesIndex(33),
  },
  {
    /* [46] */
    /* usage */ core::ParameterUsage::kCoords,
    /* matcher_indices */ MatcherIndicesIndex(23),
  },
  {
    /* [47] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(0),
  },
  {
    /* [48] */
    /* usage */ core::ParameterUsage::kNone,
    /* matcher_indices */ MatcherIndicesIndex(32),
  },
};

static_assert(ParameterIndex::CanIndex(kParameters),
              "ParameterIndex is not large enough to index kParameters");

constexpr TemplateInfo kTemplates[] = {
  {
    /* [0] */
    /* name */ "T",
    /* matcher_indices */ MatcherIndicesIndex(31),
    /* kind */ TemplateInfo::Kind::kType,
  },
  {
    /* [1] */
    /* name */ "S",
    /* matcher_indices */ MatcherIndicesIndex(32),
    /* kind */ TemplateInfo::Kind::kNumber,
  },
  {
    /* [2] */
    /* name */ "A",
    /* matcher_indices */ MatcherIndicesIndex(31),
    /* kind */ TemplateInfo::Kind::kType,
  },
};

static_assert(TemplateIndex::CanIndex(kTemplates),
              "TemplateIndex is not large enough to index kTemplates");

constexpr OverloadInfo kOverloads[] = {
  {
    /* [0] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(38),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [1] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(15),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [2] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 4,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(15),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [3] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 4,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 1,
    /* templates */ TemplateIndex(2),
    /* parameters */ ParameterIndex(5),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [4] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 5,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 1,
    /* templates */ TemplateIndex(2),
    /* parameters */ ParameterIndex(5),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [5] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(19),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [6] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 4,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(19),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [7] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(41),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [8] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 4,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 1,
    /* templates */ TemplateIndex(2),
    /* parameters */ ParameterIndex(23),
    /* return_matcher_indices */ MatcherIndicesIndex(9),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [9] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(27),
    /* return_matcher_indices */ MatcherIndicesIndex(10),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [10] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 4,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(27),
    /* return_matcher_indices */ MatcherIndicesIndex(10),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [11] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 4,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 1,
    /* templates */ TemplateIndex(2),
    /* parameters */ ParameterIndex(10),
    /* return_matcher_indices */ MatcherIndicesIndex(10),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [12] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 5,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 1,
    /* templates */ TemplateIndex(2),
    /* parameters */ ParameterIndex(10),
    /* return_matcher_indices */ MatcherIndicesIndex(10),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [13] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(44),
    /* return_matcher_indices */ MatcherIndicesIndex(10),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [14] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
    /* num_parameters */ 4,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 1,
    /* templates */ TemplateIndex(2),
    /* parameters */ ParameterIndex(31),
    /* return_matcher_indices */ MatcherIndicesIndex(10),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [15] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
    /* num_parameters */ 5,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 2,
    /* templates */ TemplateIndex(0),
    /* parameters */ ParameterIndex(0),
    /* return_matcher_indices */ MatcherIndicesIndex(4),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [16] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 2,
    /* templates */ TemplateIndex(0),
    /* parameters */ ParameterIndex(35),
    /* return_matcher_indices */ MatcherIndicesIndex(3),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [17] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
    /* num_parameters */ 2,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 2,
    /* templates */ TemplateIndex(0),
    /* parameters */ ParameterIndex(47),
    /* return_matcher_indices */ MatcherIndicesIndex(3),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [18] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
    /* num_parameters */ 3,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 2,
    /* templates */ TemplateIndex(0),
    /* parameters */ ParameterIndex(35),
    /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
  {
    /* [19] */
    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
    /* num_parameters */ 1,
    /* num_explicit_templates */ 0,
    /* num_templates   */ 0,
    /* templates */ TemplateIndex(/* invalid */),
    /* parameters */ ParameterIndex(3),
    /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
  },
};

static_assert(OverloadIndex::CanIndex(kOverloads),
              "OverloadIndex is not large enough to index kOverloads");

constexpr IntrinsicInfo kBuiltins[] = {
  {
    /* [0] */
    /* fn atomic_compare_exchange_weak_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, ptr<function, T, read_write>, T, u32, u32) -> bool */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(15),
  },
  {
    /* [1] */
    /* fn atomic_exchange_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [2] */
    /* fn atomic_fetch_add_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [3] */
    /* fn atomic_fetch_and_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [4] */
    /* fn atomic_fetch_max_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [5] */
    /* fn atomic_fetch_min_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [6] */
    /* fn atomic_fetch_or_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [7] */
    /* fn atomic_fetch_sub_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [8] */
    /* fn atomic_fetch_xor_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(16),
  },
  {
    /* [9] */
    /* fn atomic_load_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(17),
  },
  {
    /* [10] */
    /* fn atomic_store_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(18),
  },
  {
    /* [11] */
    /* fn sample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32> */
    /* fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
    /* fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
    /* fn sample[A : iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<f32> */
    /* fn sample[A : iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> vec4<f32> */
    /* fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
    /* fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32> */
    /* fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
    /* fn sample[A : iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32> */
    /* fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32 */
    /* fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32 */
    /* fn sample[A : iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A) -> f32 */
    /* fn sample[A : iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> f32 */
    /* fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32 */
    /* fn sample[A : iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> f32 */
    /* num overloads */ 15,
    /* overloads */ OverloadIndex(0),
  },
  {
    /* [12] */
    /* fn threadgroup_barrier(u32) */
    /* num overloads */ 1,
    /* overloads */ OverloadIndex(19),
  },
};

// clang-format on

}  // anonymous namespace

const core::intrinsic::TableData Dialect::kData{
    /* templates */ kTemplates,
    /* type_matcher_indices */ kMatcherIndices,
    /* type_matchers */ kTypeMatchers,
    /* number_matchers */ kNumberMatchers,
    /* parameters */ kParameters,
    /* overloads */ kOverloads,
    /* const_eval_functions */ Empty,
    /* ctor_conv */ Empty,
    /* builtins */ kBuiltins,
    /* binary '+' */ tint::core::intrinsic::kNoOverloads,
    /* binary '-' */ tint::core::intrinsic::kNoOverloads,
    /* binary '*' */ tint::core::intrinsic::kNoOverloads,
    /* binary '/' */ tint::core::intrinsic::kNoOverloads,
    /* binary '%' */ tint::core::intrinsic::kNoOverloads,
    /* binary '^' */ tint::core::intrinsic::kNoOverloads,
    /* binary '&' */ tint::core::intrinsic::kNoOverloads,
    /* binary '|' */ tint::core::intrinsic::kNoOverloads,
    /* binary '&&' */ tint::core::intrinsic::kNoOverloads,
    /* binary '||' */ tint::core::intrinsic::kNoOverloads,
    /* binary '==' */ tint::core::intrinsic::kNoOverloads,
    /* binary '!=' */ tint::core::intrinsic::kNoOverloads,
    /* binary '<' */ tint::core::intrinsic::kNoOverloads,
    /* binary '>' */ tint::core::intrinsic::kNoOverloads,
    /* binary '<=' */ tint::core::intrinsic::kNoOverloads,
    /* binary '>=' */ tint::core::intrinsic::kNoOverloads,
    /* binary '<<' */ tint::core::intrinsic::kNoOverloads,
    /* binary '>>' */ tint::core::intrinsic::kNoOverloads,
    /* unary '!' */ tint::core::intrinsic::kNoOverloads,
    /* unary '~' */ tint::core::intrinsic::kNoOverloads,
    /* unary '-' */ tint::core::intrinsic::kNoOverloads,
    /* unary '*' */ tint::core::intrinsic::kNoOverloads,
    /* unary '&' */ tint::core::intrinsic::kNoOverloads,
};

}  // namespace tint::msl::intrinsic