spirv-reader: use spirv.hpp11
In preparation for SPIRV-Tools change where its internals
use the C++11 headers.
This patch works with SPIRV-Tools using the old C header
and using the C++11 header.
This patch includes some complex machinery inside "three_sided_patch"
namespaces that can be removed after third_party/vulkan-deps/spirv-tools has
fully transitioned into using the C++11 headers.
Change-Id: I36f358fe3edcc5e613625708017fb8d7919c40c6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/108780
Reviewed-by: Alan Baker <alanbaker@google.com>
Commit-Queue: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/reader/spirv/enum_converter.cc b/src/tint/reader/spirv/enum_converter.cc
index 923f250..d590507 100644
--- a/src/tint/reader/spirv/enum_converter.cc
+++ b/src/tint/reader/spirv/enum_converter.cc
@@ -20,13 +20,13 @@
EnumConverter::~EnumConverter() = default;
-ast::PipelineStage EnumConverter::ToPipelineStage(SpvExecutionModel model) {
+ast::PipelineStage EnumConverter::ToPipelineStage(spv::ExecutionModel model) {
switch (model) {
- case SpvExecutionModelVertex:
+ case spv::ExecutionModel::Vertex:
return ast::PipelineStage::kVertex;
- case SpvExecutionModelFragment:
+ case spv::ExecutionModel::Fragment:
return ast::PipelineStage::kFragment;
- case SpvExecutionModelGLCompute:
+ case spv::ExecutionModel::GLCompute:
return ast::PipelineStage::kCompute;
default:
break;
@@ -36,23 +36,23 @@
return ast::PipelineStage::kNone;
}
-ast::AddressSpace EnumConverter::ToAddressSpace(const SpvStorageClass sc) {
+ast::AddressSpace EnumConverter::ToAddressSpace(const spv::StorageClass sc) {
switch (sc) {
- case SpvStorageClassInput:
+ case spv::StorageClass::Input:
return ast::AddressSpace::kIn;
- case SpvStorageClassOutput:
+ case spv::StorageClass::Output:
return ast::AddressSpace::kOut;
- case SpvStorageClassUniform:
+ case spv::StorageClass::Uniform:
return ast::AddressSpace::kUniform;
- case SpvStorageClassWorkgroup:
+ case spv::StorageClass::Workgroup:
return ast::AddressSpace::kWorkgroup;
- case SpvStorageClassUniformConstant:
+ case spv::StorageClass::UniformConstant:
return ast::AddressSpace::kNone;
- case SpvStorageClassStorageBuffer:
+ case spv::StorageClass::StorageBuffer:
return ast::AddressSpace::kStorage;
- case SpvStorageClassPrivate:
+ case spv::StorageClass::Private:
return ast::AddressSpace::kPrivate;
- case SpvStorageClassFunction:
+ case spv::StorageClass::Function:
return ast::AddressSpace::kFunction;
default:
break;
@@ -62,31 +62,31 @@
return ast::AddressSpace::kUndefined;
}
-ast::BuiltinValue EnumConverter::ToBuiltin(SpvBuiltIn b) {
+ast::BuiltinValue EnumConverter::ToBuiltin(spv::BuiltIn b) {
switch (b) {
- case SpvBuiltInPosition:
+ case spv::BuiltIn::Position:
return ast::BuiltinValue::kPosition;
- case SpvBuiltInVertexIndex:
+ case spv::BuiltIn::VertexIndex:
return ast::BuiltinValue::kVertexIndex;
- case SpvBuiltInInstanceIndex:
+ case spv::BuiltIn::InstanceIndex:
return ast::BuiltinValue::kInstanceIndex;
- case SpvBuiltInFrontFacing:
+ case spv::BuiltIn::FrontFacing:
return ast::BuiltinValue::kFrontFacing;
- case SpvBuiltInFragCoord:
+ case spv::BuiltIn::FragCoord:
return ast::BuiltinValue::kPosition;
- case SpvBuiltInFragDepth:
+ case spv::BuiltIn::FragDepth:
return ast::BuiltinValue::kFragDepth;
- case SpvBuiltInLocalInvocationId:
+ case spv::BuiltIn::LocalInvocationId:
return ast::BuiltinValue::kLocalInvocationId;
- case SpvBuiltInLocalInvocationIndex:
+ case spv::BuiltIn::LocalInvocationIndex:
return ast::BuiltinValue::kLocalInvocationIndex;
- case SpvBuiltInGlobalInvocationId:
+ case spv::BuiltIn::GlobalInvocationId:
return ast::BuiltinValue::kGlobalInvocationId;
- case SpvBuiltInWorkgroupId:
+ case spv::BuiltIn::WorkgroupId:
return ast::BuiltinValue::kWorkgroupId;
- case SpvBuiltInSampleId:
+ case spv::BuiltIn::SampleId:
return ast::BuiltinValue::kSampleIndex;
- case SpvBuiltInSampleMask:
+ case spv::BuiltIn::SampleMask:
return ast::BuiltinValue::kSampleMask;
default:
break;
@@ -96,12 +96,12 @@
return ast::BuiltinValue::kUndefined;
}
-ast::TextureDimension EnumConverter::ToDim(SpvDim dim, bool arrayed) {
+ast::TextureDimension EnumConverter::ToDim(spv::Dim dim, bool arrayed) {
if (arrayed) {
switch (dim) {
- case SpvDim2D:
+ case spv::Dim::Dim2D:
return ast::TextureDimension::k2dArray;
- case SpvDimCube:
+ case spv::Dim::Cube:
return ast::TextureDimension::kCubeArray;
default:
break;
@@ -111,13 +111,13 @@
}
// Assume non-arrayed
switch (dim) {
- case SpvDim1D:
+ case spv::Dim::Dim1D:
return ast::TextureDimension::k1d;
- case SpvDim2D:
+ case spv::Dim::Dim2D:
return ast::TextureDimension::k2d;
- case SpvDim3D:
+ case spv::Dim::Dim3D:
return ast::TextureDimension::k3d;
- case SpvDimCube:
+ case spv::Dim::Cube:
return ast::TextureDimension::kCube;
default:
break;
@@ -126,47 +126,47 @@
return ast::TextureDimension::kNone;
}
-ast::TexelFormat EnumConverter::ToTexelFormat(SpvImageFormat fmt) {
+ast::TexelFormat EnumConverter::ToTexelFormat(spv::ImageFormat fmt) {
switch (fmt) {
- case SpvImageFormatUnknown:
+ case spv::ImageFormat::Unknown:
return ast::TexelFormat::kUndefined;
// 8 bit channels
- case SpvImageFormatRgba8:
+ case spv::ImageFormat::Rgba8:
return ast::TexelFormat::kRgba8Unorm;
- case SpvImageFormatRgba8Snorm:
+ case spv::ImageFormat::Rgba8Snorm:
return ast::TexelFormat::kRgba8Snorm;
- case SpvImageFormatRgba8ui:
+ case spv::ImageFormat::Rgba8ui:
return ast::TexelFormat::kRgba8Uint;
- case SpvImageFormatRgba8i:
+ case spv::ImageFormat::Rgba8i:
return ast::TexelFormat::kRgba8Sint;
// 16 bit channels
- case SpvImageFormatRgba16ui:
+ case spv::ImageFormat::Rgba16ui:
return ast::TexelFormat::kRgba16Uint;
- case SpvImageFormatRgba16i:
+ case spv::ImageFormat::Rgba16i:
return ast::TexelFormat::kRgba16Sint;
- case SpvImageFormatRgba16f:
+ case spv::ImageFormat::Rgba16f:
return ast::TexelFormat::kRgba16Float;
// 32 bit channels
- case SpvImageFormatR32ui:
+ case spv::ImageFormat::R32ui:
return ast::TexelFormat::kR32Uint;
- case SpvImageFormatR32i:
+ case spv::ImageFormat::R32i:
return ast::TexelFormat::kR32Sint;
- case SpvImageFormatR32f:
+ case spv::ImageFormat::R32f:
return ast::TexelFormat::kR32Float;
- case SpvImageFormatRg32ui:
+ case spv::ImageFormat::Rg32ui:
return ast::TexelFormat::kRg32Uint;
- case SpvImageFormatRg32i:
+ case spv::ImageFormat::Rg32i:
return ast::TexelFormat::kRg32Sint;
- case SpvImageFormatRg32f:
+ case spv::ImageFormat::Rg32f:
return ast::TexelFormat::kRg32Float;
- case SpvImageFormatRgba32ui:
+ case spv::ImageFormat::Rgba32ui:
return ast::TexelFormat::kRgba32Uint;
- case SpvImageFormatRgba32i:
+ case spv::ImageFormat::Rgba32i:
return ast::TexelFormat::kRgba32Sint;
- case SpvImageFormatRgba32f:
+ case spv::ImageFormat::Rgba32f:
return ast::TexelFormat::kRgba32Float;
default:
break;
diff --git a/src/tint/reader/spirv/enum_converter.h b/src/tint/reader/spirv/enum_converter.h
index 1ebea36..0082499 100644
--- a/src/tint/reader/spirv/enum_converter.h
+++ b/src/tint/reader/spirv/enum_converter.h
@@ -16,6 +16,7 @@
#define SRC_TINT_READER_SPIRV_ENUM_CONVERTER_H_
#include "spirv/unified1/spirv.h"
+#include "spirv/unified1/spirv.hpp11"
#include "src/tint/ast/address_space.h"
#include "src/tint/ast/builtin_value.h"
#include "src/tint/ast/pipeline_stage.h"
@@ -37,32 +38,49 @@
/// On failure, logs an error and returns kNone
/// @param model the SPIR-V entry point execution model
/// @returns a Tint AST pipeline stage
- ast::PipelineStage ToPipelineStage(SpvExecutionModel model);
+ ast::PipelineStage ToPipelineStage(spv::ExecutionModel model);
/// Converts a SPIR-V storage class to a Tint address space.
/// On failure, logs an error and returns kNone
/// @param sc the SPIR-V storage class
/// @returns a Tint AST address space
- ast::AddressSpace ToAddressSpace(const SpvStorageClass sc);
+ ast::AddressSpace ToAddressSpace(const spv::StorageClass sc);
/// Converts a SPIR-V Builtin value a Tint Builtin.
/// On failure, logs an error and returns kNone
/// @param b the SPIR-V builtin
/// @returns a Tint AST builtin
- ast::BuiltinValue ToBuiltin(SpvBuiltIn b);
+ ast::BuiltinValue ToBuiltin(spv::BuiltIn b);
/// Converts a possibly arrayed SPIR-V Dim to a Tint texture dimension.
/// On failure, logs an error and returns kNone
/// @param dim the SPIR-V Dim value
/// @param arrayed true if the texture is arrayed
/// @returns a Tint AST texture dimension
- ast::TextureDimension ToDim(SpvDim dim, bool arrayed);
+ ast::TextureDimension ToDim(spv::Dim dim, bool arrayed);
+
+ /// Converts a possibly arrayed SPIR-V Dim to a Tint texture dimension.
+ /// On failure, logs an error and returns kNone
+ /// @param dim the SPIR-V Dim value
+ /// @param arrayed true if the texture is arrayed
+ /// @returns a Tint AST texture dimension
+ ast::TextureDimension ToDim(SpvDim dim, bool arrayed) {
+ return ToDim(static_cast<spv::Dim>(dim), arrayed);
+ }
/// Converts a SPIR-V Image Format to a TexelFormat
/// On failure, logs an error and returns kNone
/// @param fmt the SPIR-V format
/// @returns a Tint AST format
- ast::TexelFormat ToTexelFormat(SpvImageFormat fmt);
+ ast::TexelFormat ToTexelFormat(spv::ImageFormat fmt);
+
+ /// Converts a SPIR-V Image Format to a TexelFormat
+ /// On failure, logs an error and returns kNone
+ /// @param fmt the SPIR-V format
+ /// @returns a Tint AST format
+ ast::TexelFormat ToTexelFormat(SpvImageFormat fmt) {
+ return ToTexelFormat(static_cast<spv::ImageFormat>(fmt));
+ }
private:
/// Registers a failure and returns a stream for log diagnostics.
diff --git a/src/tint/reader/spirv/enum_converter_test.cc b/src/tint/reader/spirv/enum_converter_test.cc
index 112dc4a..b366ebd 100644
--- a/src/tint/reader/spirv/enum_converter_test.cc
+++ b/src/tint/reader/spirv/enum_converter_test.cc
@@ -24,12 +24,12 @@
// Pipeline stage
struct PipelineStageCase {
- SpvExecutionModel model;
+ spv::ExecutionModel model;
bool expect_success;
ast::PipelineStage expected;
};
inline std::ostream& operator<<(std::ostream& out, PipelineStageCase psc) {
- out << "PipelineStageCase{ SpvExecutionModel:" << int(psc.model)
+ out << "PipelineStageCase{ spv::ExecutionModel:::" << int(psc.model)
<< " expect_success?:" << int(psc.expect_success) << " expected:" << int(psc.expected)
<< "}";
return out;
@@ -65,29 +65,29 @@
INSTANTIATE_TEST_SUITE_P(EnumConverterGood,
SpvPipelineStageTest,
- testing::Values(PipelineStageCase{SpvExecutionModelVertex, true,
+ testing::Values(PipelineStageCase{spv::ExecutionModel::Vertex, true,
ast::PipelineStage::kVertex},
- PipelineStageCase{SpvExecutionModelFragment, true,
+ PipelineStageCase{spv::ExecutionModel::Fragment, true,
ast::PipelineStage::kFragment},
- PipelineStageCase{SpvExecutionModelGLCompute, true,
+ PipelineStageCase{spv::ExecutionModel::GLCompute, true,
ast::PipelineStage::kCompute}));
INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
SpvPipelineStageTest,
- testing::Values(PipelineStageCase{static_cast<SpvExecutionModel>(9999),
+ testing::Values(PipelineStageCase{static_cast<spv::ExecutionModel>(9999),
false, ast::PipelineStage::kNone},
- PipelineStageCase{SpvExecutionModelTessellationControl,
+ PipelineStageCase{spv::ExecutionModel::TessellationControl,
false, ast::PipelineStage::kNone}));
// Storage class
struct StorageClassCase {
- SpvStorageClass sc;
+ spv::StorageClass sc;
bool expect_success;
ast::AddressSpace expected;
};
inline std::ostream& operator<<(std::ostream& out, StorageClassCase scc) {
- out << "StorageClassCase{ SpvStorageClass:" << int(scc.sc)
+ out << "StorageClassCase{ spv::StorageClass:::" << int(scc.sc)
<< " expect_success?:" << int(scc.expect_success) << " expected:" << int(scc.expected)
<< "}";
return out;
@@ -125,29 +125,29 @@
EnumConverterGood,
SpvStorageClassTest,
testing::Values(
- StorageClassCase{SpvStorageClassInput, true, ast::AddressSpace::kIn},
- StorageClassCase{SpvStorageClassOutput, true, ast::AddressSpace::kOut},
- StorageClassCase{SpvStorageClassUniform, true, ast::AddressSpace::kUniform},
- StorageClassCase{SpvStorageClassWorkgroup, true, ast::AddressSpace::kWorkgroup},
- StorageClassCase{SpvStorageClassUniformConstant, true, ast::AddressSpace::kNone},
- StorageClassCase{SpvStorageClassStorageBuffer, true, ast::AddressSpace::kStorage},
- StorageClassCase{SpvStorageClassPrivate, true, ast::AddressSpace::kPrivate},
- StorageClassCase{SpvStorageClassFunction, true, ast::AddressSpace::kFunction}));
+ StorageClassCase{spv::StorageClass::Input, true, ast::AddressSpace::kIn},
+ StorageClassCase{spv::StorageClass::Output, true, ast::AddressSpace::kOut},
+ StorageClassCase{spv::StorageClass::Uniform, true, ast::AddressSpace::kUniform},
+ StorageClassCase{spv::StorageClass::Workgroup, true, ast::AddressSpace::kWorkgroup},
+ StorageClassCase{spv::StorageClass::UniformConstant, true, ast::AddressSpace::kNone},
+ StorageClassCase{spv::StorageClass::StorageBuffer, true, ast::AddressSpace::kStorage},
+ StorageClassCase{spv::StorageClass::Private, true, ast::AddressSpace::kPrivate},
+ StorageClassCase{spv::StorageClass::Function, true, ast::AddressSpace::kFunction}));
INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
SpvStorageClassTest,
- testing::Values(StorageClassCase{static_cast<SpvStorageClass>(9999), false,
- ast::AddressSpace::kUndefined}));
+ testing::Values(StorageClassCase{static_cast<spv::StorageClass>(9999),
+ false, ast::AddressSpace::kUndefined}));
// Builtin
struct BuiltinCase {
- SpvBuiltIn builtin;
+ spv::BuiltIn builtin;
bool expect_success;
ast::BuiltinValue expected;
};
inline std::ostream& operator<<(std::ostream& out, BuiltinCase bc) {
- out << "BuiltinCase{ SpvBuiltIn:" << int(bc.builtin)
+ out << "BuiltinCase{ spv::BuiltIn::" << int(bc.builtin)
<< " expect_success?:" << int(bc.expect_success) << " expected:" << int(bc.expected) << "}";
return out;
}
@@ -184,43 +184,44 @@
EnumConverterGood_Input,
SpvBuiltinTest,
testing::Values(
- BuiltinCase{SpvBuiltInPosition, true, ast::BuiltinValue::kPosition},
- BuiltinCase{SpvBuiltInInstanceIndex, true, ast::BuiltinValue::kInstanceIndex},
- BuiltinCase{SpvBuiltInFrontFacing, true, ast::BuiltinValue::kFrontFacing},
- BuiltinCase{SpvBuiltInFragCoord, true, ast::BuiltinValue::kPosition},
- BuiltinCase{SpvBuiltInLocalInvocationId, true, ast::BuiltinValue::kLocalInvocationId},
- BuiltinCase{SpvBuiltInLocalInvocationIndex, true, ast::BuiltinValue::kLocalInvocationIndex},
- BuiltinCase{SpvBuiltInGlobalInvocationId, true, ast::BuiltinValue::kGlobalInvocationId},
- BuiltinCase{SpvBuiltInWorkgroupId, true, ast::BuiltinValue::kWorkgroupId},
- BuiltinCase{SpvBuiltInSampleId, true, ast::BuiltinValue::kSampleIndex},
- BuiltinCase{SpvBuiltInSampleMask, true, ast::BuiltinValue::kSampleMask}));
+ BuiltinCase{spv::BuiltIn::Position, true, ast::BuiltinValue::kPosition},
+ BuiltinCase{spv::BuiltIn::InstanceIndex, true, ast::BuiltinValue::kInstanceIndex},
+ BuiltinCase{spv::BuiltIn::FrontFacing, true, ast::BuiltinValue::kFrontFacing},
+ BuiltinCase{spv::BuiltIn::FragCoord, true, ast::BuiltinValue::kPosition},
+ BuiltinCase{spv::BuiltIn::LocalInvocationId, true, ast::BuiltinValue::kLocalInvocationId},
+ BuiltinCase{spv::BuiltIn::LocalInvocationIndex, true,
+ ast::BuiltinValue::kLocalInvocationIndex},
+ BuiltinCase{spv::BuiltIn::GlobalInvocationId, true, ast::BuiltinValue::kGlobalInvocationId},
+ BuiltinCase{spv::BuiltIn::WorkgroupId, true, ast::BuiltinValue::kWorkgroupId},
+ BuiltinCase{spv::BuiltIn::SampleId, true, ast::BuiltinValue::kSampleIndex},
+ BuiltinCase{spv::BuiltIn::SampleMask, true, ast::BuiltinValue::kSampleMask}));
INSTANTIATE_TEST_SUITE_P(
EnumConverterGood_Output,
SpvBuiltinTest,
- testing::Values(BuiltinCase{SpvBuiltInPosition, true, ast::BuiltinValue::kPosition},
- BuiltinCase{SpvBuiltInFragDepth, true, ast::BuiltinValue::kFragDepth},
- BuiltinCase{SpvBuiltInSampleMask, true, ast::BuiltinValue::kSampleMask}));
+ testing::Values(BuiltinCase{spv::BuiltIn::Position, true, ast::BuiltinValue::kPosition},
+ BuiltinCase{spv::BuiltIn::FragDepth, true, ast::BuiltinValue::kFragDepth},
+ BuiltinCase{spv::BuiltIn::SampleMask, true, ast::BuiltinValue::kSampleMask}));
INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
SpvBuiltinTest,
- testing::Values(BuiltinCase{static_cast<SpvBuiltIn>(9999), false,
+ testing::Values(BuiltinCase{static_cast<spv::BuiltIn>(9999), false,
ast::BuiltinValue::kUndefined},
- BuiltinCase{static_cast<SpvBuiltIn>(9999), false,
+ BuiltinCase{static_cast<spv::BuiltIn>(9999), false,
ast::BuiltinValue::kUndefined},
- BuiltinCase{SpvBuiltInNumWorkgroups, false,
+ BuiltinCase{spv::BuiltIn::NumWorkgroups, false,
ast::BuiltinValue::kUndefined}));
// Dim
struct DimCase {
- SpvDim dim;
+ spv::Dim dim;
bool arrayed;
bool expect_success;
ast::TextureDimension expected;
};
inline std::ostream& operator<<(std::ostream& out, DimCase dc) {
- out << "DimCase{ SpvDim:" << int(dc.dim) << " arrayed?:" << int(dc.arrayed)
+ out << "DimCase{ spv::Dim:::" << int(dc.dim) << " arrayed?:" << int(dc.arrayed)
<< " expect_success?:" << int(dc.expect_success) << " expected:" << int(dc.expected) << "}";
return out;
}
@@ -256,40 +257,41 @@
SpvDimTest,
testing::Values(
// Non-arrayed
- DimCase{SpvDim1D, false, true, ast::TextureDimension::k1d},
- DimCase{SpvDim2D, false, true, ast::TextureDimension::k2d},
- DimCase{SpvDim3D, false, true, ast::TextureDimension::k3d},
- DimCase{SpvDimCube, false, true, ast::TextureDimension::kCube},
+ DimCase{spv::Dim::Dim1D, false, true, ast::TextureDimension::k1d},
+ DimCase{spv::Dim::Dim2D, false, true, ast::TextureDimension::k2d},
+ DimCase{spv::Dim::Dim3D, false, true, ast::TextureDimension::k3d},
+ DimCase{spv::Dim::Cube, false, true, ast::TextureDimension::kCube},
// Arrayed
- DimCase{SpvDim2D, true, true, ast::TextureDimension::k2dArray},
- DimCase{SpvDimCube, true, true, ast::TextureDimension::kCubeArray}));
+ DimCase{spv::Dim::Dim2D, true, true, ast::TextureDimension::k2dArray},
+ DimCase{spv::Dim::Cube, true, true,
+ ast::TextureDimension::kCubeArray}));
-INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
- SpvDimTest,
- testing::Values(
- // Invalid SPIR-V dimensionality.
- DimCase{SpvDimMax, false, false, ast::TextureDimension::kNone},
- DimCase{SpvDimMax, true, false, ast::TextureDimension::kNone},
- // Vulkan non-arrayed dimensionalities not supported by WGSL.
- DimCase{SpvDimRect, false, false, ast::TextureDimension::kNone},
- DimCase{SpvDimBuffer, false, false, ast::TextureDimension::kNone},
- DimCase{SpvDimSubpassData, false, false, ast::TextureDimension::kNone},
- // Arrayed dimensionalities not supported by WGSL
- DimCase{SpvDim3D, true, false, ast::TextureDimension::kNone},
- DimCase{SpvDimRect, true, false, ast::TextureDimension::kNone},
- DimCase{SpvDimBuffer, true, false, ast::TextureDimension::kNone},
- DimCase{SpvDimSubpassData, true, false,
- ast::TextureDimension::kNone}));
+INSTANTIATE_TEST_SUITE_P(
+ EnumConverterBad,
+ SpvDimTest,
+ testing::Values(
+ // Invalid SPIR-V dimensionality.
+ DimCase{spv::Dim::Max, false, false, ast::TextureDimension::kNone},
+ DimCase{spv::Dim::Max, true, false, ast::TextureDimension::kNone},
+ // Vulkan non-arrayed dimensionalities not supported by WGSL.
+ DimCase{spv::Dim::Rect, false, false, ast::TextureDimension::kNone},
+ DimCase{spv::Dim::Buffer, false, false, ast::TextureDimension::kNone},
+ DimCase{spv::Dim::SubpassData, false, false, ast::TextureDimension::kNone},
+ // Arrayed dimensionalities not supported by WGSL
+ DimCase{spv::Dim::Dim3D, true, false, ast::TextureDimension::kNone},
+ DimCase{spv::Dim::Rect, true, false, ast::TextureDimension::kNone},
+ DimCase{spv::Dim::Buffer, true, false, ast::TextureDimension::kNone},
+ DimCase{spv::Dim::SubpassData, true, false, ast::TextureDimension::kNone}));
// TexelFormat
struct TexelFormatCase {
- SpvImageFormat format;
+ spv::ImageFormat format;
bool expect_success;
ast::TexelFormat expected;
};
inline std::ostream& operator<<(std::ostream& out, TexelFormatCase ifc) {
- out << "TexelFormatCase{ SpvImageFormat:" << int(ifc.format)
+ out << "TexelFormatCase{ spv::ImageFormat:::" << int(ifc.format)
<< " expect_success?:" << int(ifc.expect_success) << " expected:" << int(ifc.expected)
<< "}";
return out;
@@ -328,52 +330,52 @@
SpvImageFormatTest,
testing::Values(
// Unknown. This is used for sampled images.
- TexelFormatCase{SpvImageFormatUnknown, true, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Unknown, true, ast::TexelFormat::kUndefined},
// 8 bit channels
- TexelFormatCase{SpvImageFormatRgba8, true, ast::TexelFormat::kRgba8Unorm},
- TexelFormatCase{SpvImageFormatRgba8Snorm, true, ast::TexelFormat::kRgba8Snorm},
- TexelFormatCase{SpvImageFormatRgba8ui, true, ast::TexelFormat::kRgba8Uint},
- TexelFormatCase{SpvImageFormatRgba8i, true, ast::TexelFormat::kRgba8Sint},
+ TexelFormatCase{spv::ImageFormat::Rgba8, true, ast::TexelFormat::kRgba8Unorm},
+ TexelFormatCase{spv::ImageFormat::Rgba8Snorm, true, ast::TexelFormat::kRgba8Snorm},
+ TexelFormatCase{spv::ImageFormat::Rgba8ui, true, ast::TexelFormat::kRgba8Uint},
+ TexelFormatCase{spv::ImageFormat::Rgba8i, true, ast::TexelFormat::kRgba8Sint},
// 16 bit channels
- TexelFormatCase{SpvImageFormatRgba16ui, true, ast::TexelFormat::kRgba16Uint},
- TexelFormatCase{SpvImageFormatRgba16i, true, ast::TexelFormat::kRgba16Sint},
- TexelFormatCase{SpvImageFormatRgba16f, true, ast::TexelFormat::kRgba16Float},
+ TexelFormatCase{spv::ImageFormat::Rgba16ui, true, ast::TexelFormat::kRgba16Uint},
+ TexelFormatCase{spv::ImageFormat::Rgba16i, true, ast::TexelFormat::kRgba16Sint},
+ TexelFormatCase{spv::ImageFormat::Rgba16f, true, ast::TexelFormat::kRgba16Float},
// 32 bit channels
// ... 1 channel
- TexelFormatCase{SpvImageFormatR32ui, true, ast::TexelFormat::kR32Uint},
- TexelFormatCase{SpvImageFormatR32i, true, ast::TexelFormat::kR32Sint},
- TexelFormatCase{SpvImageFormatR32f, true, ast::TexelFormat::kR32Float},
+ TexelFormatCase{spv::ImageFormat::R32ui, true, ast::TexelFormat::kR32Uint},
+ TexelFormatCase{spv::ImageFormat::R32i, true, ast::TexelFormat::kR32Sint},
+ TexelFormatCase{spv::ImageFormat::R32f, true, ast::TexelFormat::kR32Float},
// ... 2 channels
- TexelFormatCase{SpvImageFormatRg32ui, true, ast::TexelFormat::kRg32Uint},
- TexelFormatCase{SpvImageFormatRg32i, true, ast::TexelFormat::kRg32Sint},
- TexelFormatCase{SpvImageFormatRg32f, true, ast::TexelFormat::kRg32Float},
+ TexelFormatCase{spv::ImageFormat::Rg32ui, true, ast::TexelFormat::kRg32Uint},
+ TexelFormatCase{spv::ImageFormat::Rg32i, true, ast::TexelFormat::kRg32Sint},
+ TexelFormatCase{spv::ImageFormat::Rg32f, true, ast::TexelFormat::kRg32Float},
// ... 4 channels
- TexelFormatCase{SpvImageFormatRgba32ui, true, ast::TexelFormat::kRgba32Uint},
- TexelFormatCase{SpvImageFormatRgba32i, true, ast::TexelFormat::kRgba32Sint},
- TexelFormatCase{SpvImageFormatRgba32f, true, ast::TexelFormat::kRgba32Float}));
+ TexelFormatCase{spv::ImageFormat::Rgba32ui, true, ast::TexelFormat::kRgba32Uint},
+ TexelFormatCase{spv::ImageFormat::Rgba32i, true, ast::TexelFormat::kRgba32Sint},
+ TexelFormatCase{spv::ImageFormat::Rgba32f, true, ast::TexelFormat::kRgba32Float}));
INSTANTIATE_TEST_SUITE_P(
EnumConverterBad,
SpvImageFormatTest,
testing::Values(
// Scanning in order from the SPIR-V spec.
- TexelFormatCase{SpvImageFormatRg16f, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatR11fG11fB10f, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatR16f, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRgb10A2, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg16, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg8, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatR16, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatR8, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRgba16Snorm, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg16Snorm, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg8Snorm, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg16i, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg8i, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatR8i, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRgb10a2ui, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg16ui, false, ast::TexelFormat::kUndefined},
- TexelFormatCase{SpvImageFormatRg8ui, false, ast::TexelFormat::kUndefined}));
+ TexelFormatCase{spv::ImageFormat::Rg16f, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::R11fG11fB10f, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::R16f, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rgb10A2, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg16, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg8, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::R16, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::R8, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rgba16Snorm, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg16Snorm, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg8Snorm, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg16i, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg8i, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::R8i, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rgb10a2ui, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg16ui, false, ast::TexelFormat::kUndefined},
+ TexelFormatCase{spv::ImageFormat::Rg8ui, false, ast::TexelFormat::kUndefined}));
} // namespace
} // namespace tint::reader::spirv
diff --git a/src/tint/reader/spirv/function.cc b/src/tint/reader/spirv/function.cc
index d3febf2..0aafdf9 100644
--- a/src/tint/reader/spirv/function.cc
+++ b/src/tint/reader/spirv/function.cc
@@ -144,24 +144,99 @@
namespace tint::reader::spirv {
+namespace three_sided_patch_function_cc {
+// This machinery is only used while SPIRV-Tools is in transition before it fully
+// uses the C++11 header spirv.hpp11
+
+/// Typedef for pointer to member function while the API call uses
+/// SpvStorageClass for its second argument.
+typedef uint32_t (
+ spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClass)(uint32_t, SpvStorageClass);
+/// Typedef for pointer to member function while the API call uses
+/// spv::StorageClass for its second argument.
+typedef uint32_t (spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClassCpp11)(
+ uint32_t,
+ spv::StorageClass);
+
+/// @param type_manager the SPIRV-Tools optimizer's type manager
+/// @param finder a pointer to member function in the type manager that does the
+/// actual lookup
+/// @param pointee_type_id the ID of the pointee type
+/// @param sc the storage class. SC can be SpvStorageClass or spv::StorageClass
+/// @returns the ID for a SPIR-V pointer to pointee_type_id in storage class sc
+template <typename FinderType, typename SC>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_manager,
+ FinderType finder,
+ uint32_t pointee_type_id,
+ SC sc);
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClass finder,
+ uint32_t pointee_type_id,
+ SpvStorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, sc);
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClass finder,
+ uint32_t pointee_type_id,
+ spv::StorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, static_cast<SpvStorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClassCpp11 finder,
+ uint32_t pointee_type_id,
+ SpvStorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, static_cast<spv::StorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClassCpp11 finder,
+ uint32_t pointee_type_id,
+ spv::StorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, sc);
+}
+} // namespace three_sided_patch_function_cc
+
namespace {
constexpr uint32_t kMaxVectorLen = 4;
+template <typename FromOpcodeType>
+spv::Op ToOpcode(FromOpcodeType oc) {
+ return static_cast<spv::Op>(oc);
+}
+
+/// @param inst a SPIR-V instruction
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction& inst) {
+ return ToOpcode(inst.opcode());
+}
+/// @param inst a SPIR-V instruction pointer
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction* inst) {
+ return ToOpcode(inst->opcode());
+}
+
// Gets the AST unary opcode for the given SPIR-V opcode, if any
// @param opcode SPIR-V opcode
// @param ast_unary_op return parameter
// @returns true if it was a unary operation
-bool GetUnaryOp(SpvOp opcode, ast::UnaryOp* ast_unary_op) {
+bool GetUnaryOp(spv::Op opcode, ast::UnaryOp* ast_unary_op) {
switch (opcode) {
- case SpvOpSNegate:
- case SpvOpFNegate:
+ case spv::Op::OpSNegate:
+ case spv::Op::OpFNegate:
*ast_unary_op = ast::UnaryOp::kNegation;
return true;
- case SpvOpLogicalNot:
+ case spv::Op::OpLogicalNot:
*ast_unary_op = ast::UnaryOp::kNot;
return true;
- case SpvOpNot:
+ case spv::Op::OpNot:
*ast_unary_op = ast::UnaryOp::kComplement;
return true;
default:
@@ -173,17 +248,17 @@
/// Converts a SPIR-V opcode for a WGSL builtin function, if there is a
/// direct translation. Returns nullptr otherwise.
/// @returns the WGSL builtin function name for the given opcode, or nullptr.
-const char* GetUnaryBuiltInFunctionName(SpvOp opcode) {
+const char* GetUnaryBuiltInFunctionName(spv::Op opcode) {
switch (opcode) {
- case SpvOpAny:
+ case spv::Op::OpAny:
return "any";
- case SpvOpAll:
+ case spv::Op::OpAll:
return "all";
- case SpvOpIsNan:
+ case spv::Op::OpIsNan:
return "isNan";
- case SpvOpIsInf:
+ case spv::Op::OpIsInf:
return "isInf";
- case SpvOpTranspose:
+ case spv::Op::OpTranspose:
return "transpose";
default:
break;
@@ -194,63 +269,63 @@
// Converts a SPIR-V opcode to its corresponding AST binary opcode, if any
// @param opcode SPIR-V opcode
// @returns the AST binary op for the given opcode, or kNone
-ast::BinaryOp ConvertBinaryOp(SpvOp opcode) {
+ast::BinaryOp ConvertBinaryOp(spv::Op opcode) {
switch (opcode) {
- case SpvOpIAdd:
- case SpvOpFAdd:
+ case spv::Op::OpIAdd:
+ case spv::Op::OpFAdd:
return ast::BinaryOp::kAdd;
- case SpvOpISub:
- case SpvOpFSub:
+ case spv::Op::OpISub:
+ case spv::Op::OpFSub:
return ast::BinaryOp::kSubtract;
- case SpvOpIMul:
- case SpvOpFMul:
- case SpvOpVectorTimesScalar:
- case SpvOpMatrixTimesScalar:
- case SpvOpVectorTimesMatrix:
- case SpvOpMatrixTimesVector:
- case SpvOpMatrixTimesMatrix:
+ case spv::Op::OpIMul:
+ case spv::Op::OpFMul:
+ case spv::Op::OpVectorTimesScalar:
+ case spv::Op::OpMatrixTimesScalar:
+ case spv::Op::OpVectorTimesMatrix:
+ case spv::Op::OpMatrixTimesVector:
+ case spv::Op::OpMatrixTimesMatrix:
return ast::BinaryOp::kMultiply;
- case SpvOpUDiv:
- case SpvOpSDiv:
- case SpvOpFDiv:
+ case spv::Op::OpUDiv:
+ case spv::Op::OpSDiv:
+ case spv::Op::OpFDiv:
return ast::BinaryOp::kDivide;
- case SpvOpUMod:
- case SpvOpSMod:
- case SpvOpFRem:
+ case spv::Op::OpUMod:
+ case spv::Op::OpSMod:
+ case spv::Op::OpFRem:
return ast::BinaryOp::kModulo;
- case SpvOpLogicalEqual:
- case SpvOpIEqual:
- case SpvOpFOrdEqual:
+ case spv::Op::OpLogicalEqual:
+ case spv::Op::OpIEqual:
+ case spv::Op::OpFOrdEqual:
return ast::BinaryOp::kEqual;
- case SpvOpLogicalNotEqual:
- case SpvOpINotEqual:
- case SpvOpFOrdNotEqual:
+ case spv::Op::OpLogicalNotEqual:
+ case spv::Op::OpINotEqual:
+ case spv::Op::OpFOrdNotEqual:
return ast::BinaryOp::kNotEqual;
- case SpvOpBitwiseAnd:
+ case spv::Op::OpBitwiseAnd:
return ast::BinaryOp::kAnd;
- case SpvOpBitwiseOr:
+ case spv::Op::OpBitwiseOr:
return ast::BinaryOp::kOr;
- case SpvOpBitwiseXor:
+ case spv::Op::OpBitwiseXor:
return ast::BinaryOp::kXor;
- case SpvOpLogicalAnd:
+ case spv::Op::OpLogicalAnd:
return ast::BinaryOp::kAnd;
- case SpvOpLogicalOr:
+ case spv::Op::OpLogicalOr:
return ast::BinaryOp::kOr;
- case SpvOpUGreaterThan:
- case SpvOpSGreaterThan:
- case SpvOpFOrdGreaterThan:
+ case spv::Op::OpUGreaterThan:
+ case spv::Op::OpSGreaterThan:
+ case spv::Op::OpFOrdGreaterThan:
return ast::BinaryOp::kGreaterThan;
- case SpvOpUGreaterThanEqual:
- case SpvOpSGreaterThanEqual:
- case SpvOpFOrdGreaterThanEqual:
+ case spv::Op::OpUGreaterThanEqual:
+ case spv::Op::OpSGreaterThanEqual:
+ case spv::Op::OpFOrdGreaterThanEqual:
return ast::BinaryOp::kGreaterThanEqual;
- case SpvOpULessThan:
- case SpvOpSLessThan:
- case SpvOpFOrdLessThan:
+ case spv::Op::OpULessThan:
+ case spv::Op::OpSLessThan:
+ case spv::Op::OpFOrdLessThan:
return ast::BinaryOp::kLessThan;
- case SpvOpULessThanEqual:
- case SpvOpSLessThanEqual:
- case SpvOpFOrdLessThanEqual:
+ case spv::Op::OpULessThanEqual:
+ case spv::Op::OpSLessThanEqual:
+ case spv::Op::OpFOrdLessThanEqual:
return ast::BinaryOp::kLessThanEqual;
default:
break;
@@ -265,19 +340,19 @@
// Othewrise returns BinaryOp::kNone.
// @param opcode SPIR-V opcode
// @returns operation corresponding to negated version of the SPIR-V opcode
-ast::BinaryOp NegatedFloatCompare(SpvOp opcode) {
+ast::BinaryOp NegatedFloatCompare(spv::Op opcode) {
switch (opcode) {
- case SpvOpFUnordEqual:
+ case spv::Op::OpFUnordEqual:
return ast::BinaryOp::kNotEqual;
- case SpvOpFUnordNotEqual:
+ case spv::Op::OpFUnordNotEqual:
return ast::BinaryOp::kEqual;
- case SpvOpFUnordLessThan:
+ case spv::Op::OpFUnordLessThan:
return ast::BinaryOp::kGreaterThanEqual;
- case SpvOpFUnordLessThanEqual:
+ case spv::Op::OpFUnordLessThanEqual:
return ast::BinaryOp::kGreaterThan;
- case SpvOpFUnordGreaterThan:
+ case spv::Op::OpFUnordGreaterThan:
return ast::BinaryOp::kLessThanEqual;
- case SpvOpFUnordGreaterThanEqual:
+ case spv::Op::OpFUnordGreaterThanEqual:
return ast::BinaryOp::kLessThan;
default:
break;
@@ -443,36 +518,36 @@
// Returns the WGSL standard library function builtin for the
// given instruction, or sem::BuiltinType::kNone
-sem::BuiltinType GetBuiltin(SpvOp opcode) {
+sem::BuiltinType GetBuiltin(spv::Op opcode) {
switch (opcode) {
- case SpvOpBitCount:
+ case spv::Op::OpBitCount:
return sem::BuiltinType::kCountOneBits;
- case SpvOpBitFieldInsert:
+ case spv::Op::OpBitFieldInsert:
return sem::BuiltinType::kInsertBits;
- case SpvOpBitFieldSExtract:
- case SpvOpBitFieldUExtract:
+ case spv::Op::OpBitFieldSExtract:
+ case spv::Op::OpBitFieldUExtract:
return sem::BuiltinType::kExtractBits;
- case SpvOpBitReverse:
+ case spv::Op::OpBitReverse:
return sem::BuiltinType::kReverseBits;
- case SpvOpDot:
+ case spv::Op::OpDot:
return sem::BuiltinType::kDot;
- case SpvOpDPdx:
+ case spv::Op::OpDPdx:
return sem::BuiltinType::kDpdx;
- case SpvOpDPdy:
+ case spv::Op::OpDPdy:
return sem::BuiltinType::kDpdy;
- case SpvOpFwidth:
+ case spv::Op::OpFwidth:
return sem::BuiltinType::kFwidth;
- case SpvOpDPdxFine:
+ case spv::Op::OpDPdxFine:
return sem::BuiltinType::kDpdxFine;
- case SpvOpDPdyFine:
+ case spv::Op::OpDPdyFine:
return sem::BuiltinType::kDpdyFine;
- case SpvOpFwidthFine:
+ case spv::Op::OpFwidthFine:
return sem::BuiltinType::kFwidthFine;
- case SpvOpDPdxCoarse:
+ case spv::Op::OpDPdxCoarse:
return sem::BuiltinType::kDpdxCoarse;
- case SpvOpDPdyCoarse:
+ case spv::Op::OpDPdyCoarse:
return sem::BuiltinType::kDpdyCoarse;
- case SpvOpFwidthCoarse:
+ case spv::Op::OpFwidthCoarse:
return sem::BuiltinType::kFwidthCoarse;
default:
break;
@@ -483,20 +558,20 @@
// @param opcode a SPIR-V opcode
// @returns true if the given instruction is an image access instruction
// whose first input operand is an OpSampledImage value.
-bool IsSampledImageAccess(SpvOp opcode) {
+bool IsSampledImageAccess(spv::Op opcode) {
switch (opcode) {
- case SpvOpImageSampleImplicitLod:
- case SpvOpImageSampleExplicitLod:
- case SpvOpImageSampleDrefImplicitLod:
- case SpvOpImageSampleDrefExplicitLod:
+ case spv::Op::OpImageSampleImplicitLod:
+ case spv::Op::OpImageSampleExplicitLod:
+ case spv::Op::OpImageSampleDrefImplicitLod:
+ case spv::Op::OpImageSampleDrefExplicitLod:
// WGSL doesn't have *Proj* texturing; spirv reader emulates it.
- case SpvOpImageSampleProjImplicitLod:
- case SpvOpImageSampleProjExplicitLod:
- case SpvOpImageSampleProjDrefImplicitLod:
- case SpvOpImageSampleProjDrefExplicitLod:
- case SpvOpImageGather:
- case SpvOpImageDrefGather:
- case SpvOpImageQueryLod:
+ case spv::Op::OpImageSampleProjImplicitLod:
+ case spv::Op::OpImageSampleProjExplicitLod:
+ case spv::Op::OpImageSampleProjDrefImplicitLod:
+ case spv::Op::OpImageSampleProjDrefExplicitLod:
+ case spv::Op::OpImageGather:
+ case spv::Op::OpImageDrefGather:
+ case spv::Op::OpImageQueryLod:
return true;
default:
break;
@@ -506,29 +581,29 @@
// @param opcode a SPIR-V opcode
// @returns true if the given instruction is an atomic operation.
-bool IsAtomicOp(SpvOp opcode) {
+bool IsAtomicOp(spv::Op opcode) {
switch (opcode) {
- case SpvOpAtomicLoad:
- case SpvOpAtomicStore:
- case SpvOpAtomicExchange:
- case SpvOpAtomicCompareExchange:
- case SpvOpAtomicCompareExchangeWeak:
- case SpvOpAtomicIIncrement:
- case SpvOpAtomicIDecrement:
- case SpvOpAtomicIAdd:
- case SpvOpAtomicISub:
- case SpvOpAtomicSMin:
- case SpvOpAtomicUMin:
- case SpvOpAtomicSMax:
- case SpvOpAtomicUMax:
- case SpvOpAtomicAnd:
- case SpvOpAtomicOr:
- case SpvOpAtomicXor:
- case SpvOpAtomicFlagTestAndSet:
- case SpvOpAtomicFlagClear:
- case SpvOpAtomicFMinEXT:
- case SpvOpAtomicFMaxEXT:
- case SpvOpAtomicFAddEXT:
+ case spv::Op::OpAtomicLoad:
+ case spv::Op::OpAtomicStore:
+ case spv::Op::OpAtomicExchange:
+ case spv::Op::OpAtomicCompareExchange:
+ case spv::Op::OpAtomicCompareExchangeWeak:
+ case spv::Op::OpAtomicIIncrement:
+ case spv::Op::OpAtomicIDecrement:
+ case spv::Op::OpAtomicIAdd:
+ case spv::Op::OpAtomicISub:
+ case spv::Op::OpAtomicSMin:
+ case spv::Op::OpAtomicUMin:
+ case spv::Op::OpAtomicSMax:
+ case spv::Op::OpAtomicUMax:
+ case spv::Op::OpAtomicAnd:
+ case spv::Op::OpAtomicOr:
+ case spv::Op::OpAtomicXor:
+ case spv::Op::OpAtomicFlagTestAndSet:
+ case spv::Op::OpAtomicFlagClear:
+ case spv::Op::OpAtomicFMinEXT:
+ case spv::Op::OpAtomicFMaxEXT:
+ case spv::Op::OpAtomicFAddEXT:
return true;
default:
break;
@@ -539,19 +614,19 @@
// @param opcode a SPIR-V opcode
// @returns true if the given instruction is an image sampling, gather,
// or gather-compare operation.
-bool IsImageSamplingOrGatherOrDrefGather(SpvOp opcode) {
+bool IsImageSamplingOrGatherOrDrefGather(spv::Op opcode) {
switch (opcode) {
- case SpvOpImageSampleImplicitLod:
- case SpvOpImageSampleExplicitLod:
- case SpvOpImageSampleDrefImplicitLod:
- case SpvOpImageSampleDrefExplicitLod:
+ case spv::Op::OpImageSampleImplicitLod:
+ case spv::Op::OpImageSampleExplicitLod:
+ case spv::Op::OpImageSampleDrefImplicitLod:
+ case spv::Op::OpImageSampleDrefExplicitLod:
// WGSL doesn't have *Proj* texturing; spirv reader emulates it.
- case SpvOpImageSampleProjImplicitLod:
- case SpvOpImageSampleProjExplicitLod:
- case SpvOpImageSampleProjDrefImplicitLod:
- case SpvOpImageSampleProjDrefExplicitLod:
- case SpvOpImageGather:
- case SpvOpImageDrefGather:
+ case spv::Op::OpImageSampleProjImplicitLod:
+ case spv::Op::OpImageSampleProjExplicitLod:
+ case spv::Op::OpImageSampleProjDrefImplicitLod:
+ case spv::Op::OpImageSampleProjDrefExplicitLod:
+ case spv::Op::OpImageGather:
+ case spv::Op::OpImageDrefGather:
return true;
default:
break;
@@ -562,11 +637,11 @@
// @param opcode a SPIR-V opcode
// @returns true if the given instruction is an image access instruction
// whose first input operand is an OpImage value.
-bool IsRawImageAccess(SpvOp opcode) {
+bool IsRawImageAccess(spv::Op opcode) {
switch (opcode) {
- case SpvOpImageRead:
- case SpvOpImageWrite:
- case SpvOpImageFetch:
+ case spv::Op::OpImageRead:
+ case spv::Op::OpImageWrite:
+ case spv::Op::OpImageFetch:
return true;
default:
break;
@@ -576,13 +651,13 @@
// @param opcode a SPIR-V opcode
// @returns true if the given instruction is an image query instruction
-bool IsImageQuery(SpvOp opcode) {
+bool IsImageQuery(spv::Op opcode) {
switch (opcode) {
- case SpvOpImageQuerySize:
- case SpvOpImageQuerySizeLod:
- case SpvOpImageQueryLevels:
- case SpvOpImageQuerySamples:
- case SpvOpImageQueryLod:
+ case spv::Op::OpImageQuerySize:
+ case spv::Op::OpImageQuerySizeLod:
+ case spv::Op::OpImageQueryLevels:
+ case spv::Op::OpImageQuerySamples:
+ case spv::Op::OpImageQueryLod:
return true;
default:
break;
@@ -653,15 +728,15 @@
// Visit successors. We will naturally skip the continue target and merge
// blocks.
auto* terminator = bb->terminator();
- auto opcode = terminator->opcode();
- if (opcode == SpvOpBranchConditional) {
+ auto opcode = ToOpcode(terminator->opcode());
+ if (opcode == spv::Op::OpBranchConditional) {
// Visit the false branch, then the true branch, to make them come
// out in the natural order for an "if".
VisitBackward(terminator->GetSingleWordInOperand(2));
VisitBackward(terminator->GetSingleWordInOperand(1));
- } else if (opcode == SpvOpBranch) {
+ } else if (opcode == spv::Op::OpBranch) {
VisitBackward(terminator->GetSingleWordInOperand(0));
- } else if (opcode == SpvOpSwitch) {
+ } else if (opcode == spv::Op::OpSwitch) {
// TODO(dneto): Consider visiting the labels in literal-value order.
utils::Vector<uint32_t, 32> successors;
bb->ForEachSuccessorLabel(
@@ -1265,7 +1340,7 @@
for (uint32_t var_id : ep_info_->inputs) {
const auto* var = def_use_mgr_->GetDef(var_id);
TINT_ASSERT(Reader, var != nullptr);
- TINT_ASSERT(Reader, var->opcode() == SpvOpVariable);
+ TINT_ASSERT(Reader, opcode(var) == spv::Op::OpVariable);
auto* store_type = GetVariableStoreType(*var);
auto* forced_param_type = store_type;
AttributeList param_decos;
@@ -1344,7 +1419,7 @@
} else {
const auto* var = def_use_mgr_->GetDef(var_id);
TINT_ASSERT(Reader, var != nullptr);
- TINT_ASSERT(Reader, var->opcode() == SpvOpVariable);
+ TINT_ASSERT(Reader, opcode(var) == spv::Op::OpVariable);
const Type* store_type = GetVariableStoreType(*var);
const Type* forced_member_type = store_type;
AttributeList out_decos;
@@ -1570,19 +1645,19 @@
}
if (const auto* inst = block.GetMergeInst()) {
- auto terminator_opcode = block.terminator()->opcode();
- switch (inst->opcode()) {
- case SpvOpSelectionMerge:
- if ((terminator_opcode != SpvOpBranchConditional) &&
- (terminator_opcode != SpvOpSwitch)) {
+ auto terminator_opcode = opcode(block.terminator());
+ switch (opcode(inst)) {
+ case spv::Op::OpSelectionMerge:
+ if ((terminator_opcode != spv::Op::OpBranchConditional) &&
+ (terminator_opcode != spv::Op::OpSwitch)) {
return Fail() << "Selection header " << block_id
<< " does not end in an OpBranchConditional or "
"OpSwitch instruction";
}
break;
- case SpvOpLoopMerge:
- if ((terminator_opcode != SpvOpBranchConditional) &&
- (terminator_opcode != SpvOpBranch)) {
+ case spv::Op::OpLoopMerge:
+ if ((terminator_opcode != spv::Op::OpBranchConditional) &&
+ (terminator_opcode != spv::Op::OpBranch)) {
return Fail() << "Loop header " << block_id
<< " does not end in an OpBranch or "
"OpBranchConditional instruction";
@@ -1612,7 +1687,7 @@
merge_info->header_for_merge = header;
header_info->merge_for_header = merge;
- if (inst->opcode() == SpvOpLoopMerge) {
+ if (opcode(inst) == spv::Op::OpLoopMerge) {
if (header == entry_id) {
return Fail() << "Function entry block " << entry_id
<< " cannot be a loop header";
@@ -1867,8 +1942,8 @@
// From the interval rule, the selection construct consists of blocks
// in the block order, starting at the header, until just before the
// merge block.
- const auto branch_opcode = header_info->basic_block->terminator()->opcode();
- const auto kind = (branch_opcode == SpvOpBranchConditional)
+ const auto branch_opcode = opcode(header_info->basic_block->terminator());
+ const auto kind = (branch_opcode == spv::Op::OpBranchConditional)
? Construct::kIfSelection
: Construct::kSwitchSelection;
top = push_construct(depth, kind, header, merge);
@@ -2485,7 +2560,7 @@
return false;
}
for (auto& inst : *function_.entry()) {
- if (inst.opcode() != SpvOpVariable) {
+ if (opcode(inst) != spv::Op::OpVariable) {
continue;
}
auto* var_store_type = GetVariableStoreType(inst);
@@ -2586,8 +2661,8 @@
Fail() << "ID " << id << " does not have a defining SPIR-V instruction";
return {};
}
- switch (inst->opcode()) {
- case SpvOpVariable: {
+ switch (opcode(inst)) {
+ case spv::Op::OpVariable: {
// This occurs for module-scope variables.
auto name = namer_.Name(id);
// Construct the reference type, mapping storage class correctly.
@@ -2596,7 +2671,7 @@
return TypedExpression{type, create<ast::IdentifierExpression>(
Source{}, builder_.Symbols().Register(name))};
}
- case SpvOpUndef:
+ case spv::Op::OpUndef:
// Substitute a null value for undef.
// This case occurs when OpUndef appears at module scope, as if it were
// a constant.
@@ -3102,11 +3177,11 @@
bool FunctionEmitter::EmitNormalTerminator(const BlockInfo& block_info) {
const auto& terminator = *(block_info.basic_block->terminator());
- switch (terminator.opcode()) {
- case SpvOpReturn:
+ switch (opcode(terminator)) {
+ case spv::Op::OpReturn:
AddStatement(create<ast::ReturnStatement>(Source{}));
return true;
- case SpvOpReturnValue: {
+ case spv::Op::OpReturnValue: {
auto value = MakeExpression(terminator.GetSingleWordInOperand(0));
if (!value) {
return false;
@@ -3114,12 +3189,12 @@
AddStatement(create<ast::ReturnStatement>(Source{}, value.expr));
}
return true;
- case SpvOpKill:
+ case spv::Op::OpKill:
// For now, assume SPIR-V OpKill has same semantics as WGSL discard.
// TODO(dneto): https://github.com/gpuweb/gpuweb/issues/676
AddStatement(create<ast::DiscardStatement>(Source{}));
return true;
- case SpvOpUnreachable:
+ case spv::Op::OpUnreachable:
// Translate as if it's a return. This avoids the problem where WGSL
// requires a return statement at the end of the function body.
{
@@ -3133,12 +3208,12 @@
}
}
return true;
- case SpvOpBranch: {
+ case spv::Op::OpBranch: {
const auto dest_id = terminator.GetSingleWordInOperand(0);
AddStatement(MakeBranch(block_info, *GetBlockInfo(dest_id)));
return true;
}
- case SpvOpBranchConditional: {
+ case spv::Op::OpBranchConditional: {
// If both destinations are the same, then do the same as we would
// for an unconditional branch (OpBranch).
const auto true_dest = terminator.GetSingleWordInOperand(1);
@@ -3191,7 +3266,7 @@
}
return true;
}
- case SpvOpSwitch:
+ case spv::Op::OpSwitch:
// An OpSelectionMerge must precede an OpSwitch. That is clarified
// in the resolution to Khronos-internal SPIR-V issue 115.
// A new enough version of the SPIR-V validator checks this case.
@@ -3377,8 +3452,8 @@
const auto* terminator = bb.terminator();
const auto* merge = bb.GetMergeInst(); // Might be nullptr
for (auto& inst : bb) {
- if (&inst == terminator || &inst == merge || inst.opcode() == SpvOpLabel ||
- inst.opcode() == SpvOpVariable) {
+ if (&inst == terminator || &inst == merge || opcode(inst) == spv::Op::OpLabel ||
+ opcode(inst) == spv::Op::OpVariable) {
continue;
}
if (!EmitStatement(inst)) {
@@ -3546,23 +3621,23 @@
return false;
}
- if (IsImageQuery(inst.opcode())) {
+ if (IsImageQuery(opcode(inst))) {
return EmitImageQuery(inst);
}
- if (IsSampledImageAccess(inst.opcode()) || IsRawImageAccess(inst.opcode())) {
+ if (IsSampledImageAccess(opcode(inst)) || IsRawImageAccess(opcode(inst))) {
return EmitImageAccess(inst);
}
- if (IsAtomicOp(inst.opcode())) {
+ if (IsAtomicOp(opcode(inst))) {
return EmitAtomicOp(inst);
}
- switch (inst.opcode()) {
- case SpvOpNop:
+ switch (opcode(inst)) {
+ case spv::Op::OpNop:
return true;
- case SpvOpStore: {
+ case spv::Op::OpStore: {
auto ptr_id = inst.GetSingleWordInOperand(0);
const auto value_id = inst.GetSingleWordInOperand(1);
@@ -3633,7 +3708,7 @@
return success();
}
- case SpvOpLoad: {
+ case spv::Op::OpLoad: {
// Memory accesses must be issued in SPIR-V program order.
// So represent a load by a new const definition.
const auto ptr_id = inst.GetSingleWordInOperand(0);
@@ -3684,7 +3759,7 @@
return EmitConstDefOrWriteToHoistedVar(inst, expr);
}
- case SpvOpCopyMemory: {
+ case spv::Op::OpCopyMemory: {
// Generate an assignment.
auto lhs = MakeOperand(inst, 0);
auto rhs = MakeOperand(inst, 1);
@@ -3713,7 +3788,7 @@
return success();
}
- case SpvOpCopyObject: {
+ case spv::Op::OpCopyObject: {
// Arguably, OpCopyObject is purely combinatorial. On the other hand,
// it exists to make a new name for something. So we choose to make
// a new named constant definition.
@@ -3732,39 +3807,39 @@
return EmitConstDefOrWriteToHoistedVar(inst, expr);
}
- case SpvOpPhi: {
+ case spv::Op::OpPhi: {
// The value will be in scope, available for reading from the phi ID.
return true;
}
- case SpvOpOuterProduct:
+ case spv::Op::OpOuterProduct:
// Synthesize an outer product expression in its own statement.
return EmitConstDefOrWriteToHoistedVar(inst, MakeOuterProduct(inst));
- case SpvOpVectorInsertDynamic:
+ case spv::Op::OpVectorInsertDynamic:
// Synthesize a vector insertion in its own statements.
return MakeVectorInsertDynamic(inst);
- case SpvOpCompositeInsert:
+ case spv::Op::OpCompositeInsert:
// Synthesize a composite insertion in its own statements.
return MakeCompositeInsert(inst);
- case SpvOpFunctionCall:
+ case spv::Op::OpFunctionCall:
return EmitFunctionCall(inst);
- case SpvOpControlBarrier:
+ case spv::Op::OpControlBarrier:
return EmitControlBarrier(inst);
- case SpvOpExtInst:
+ case spv::Op::OpExtInst:
if (parser_impl_.IsIgnoredExtendedInstruction(inst)) {
return true;
}
break;
- case SpvOpIAddCarry:
- case SpvOpISubBorrow:
- case SpvOpUMulExtended:
- case SpvOpSMulExtended:
+ case spv::Op::OpIAddCarry:
+ case spv::Op::OpISubBorrow:
+ case spv::Op::OpUMulExtended:
+ case spv::Op::OpSMulExtended:
return Fail() << "extended arithmetic is not finalized for WGSL: "
"https://github.com/gpuweb/gpuweb/issues/1565: "
<< inst.PrettyPrint();
@@ -3772,7 +3847,7 @@
default:
break;
}
- return Fail() << "unhandled instruction with opcode " << inst.opcode() << ": "
+ return Fail() << "unhandled instruction with opcode " << uint32_t(opcode(inst)) << ": "
<< inst.PrettyPrint();
}
@@ -3791,7 +3866,7 @@
return {};
}
- const auto opcode = inst.opcode();
+ const auto op = opcode(inst);
const Type* ast_type = nullptr;
if (inst.type_id()) {
@@ -3802,7 +3877,7 @@
}
}
- auto binary_op = ConvertBinaryOp(opcode);
+ auto binary_op = ConvertBinaryOp(op);
if (binary_op != ast::BinaryOp::kNone) {
auto arg0 = MakeOperand(inst, 0);
auto arg1 =
@@ -3817,14 +3892,14 @@
}
auto unary_op = ast::UnaryOp::kNegation;
- if (GetUnaryOp(opcode, &unary_op)) {
+ if (GetUnaryOp(op, &unary_op)) {
auto arg0 = MakeOperand(inst, 0);
auto* unary_expr = create<ast::UnaryOpExpression>(Source{}, unary_op, arg0.expr);
TypedExpression result{ast_type, unary_expr};
return parser_impl_.RectifyForcedResultType(result, inst, arg0.type);
}
- const char* unary_builtin_name = GetUnaryBuiltInFunctionName(opcode);
+ const char* unary_builtin_name = GetUnaryBuiltInFunctionName(op);
if (unary_builtin_name != nullptr) {
ExpressionList params;
params.Push(MakeOperand(inst, 0).expr);
@@ -3835,40 +3910,40 @@
std::move(params))};
}
- const auto builtin = GetBuiltin(opcode);
+ const auto builtin = GetBuiltin(op);
if (builtin != sem::BuiltinType::kNone) {
return MakeBuiltinCall(inst);
}
- if (opcode == SpvOpFMod) {
+ if (op == spv::Op::OpFMod) {
return MakeFMod(inst);
}
- if (opcode == SpvOpAccessChain || opcode == SpvOpInBoundsAccessChain) {
+ if (op == spv::Op::OpAccessChain || op == spv::Op::OpInBoundsAccessChain) {
return MakeAccessChain(inst);
}
- if (opcode == SpvOpBitcast) {
+ if (op == spv::Op::OpBitcast) {
return {ast_type, create<ast::BitcastExpression>(Source{}, ast_type->Build(builder_),
MakeOperand(inst, 0).expr)};
}
- if (opcode == SpvOpShiftLeftLogical || opcode == SpvOpShiftRightLogical ||
- opcode == SpvOpShiftRightArithmetic) {
+ if (op == spv::Op::OpShiftLeftLogical || op == spv::Op::OpShiftRightLogical ||
+ op == spv::Op::OpShiftRightArithmetic) {
auto arg0 = MakeOperand(inst, 0);
// The second operand must be unsigned. It's ok to wrap the shift amount
// since the shift is modulo the bit width of the first operand.
auto arg1 = parser_impl_.AsUnsigned(MakeOperand(inst, 1));
- switch (opcode) {
- case SpvOpShiftLeftLogical:
+ switch (op) {
+ case spv::Op::OpShiftLeftLogical:
binary_op = ast::BinaryOp::kShiftLeft;
break;
- case SpvOpShiftRightLogical:
+ case spv::Op::OpShiftRightLogical:
arg0 = parser_impl_.AsUnsigned(arg0);
binary_op = ast::BinaryOp::kShiftRight;
break;
- case SpvOpShiftRightArithmetic:
+ case spv::Op::OpShiftRightArithmetic:
arg0 = parser_impl_.AsSigned(arg0);
binary_op = ast::BinaryOp::kShiftRight;
break;
@@ -3880,7 +3955,7 @@
return parser_impl_.RectifyForcedResultType(result, inst, arg0.type);
}
- auto negated_op = NegatedFloatCompare(opcode);
+ auto negated_op = NegatedFloatCompare(op);
if (negated_op != ast::BinaryOp::kNone) {
auto arg0 = MakeOperand(inst, 0);
auto arg1 = MakeOperand(inst, 1);
@@ -3891,7 +3966,7 @@
return {ast_type, negated_expr};
}
- if (opcode == SpvOpExtInst) {
+ if (op == spv::Op::OpExtInst) {
if (parser_impl_.IsIgnoredExtendedInstruction(inst)) {
// Ignore it but don't error out.
return {};
@@ -3904,7 +3979,7 @@
return EmitGlslStd450ExtInst(inst);
}
- if (opcode == SpvOpCompositeConstruct) {
+ if (op == spv::Op::OpCompositeConstruct) {
ExpressionList operands;
for (uint32_t iarg = 0; iarg < inst.NumInOperands(); ++iarg) {
operands.Push(MakeOperand(inst, iarg).expr);
@@ -3913,34 +3988,34 @@
builder_.Construct(Source{}, ast_type->Build(builder_), std::move(operands))};
}
- if (opcode == SpvOpCompositeExtract) {
+ if (op == spv::Op::OpCompositeExtract) {
return MakeCompositeExtract(inst);
}
- if (opcode == SpvOpVectorShuffle) {
+ if (op == spv::Op::OpVectorShuffle) {
return MakeVectorShuffle(inst);
}
- if (opcode == SpvOpVectorExtractDynamic) {
+ if (op == spv::Op::OpVectorExtractDynamic) {
return {ast_type, create<ast::IndexAccessorExpression>(Source{}, MakeOperand(inst, 0).expr,
MakeOperand(inst, 1).expr)};
}
- if (opcode == SpvOpConvertSToF || opcode == SpvOpConvertUToF || opcode == SpvOpConvertFToS ||
- opcode == SpvOpConvertFToU) {
+ if (op == spv::Op::OpConvertSToF || op == spv::Op::OpConvertUToF ||
+ op == spv::Op::OpConvertFToS || op == spv::Op::OpConvertFToU) {
return MakeNumericConversion(inst);
}
- if (opcode == SpvOpUndef) {
+ if (op == spv::Op::OpUndef) {
// Replace undef with the null value.
return parser_impl_.MakeNullExpression(ast_type);
}
- if (opcode == SpvOpSelect) {
+ if (op == spv::Op::OpSelect) {
return MakeSimpleSelect(inst);
}
- if (opcode == SpvOpArrayLength) {
+ if (op == spv::Op::OpArrayLength) {
return MakeArrayLength(inst);
}
@@ -4424,12 +4499,12 @@
const auto constants = constant_mgr_->GetOperandConstants(&inst);
const auto* ptr_type_inst = def_use_mgr_->GetDef(ptr_ty_id);
- if (!ptr_type_inst || (ptr_type_inst->opcode() != SpvOpTypePointer)) {
+ if (!ptr_type_inst || (opcode(ptr_type_inst) != spv::Op::OpTypePointer)) {
Fail() << "Access chain %" << inst.result_id() << " base pointer is not of pointer type";
return {};
}
- SpvStorageClass address_space =
- static_cast<SpvStorageClass>(ptr_type_inst->GetSingleWordInOperand(0));
+ spv::StorageClass address_space =
+ static_cast<spv::StorageClass>(ptr_type_inst->GetSingleWordInOperand(0));
uint32_t pointee_type_id = ptr_type_inst->GetSingleWordInOperand(1);
// Build up a nested expression for the access chain by walking down the type
@@ -4446,8 +4521,8 @@
<< (index - first_index) << " indices: " << inst.PrettyPrint();
return {};
}
- switch (pointee_type_inst->opcode()) {
- case SpvOpTypeVector:
+ switch (opcode(pointee_type_inst)) {
+ case spv::Op::OpTypeVector:
if (index_const) {
// Try generating a MemberAccessor expression
const auto num_elems = pointee_type_inst->GetSingleWordInOperand(1);
@@ -4473,24 +4548,24 @@
// Sink pointers to vector components.
sink_pointer = true;
break;
- case SpvOpTypeMatrix:
+ case spv::Op::OpTypeMatrix:
// Use array syntax.
next_expr = create<ast::IndexAccessorExpression>(Source{}, current_expr.expr,
MakeOperand(inst, index).expr);
// All matrix components are the same type.
pointee_type_id = pointee_type_inst->GetSingleWordInOperand(0);
break;
- case SpvOpTypeArray:
+ case spv::Op::OpTypeArray:
next_expr = create<ast::IndexAccessorExpression>(Source{}, current_expr.expr,
MakeOperand(inst, index).expr);
pointee_type_id = pointee_type_inst->GetSingleWordInOperand(0);
break;
- case SpvOpTypeRuntimeArray:
+ case spv::Op::OpTypeRuntimeArray:
next_expr = create<ast::IndexAccessorExpression>(Source{}, current_expr.expr,
MakeOperand(inst, index).expr);
pointee_type_id = pointee_type_inst->GetSingleWordInOperand(0);
break;
- case SpvOpTypeStruct: {
+ case spv::Op::OpTypeStruct: {
if (!index_const) {
Fail() << "Access chain %" << inst.result_id() << " index %"
<< inst.GetSingleWordInOperand(index)
@@ -4519,7 +4594,9 @@
<< ": " << pointee_type_inst->PrettyPrint();
return {};
}
- const auto pointer_type_id = type_mgr_->FindPointerToType(pointee_type_id, address_space);
+ const auto pointer_type_id = three_sided_patch_function_cc::FindPointerToType(
+ type_mgr_, &spvtools::opt::analysis::TypeManager::FindPointerToType, pointee_type_id,
+ address_space);
auto* type = parser_impl_.ConvertType(pointer_type_id, PtrAs::Ref);
TINT_ASSERT(Reader, type && type->Is<Reference>());
current_expr = TypedExpression{type, next_expr};
@@ -4594,11 +4671,11 @@
return {};
}
const char* operation_name = nullptr;
- switch (inst.opcode()) {
- case SpvOpCompositeExtract:
+ switch (opcode(inst)) {
+ case spv::Op::OpCompositeExtract:
operation_name = "OpCompositeExtract";
break;
- case SpvOpCompositeInsert:
+ case spv::Op::OpCompositeInsert:
operation_name = "OpCompositeInsert";
break;
default:
@@ -4606,8 +4683,8 @@
return {};
}
const ast::Expression* next_expr = nullptr;
- switch (current_type_inst->opcode()) {
- case SpvOpTypeVector: {
+ switch (opcode(current_type_inst)) {
+ case spv::Op::OpTypeVector: {
// Try generating a MemberAccessor expression. That result in something
// like "foo.z", which is more idiomatic than "foo[2]".
const auto num_elems = current_type_inst->GetSingleWordInOperand(1);
@@ -4628,7 +4705,7 @@
current_type_id = current_type_inst->GetSingleWordInOperand(0);
break;
}
- case SpvOpTypeMatrix: {
+ case spv::Op::OpTypeMatrix: {
// Check bounds
const auto num_elems = current_type_inst->GetSingleWordInOperand(1);
if (num_elems <= index_val) {
@@ -4648,7 +4725,7 @@
current_type_id = current_type_inst->GetSingleWordInOperand(0);
break;
}
- case SpvOpTypeArray:
+ case spv::Op::OpTypeArray:
// The array size could be a spec constant, and so it's not always
// statically checkable. Instead, rely on a runtime index clamp
// or runtime check to keep this safe.
@@ -4656,11 +4733,11 @@
make_index(index_val));
current_type_id = current_type_inst->GetSingleWordInOperand(0);
break;
- case SpvOpTypeRuntimeArray:
+ case spv::Op::OpTypeRuntimeArray:
Fail() << "can't do " << operation_name
<< " on a runtime array: " << inst.PrettyPrint();
return {};
- case SpvOpTypeStruct: {
+ case spv::Op::OpTypeStruct: {
const auto num_members = current_type_inst->NumInOperands();
if (num_members <= index_val) {
Fail() << operation_name << " %" << inst.result_id() << " index value "
@@ -4751,14 +4828,14 @@
auto& def = def_info_[id];
// Builtins are always defined outside the function.
switch (builtin) {
- case SpvBuiltInPointSize:
+ case spv::BuiltIn::PointSize:
def->skip = SkipReason::kPointSizeBuiltinPointer;
break;
- case SpvBuiltInSampleMask: {
+ case spv::BuiltIn::SampleMask: {
// Distinguish between input and output variable.
const auto storage_class =
- static_cast<SpvStorageClass>(var->GetSingleWordInOperand(0));
- if (storage_class == SpvStorageClassInput) {
+ static_cast<spv::StorageClass>(var->GetSingleWordInOperand(0));
+ if (storage_class == spv::StorageClass::Input) {
sample_mask_in_id = id;
def->skip = SkipReason::kSampleMaskInBuiltinPointer;
} else {
@@ -4767,14 +4844,14 @@
}
break;
}
- case SpvBuiltInSampleId:
- case SpvBuiltInInstanceIndex:
- case SpvBuiltInVertexIndex:
- case SpvBuiltInLocalInvocationIndex:
- case SpvBuiltInLocalInvocationId:
- case SpvBuiltInGlobalInvocationId:
- case SpvBuiltInWorkgroupId:
- case SpvBuiltInNumWorkgroups:
+ case spv::BuiltIn::SampleId:
+ case spv::BuiltIn::InstanceIndex:
+ case spv::BuiltIn::VertexIndex:
+ case spv::BuiltIn::LocalInvocationIndex:
+ case spv::BuiltIn::LocalInvocationId:
+ case spv::BuiltIn::GlobalInvocationId:
+ case spv::BuiltIn::WorkgroupId:
+ case spv::BuiltIn::NumWorkgroups:
break;
default:
return Fail() << "unrecognized special builtin: " << int(builtin);
@@ -4791,7 +4868,7 @@
const auto block_pos = block_info->pos;
for (const auto& inst : *(block_info->basic_block)) {
const auto result_id = inst.result_id();
- if ((result_id == 0) || inst.opcode() == SpvOpLabel) {
+ if ((result_id == 0) || opcode(inst) == spv::Op::OpLabel) {
continue;
}
def_info_[result_id] = std::make_unique<DefInfo>(index, inst, block_pos);
@@ -4807,15 +4884,15 @@
// pointer definitions are SSA values, and their definitions must be
// visited before their uses.
if (type->AsPointer()) {
- switch (inst.opcode()) {
- case SpvOpUndef:
+ switch (opcode(inst)) {
+ case spv::Op::OpUndef:
return Fail() << "undef pointer is not valid: " << inst.PrettyPrint();
- case SpvOpVariable:
+ case spv::Op::OpVariable:
info->pointer = GetPointerInfo(result_id);
break;
- case SpvOpAccessChain:
- case SpvOpInBoundsAccessChain:
- case SpvOpCopyObject:
+ case spv::Op::OpAccessChain:
+ case spv::Op::OpInBoundsAccessChain:
+ case spv::Op::OpCopyObject:
// Inherit from the first operand. We need this so we can pick up
// a remapped storage buffer.
info->pointer = GetPointerInfo(inst.GetSingleWordInOperand(0));
@@ -4846,8 +4923,8 @@
[&](const spvtools::opt::Instruction& inst) -> DefInfo::Pointer {
// WGSL root identifiers (or SPIR-V "memory object declarations") are
// either variables or function parameters.
- switch (inst.opcode()) {
- case SpvOpVariable: {
+ switch (opcode(inst)) {
+ case spv::Op::OpVariable: {
if (const auto* module_var = parser_impl_.GetModuleVariable(id)) {
return DefInfo::Pointer{module_var->declared_address_space,
module_var->declared_access};
@@ -4856,7 +4933,7 @@
// access mode.
return DefInfo::Pointer{ast::AddressSpace::kFunction, ast::Access::kUndefined};
}
- case SpvOpFunctionParameter: {
+ case spv::Op::OpFunctionParameter: {
const auto* type = As<Pointer>(parser_impl_.ConvertType(inst.type_id()));
// For access mode, kUndefined is ok for now, since the
// only non-default access mode on a pointer would be for a storage
@@ -4879,7 +4956,7 @@
auto where = def_info_.find(id);
if (where != def_info_.end()) {
const auto& info = where->second;
- if (info->inst.opcode() == SpvOpVariable) {
+ if (opcode(info->inst) == spv::Op::OpVariable) {
// Ignore the cache in this case and compute it from scratch.
// That's because for a function-scope OpVariable is a
// locally-defined value. So its cache entry has been created
@@ -4925,8 +5002,8 @@
};
for (auto& id_def_info_pair : def_info_) {
const auto& inst = id_def_info_pair.second->inst;
- const auto opcode = inst.opcode();
- if ((opcode == SpvOpVectorShuffle) || (opcode == SpvOpOuterProduct)) {
+ const auto op = opcode(inst);
+ if ((op == spv::Op::OpVectorShuffle) || (op == spv::Op::OpOuterProduct)) {
// We might access the vector operands multiple times. Make sure they
// are evaluated only once.
require_named_const_def(inst, 0);
@@ -4979,7 +5056,7 @@
const auto* block_info = GetBlockInfo(block_id);
for (const auto& inst : *(block_info->basic_block)) {
// Update bookkeeping for locally-defined IDs used by this instruction.
- if (inst.opcode() == SpvOpPhi) {
+ if (opcode(inst) == spv::Op::OpPhi) {
// For an OpPhi defining value P, an incoming value V from parent block B is
// counted as being "used" at block B, not at the block containing the Phi.
// That's because we will create a variable PHI_P to hold the phi value, and
@@ -5128,7 +5205,7 @@
}
TypedExpression FunctionEmitter::MakeNumericConversion(const spvtools::opt::Instruction& inst) {
- const auto opcode = inst.opcode();
+ const auto op = opcode(inst);
auto* requested_type = parser_impl_.ConvertType(inst.type_id());
auto arg_expr = MakeOperand(inst, 0);
if (!arg_expr) {
@@ -5137,7 +5214,7 @@
arg_expr.type = arg_expr.type->UnwrapRef();
const Type* expr_type = nullptr;
- if ((opcode == SpvOpConvertSToF) || (opcode == SpvOpConvertUToF)) {
+ if ((op == spv::Op::OpConvertSToF) || (op == spv::Op::OpConvertUToF)) {
if (arg_expr.type->IsIntegerScalarOrVector()) {
expr_type = requested_type;
} else {
@@ -5145,7 +5222,7 @@
"scalar or vector: "
<< inst.PrettyPrint();
}
- } else if (inst.opcode() == SpvOpConvertFToU) {
+ } else if (op == spv::Op::OpConvertFToU) {
if (arg_expr.type->IsFloatScalarOrVector()) {
expr_type = parser_impl_.GetUnsignedIntMatchingShape(arg_expr.type);
} else {
@@ -5153,7 +5230,7 @@
"point scalar or vector: "
<< inst.PrettyPrint();
}
- } else if (inst.opcode() == SpvOpConvertFToS) {
+ } else if (op == spv::Op::OpConvertFToS) {
if (arg_expr.type->IsFloatScalarOrVector()) {
expr_type = parser_impl_.GetSignedIntMatchingShape(arg_expr.type);
} else {
@@ -5228,28 +5305,28 @@
uint32_t memory = operands[1];
uint32_t semantics = operands[2];
- if (execution != SpvScopeWorkgroup) {
+ if (execution != uint32_t(spv::Scope::Workgroup)) {
return Fail() << "unsupported control barrier execution scope: "
<< "expected Workgroup (2), got: " << execution;
}
- if (semantics & SpvMemorySemanticsAcquireReleaseMask) {
- semantics &= ~static_cast<uint32_t>(SpvMemorySemanticsAcquireReleaseMask);
+ if (semantics & uint32_t(spv::MemorySemanticsMask::AcquireRelease)) {
+ semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::AcquireRelease);
} else {
return Fail() << "control barrier semantics requires acquire and release";
}
- if (semantics & SpvMemorySemanticsWorkgroupMemoryMask) {
- if (memory != SpvScopeWorkgroup) {
+ if (semantics & uint32_t(spv::MemorySemanticsMask::WorkgroupMemory)) {
+ if (memory != uint32_t(spv::Scope::Workgroup)) {
return Fail() << "workgroupBarrier requires workgroup memory scope";
}
AddStatement(create<ast::CallStatement>(builder_.Call("workgroupBarrier")));
- semantics &= ~static_cast<uint32_t>(SpvMemorySemanticsWorkgroupMemoryMask);
+ semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::WorkgroupMemory);
}
- if (semantics & SpvMemorySemanticsUniformMemoryMask) {
- if (memory != SpvScopeDevice) {
+ if (semantics & uint32_t(spv::MemorySemanticsMask::UniformMemory)) {
+ if (memory != uint32_t(spv::Scope::Device)) {
return Fail() << "storageBarrier requires device memory scope";
}
AddStatement(create<ast::CallStatement>(builder_.Call("storageBarrier")));
- semantics &= ~static_cast<uint32_t>(SpvMemorySemanticsUniformMemoryMask);
+ semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::UniformMemory);
}
if (semantics) {
return Fail() << "unsupported control barrier semantics: " << semantics;
@@ -5258,7 +5335,7 @@
}
TypedExpression FunctionEmitter::MakeBuiltinCall(const spvtools::opt::Instruction& inst) {
- const auto builtin = GetBuiltin(inst.opcode());
+ const auto builtin = GetBuiltin(opcode(inst));
auto* name = sem::str(builtin);
auto* ident = create<ast::IdentifierExpression>(Source{}, builder_.Symbols().Register(name));
@@ -5374,7 +5451,7 @@
bool FunctionEmitter::EmitImageAccess(const spvtools::opt::Instruction& inst) {
ExpressionList args;
- const auto opcode = inst.opcode();
+ const auto op = opcode(inst);
// Form the texture operand.
const spvtools::opt::Instruction* image = GetImage(inst);
@@ -5384,7 +5461,7 @@
args.Push(GetImageExpression(inst));
// Form the sampler operand, if needed.
- if (IsSampledImageAccess(opcode)) {
+ if (IsSampledImageAccess(op)) {
// Form the sampler operand.
if (auto* sampler = GetSamplerExpression(inst)) {
args.Push(sampler);
@@ -5438,25 +5515,25 @@
bool is_dref_sample = false;
bool is_gather_or_dref_gather = false;
bool is_non_dref_sample = false;
- switch (opcode) {
- case SpvOpImageSampleImplicitLod:
- case SpvOpImageSampleExplicitLod:
- case SpvOpImageSampleProjImplicitLod:
- case SpvOpImageSampleProjExplicitLod:
+ switch (op) {
+ case spv::Op::OpImageSampleImplicitLod:
+ case spv::Op::OpImageSampleExplicitLod:
+ case spv::Op::OpImageSampleProjImplicitLod:
+ case spv::Op::OpImageSampleProjExplicitLod:
is_non_dref_sample = true;
builtin_name = "textureSample";
break;
- case SpvOpImageSampleDrefImplicitLod:
- case SpvOpImageSampleDrefExplicitLod:
- case SpvOpImageSampleProjDrefImplicitLod:
- case SpvOpImageSampleProjDrefExplicitLod:
+ case spv::Op::OpImageSampleDrefImplicitLod:
+ case spv::Op::OpImageSampleDrefExplicitLod:
+ case spv::Op::OpImageSampleProjDrefImplicitLod:
+ case spv::Op::OpImageSampleProjDrefExplicitLod:
is_dref_sample = true;
builtin_name = "textureSampleCompare";
if (!consume_dref()) {
return false;
}
break;
- case SpvOpImageGather:
+ case spv::Op::OpImageGather:
is_gather_or_dref_gather = true;
builtin_name = "textureGather";
if (!texture_type->Is<DepthTexture>()) {
@@ -5471,20 +5548,20 @@
// Skip over the component operand, even for depth textures.
arg_index++;
break;
- case SpvOpImageDrefGather:
+ case spv::Op::OpImageDrefGather:
is_gather_or_dref_gather = true;
builtin_name = "textureGatherCompare";
if (!consume_dref()) {
return false;
}
break;
- case SpvOpImageFetch:
- case SpvOpImageRead:
+ case spv::Op::OpImageFetch:
+ case spv::Op::OpImageRead:
// Read a single texel from a sampled or storage image.
builtin_name = "textureLoad";
use_level_of_detail_suffix = false;
break;
- case SpvOpImageWrite:
+ case spv::Op::OpImageWrite:
builtin_name = "textureStore";
use_level_of_detail_suffix = false;
if (arg_index < num_args) {
@@ -5511,7 +5588,7 @@
image_operands_mask = inst.GetSingleWordInOperand(arg_index);
arg_index++;
}
- if (arg_index < num_args && (image_operands_mask & SpvImageOperandsBiasMask)) {
+ if (arg_index < num_args && (image_operands_mask & uint32_t(spv::ImageOperandsMask::Bias))) {
if (is_dref_sample) {
return Fail() << "WGSL does not support depth-reference sampling with "
"level-of-detail bias: "
@@ -5524,10 +5601,10 @@
}
builtin_name += "Bias";
args.Push(MakeOperand(inst, arg_index).expr);
- image_operands_mask ^= SpvImageOperandsBiasMask;
+ image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Bias);
arg_index++;
}
- if (arg_index < num_args && (image_operands_mask & SpvImageOperandsLodMask)) {
+ if (arg_index < num_args && (image_operands_mask & uint32_t(spv::ImageOperandsMask::Lod))) {
if (use_level_of_detail_suffix) {
builtin_name += "Level";
}
@@ -5553,15 +5630,16 @@
args.Push(lod.expr);
}
- image_operands_mask ^= SpvImageOperandsLodMask;
+ image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Lod);
arg_index++;
- } else if ((opcode == SpvOpImageFetch || opcode == SpvOpImageRead) &&
+ } else if ((op == spv::Op::OpImageFetch || op == spv::Op::OpImageRead) &&
!texture_type->IsAnyOf<DepthMultisampledTexture, MultisampledTexture>()) {
// textureLoad requires an explicit level-of-detail parameter for
// non-multisampled texture types.
args.Push(parser_impl_.MakeNullValue(ty_.I32()));
}
- if (arg_index + 1 < num_args && (image_operands_mask & SpvImageOperandsGradMask)) {
+ if (arg_index + 1 < num_args &&
+ (image_operands_mask & uint32_t(spv::ImageOperandsMask::Grad))) {
if (is_dref_sample) {
return Fail() << "WGSL does not support depth-reference sampling with "
"explicit gradient: "
@@ -5575,11 +5653,12 @@
builtin_name += "Grad";
args.Push(MakeOperand(inst, arg_index).expr);
args.Push(MakeOperand(inst, arg_index + 1).expr);
- image_operands_mask ^= SpvImageOperandsGradMask;
+ image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Grad);
arg_index += 2;
}
- if (arg_index < num_args && (image_operands_mask & SpvImageOperandsConstOffsetMask)) {
- if (!IsImageSamplingOrGatherOrDrefGather(opcode)) {
+ if (arg_index < num_args &&
+ (image_operands_mask & uint32_t(spv::ImageOperandsMask::ConstOffset))) {
+ if (!IsImageSamplingOrGatherOrDrefGather(op)) {
return Fail() << "ConstOffset is only permitted for sampling, gather, or "
"depth-reference gather operations: "
<< inst.PrettyPrint();
@@ -5596,13 +5675,13 @@
}
args.Push(ToSignedIfUnsigned(MakeOperand(inst, arg_index)).expr);
- image_operands_mask ^= SpvImageOperandsConstOffsetMask;
+ image_operands_mask ^= uint32_t(spv::ImageOperandsMask::ConstOffset);
arg_index++;
}
- if (arg_index < num_args && (image_operands_mask & SpvImageOperandsSampleMask)) {
+ if (arg_index < num_args && (image_operands_mask & uint32_t(spv::ImageOperandsMask::Sample))) {
// TODO(dneto): only permitted with ImageFetch
args.Push(ToI32(MakeOperand(inst, arg_index)).expr);
- image_operands_mask ^= SpvImageOperandsSampleMask;
+ image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Sample);
arg_index++;
}
if (image_operands_mask) {
@@ -5642,7 +5721,7 @@
// Construct a 4-element vector with the result from the builtin in the
// first component.
if (texture_type->IsAnyOf<DepthTexture, DepthMultisampledTexture>()) {
- if (is_non_dref_sample || (opcode == SpvOpImageFetch)) {
+ if (is_non_dref_sample || (op == spv::Op::OpImageFetch)) {
value = builder_.Construct(Source{},
result_type->Build(builder_), // a vec4
utils::Vector{
@@ -5658,7 +5737,7 @@
// result type. Compare the SPIR-V image's sampled component type with the
// component of the result type of the SPIR-V instruction.
auto* spirv_image_type = parser_impl_.GetSpirvTypeForHandleMemoryObjectDeclaration(*image);
- if (!spirv_image_type || (spirv_image_type->opcode() != SpvOpTypeImage)) {
+ if (!spirv_image_type || (opcode(spirv_image_type) != spv::Op::OpTypeImage)) {
return Fail() << "invalid image type for image memory object declaration "
<< image->PrettyPrint();
}
@@ -5670,7 +5749,7 @@
value =
create<ast::BitcastExpression>(Source{}, result_type->Build(builder_), call_expr);
}
- if (!expected_component_type->Is<F32>() && IsSampledImageAccess(opcode)) {
+ if (!expected_component_type->Is<F32>() && IsSampledImageAccess(op)) {
// WGSL permits sampled image access only on float textures.
// Reject this case in the SPIR-V reader, at least until SPIR-V validation
// catches up with this rule and can reject it earlier in the workflow.
@@ -5697,10 +5776,10 @@
return false;
}
- const auto opcode = inst.opcode();
- switch (opcode) {
- case SpvOpImageQuerySize:
- case SpvOpImageQuerySizeLod: {
+ const auto op = opcode(inst);
+ switch (op) {
+ case spv::Op::OpImageQuerySize:
+ case spv::Op::OpImageQuerySizeLod: {
ExpressionList exprs;
// Invoke textureDimensions.
// If the texture is arrayed, combine with the result from
@@ -5708,7 +5787,7 @@
auto* dims_ident = create<ast::IdentifierExpression>(
Source{}, builder_.Symbols().Register("textureDimensions"));
ExpressionList dims_args{GetImageExpression(inst)};
- if (opcode == SpvOpImageQuerySizeLod) {
+ if (op == spv::Op::OpImageQuerySizeLod) {
dims_args.Push(MakeOperand(inst, 1).expr);
}
const ast::Expression* dims_call =
@@ -5745,13 +5824,13 @@
return EmitConstDefOrWriteToHoistedVar(inst, expr);
}
- case SpvOpImageQueryLod:
+ case spv::Op::OpImageQueryLod:
return Fail() << "WGSL does not support querying the level of detail of an image: "
<< inst.PrettyPrint();
- case SpvOpImageQueryLevels:
- case SpvOpImageQuerySamples: {
+ case spv::Op::OpImageQueryLevels:
+ case spv::Op::OpImageQuerySamples: {
const auto* name =
- (opcode == SpvOpImageQueryLevels) ? "textureNumLevels" : "textureNumSamples";
+ (op == spv::Op::OpImageQueryLevels) ? "textureNumLevels" : "textureNumSamples";
auto* levels_ident =
create<ast::IdentifierExpression>(Source{}, builder_.Symbols().Register(name));
const ast::Expression* ast_expr = create<ast::CallExpression>(
@@ -5834,46 +5913,46 @@
return {};
};
- switch (inst.opcode()) {
- case SpvOpAtomicLoad:
+ switch (opcode(inst)) {
+ case spv::Op::OpAtomicLoad:
return emit_atomic(sem::BuiltinType::kAtomicLoad, {oper(/*ptr*/ 0)});
- case SpvOpAtomicStore:
+ case spv::Op::OpAtomicStore:
return emit_atomic(sem::BuiltinType::kAtomicStore,
{oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicExchange:
+ case spv::Op::OpAtomicExchange:
return emit_atomic(sem::BuiltinType::kAtomicExchange,
{oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicCompareExchange:
- case SpvOpAtomicCompareExchangeWeak:
+ case spv::Op::OpAtomicCompareExchange:
+ case spv::Op::OpAtomicCompareExchangeWeak:
return emit_atomic(sem::BuiltinType::kAtomicCompareExchangeWeak,
{oper(/*ptr*/ 0), /*value*/ oper(5), /*comparator*/ oper(4)});
- case SpvOpAtomicIIncrement:
+ case spv::Op::OpAtomicIIncrement:
return emit_atomic(sem::BuiltinType::kAtomicAdd, {oper(/*ptr*/ 0), lit(1)});
- case SpvOpAtomicIDecrement:
+ case spv::Op::OpAtomicIDecrement:
return emit_atomic(sem::BuiltinType::kAtomicSub, {oper(/*ptr*/ 0), lit(1)});
- case SpvOpAtomicIAdd:
+ case spv::Op::OpAtomicIAdd:
return emit_atomic(sem::BuiltinType::kAtomicAdd, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicISub:
+ case spv::Op::OpAtomicISub:
return emit_atomic(sem::BuiltinType::kAtomicSub, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicSMin:
+ case spv::Op::OpAtomicSMin:
return emit_atomic(sem::BuiltinType::kAtomicMin, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicUMin:
+ case spv::Op::OpAtomicUMin:
return emit_atomic(sem::BuiltinType::kAtomicMin, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicSMax:
+ case spv::Op::OpAtomicSMax:
return emit_atomic(sem::BuiltinType::kAtomicMax, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicUMax:
+ case spv::Op::OpAtomicUMax:
return emit_atomic(sem::BuiltinType::kAtomicMax, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicAnd:
+ case spv::Op::OpAtomicAnd:
return emit_atomic(sem::BuiltinType::kAtomicAnd, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicOr:
+ case spv::Op::OpAtomicOr:
return emit_atomic(sem::BuiltinType::kAtomicOr, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicXor:
+ case spv::Op::OpAtomicXor:
return emit_atomic(sem::BuiltinType::kAtomicXor, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
- case SpvOpAtomicFlagTestAndSet:
- case SpvOpAtomicFlagClear:
- case SpvOpAtomicFMinEXT:
- case SpvOpAtomicFMaxEXT:
- case SpvOpAtomicFAddEXT:
+ case spv::Op::OpAtomicFlagTestAndSet:
+ case spv::Op::OpAtomicFlagClear:
+ case spv::Op::OpAtomicFMinEXT:
+ case spv::Op::OpAtomicFMaxEXT:
+ case spv::Op::OpAtomicFAddEXT:
return Fail() << "unsupported atomic op: " << inst.PrettyPrint();
default:
@@ -5926,11 +6005,11 @@
<< " prompted by " << inst.PrettyPrint();
}
bool is_proj = false;
- switch (inst.opcode()) {
- case SpvOpImageSampleProjImplicitLod:
- case SpvOpImageSampleProjExplicitLod:
- case SpvOpImageSampleProjDrefImplicitLod:
- case SpvOpImageSampleProjDrefExplicitLod:
+ switch (opcode(inst)) {
+ case spv::Op::OpImageSampleProjImplicitLod:
+ case spv::Op::OpImageSampleProjExplicitLod:
+ case spv::Op::OpImageSampleProjDrefImplicitLod:
+ case spv::Op::OpImageSampleProjDrefExplicitLod:
is_proj = true;
break;
default:
diff --git a/src/tint/reader/spirv/parser_impl.cc b/src/tint/reader/spirv/parser_impl.cc
index 13b4132..f055f4c 100644
--- a/src/tint/reader/spirv/parser_impl.cc
+++ b/src/tint/reader/spirv/parser_impl.cc
@@ -34,6 +34,65 @@
namespace tint::reader::spirv {
+namespace three_sided_patch {
+// This machinery is only used while SPIRV-Tools is in transition before it fully
+// uses the C++11 header spirv.hpp11
+
+/// Typedef for pointer to member function while the API call uses
+/// SpvStorageClass for its second argument.
+typedef uint32_t (
+ spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClass)(uint32_t, SpvStorageClass);
+/// Typedef for pointer to member function while the API call uses
+/// spv::StorageClass for its second argument.
+typedef uint32_t (spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClassCpp11)(
+ uint32_t,
+ spv::StorageClass);
+
+/// @param type_manager the SPIRV-Tools optimizer's type manager
+/// @param finder a pointer to member function in the type manager that does the
+/// actual lookup
+/// @param pointee_type_id the ID of the pointee type
+/// @param sc the storage class. SC can be SpvStorageClass or spv::StorageClass
+/// @returns the ID for a SPIR-V pointer to pointee_type_id in storage class sc
+template <typename FinderType, typename SC>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_manager,
+ FinderType finder,
+ uint32_t pointee_type_id,
+ SC sc);
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClass finder,
+ uint32_t pointee_type_id,
+ SpvStorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, sc);
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClass finder,
+ uint32_t pointee_type_id,
+ spv::StorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, static_cast<SpvStorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClassCpp11 finder,
+ uint32_t pointee_type_id,
+ SpvStorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, static_cast<spv::StorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+ PointerFinderSpvStorageClassCpp11 finder,
+ uint32_t pointee_type_id,
+ spv::StorageClass sc) {
+ return (type_mgr->*finder)(pointee_type_id, sc);
+}
+} // namespace three_sided_patch
+
namespace {
// Input SPIR-V needs only to conform to Vulkan 1.1 requirements.
@@ -42,6 +101,17 @@
// will satisfy SPV_ENV_WEBGPU_0 validation.
const spv_target_env kInputEnv = SPV_ENV_VULKAN_1_1;
+/// @param inst a SPIR-V instruction
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction& inst) {
+ return static_cast<spv::Op>(inst.opcode());
+}
+/// @param inst a SPIR-V instruction pointer
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction* inst) {
+ return static_cast<spv::Op>(inst->opcode());
+}
+
// A FunctionTraverser is used to compute an ordering of functions in the
// module such that callees precede callers.
class FunctionTraverser {
@@ -70,7 +140,7 @@
visited_.insert(&f);
for (const auto& bb : f) {
for (const auto& inst : bb) {
- if (inst.opcode() != SpvOpFunctionCall) {
+ if (opcode(inst) != spv::Op::OpFunctionCall) {
continue;
}
const auto* callee = id_to_func_[inst.GetSingleWordInOperand(0)];
@@ -89,17 +159,17 @@
};
// Returns true if the opcode operates as if its operands are signed integral.
-bool AssumesSignedOperands(SpvOp opcode) {
+bool AssumesSignedOperands(spv::Op opcode) {
switch (opcode) {
- case SpvOpSNegate:
- case SpvOpSDiv:
- case SpvOpSRem:
- case SpvOpSMod:
- case SpvOpSLessThan:
- case SpvOpSLessThanEqual:
- case SpvOpSGreaterThan:
- case SpvOpSGreaterThanEqual:
- case SpvOpConvertSToF:
+ case spv::Op::OpSNegate:
+ case spv::Op::OpSDiv:
+ case spv::Op::OpSRem:
+ case spv::Op::OpSMod:
+ case spv::Op::OpSLessThan:
+ case spv::Op::OpSLessThanEqual:
+ case spv::Op::OpSGreaterThan:
+ case spv::Op::OpSGreaterThanEqual:
+ case spv::Op::OpConvertSToF:
return true;
default:
break;
@@ -126,15 +196,15 @@
}
// Returns true if the opcode operates as if its operands are unsigned integral.
-bool AssumesUnsignedOperands(SpvOp opcode) {
+bool AssumesUnsignedOperands(spv::Op opcode) {
switch (opcode) {
- case SpvOpUDiv:
- case SpvOpUMod:
- case SpvOpULessThan:
- case SpvOpULessThanEqual:
- case SpvOpUGreaterThan:
- case SpvOpUGreaterThanEqual:
- case SpvOpConvertUToF:
+ case spv::Op::OpUDiv:
+ case spv::Op::OpUMod:
+ case spv::Op::OpULessThan:
+ case spv::Op::OpULessThanEqual:
+ case spv::Op::OpUGreaterThan:
+ case spv::Op::OpUGreaterThanEqual:
+ case spv::Op::OpConvertUToF:
return true;
default:
break;
@@ -163,18 +233,18 @@
// the signedness of the second operand to match the signedness of the
// first operand, and it's not one of the OpU* or OpS* instructions.
// (Those are handled via MakeOperand.)
-bool AssumesSecondOperandSignednessMatchesFirstOperand(SpvOp opcode) {
+bool AssumesSecondOperandSignednessMatchesFirstOperand(spv::Op opcode) {
switch (opcode) {
// All the OpI* integer binary operations.
- case SpvOpIAdd:
- case SpvOpISub:
- case SpvOpIMul:
- case SpvOpIEqual:
- case SpvOpINotEqual:
+ case spv::Op::OpIAdd:
+ case spv::Op::OpISub:
+ case spv::Op::OpIMul:
+ case spv::Op::OpIEqual:
+ case spv::Op::OpINotEqual:
// All the bitwise integer binary operations.
- case SpvOpBitwiseAnd:
- case SpvOpBitwiseOr:
- case SpvOpBitwiseXor:
+ case spv::Op::OpBitwiseAnd:
+ case spv::Op::OpBitwiseOr:
+ case spv::Op::OpBitwiseXor:
return true;
default:
break;
@@ -184,24 +254,24 @@
// Returns true if the corresponding WGSL operation requires
// the signedness of the result to match the signedness of the first operand.
-bool AssumesResultSignednessMatchesFirstOperand(SpvOp opcode) {
+bool AssumesResultSignednessMatchesFirstOperand(spv::Op opcode) {
switch (opcode) {
- case SpvOpNot:
- case SpvOpSNegate:
- case SpvOpBitCount:
- case SpvOpBitReverse:
- case SpvOpSDiv:
- case SpvOpSMod:
- case SpvOpSRem:
- case SpvOpIAdd:
- case SpvOpISub:
- case SpvOpIMul:
- case SpvOpBitwiseAnd:
- case SpvOpBitwiseOr:
- case SpvOpBitwiseXor:
- case SpvOpShiftLeftLogical:
- case SpvOpShiftRightLogical:
- case SpvOpShiftRightArithmetic:
+ case spv::Op::OpNot:
+ case spv::Op::OpSNegate:
+ case spv::Op::OpBitCount:
+ case spv::Op::OpBitReverse:
+ case spv::Op::OpSDiv:
+ case spv::Op::OpSMod:
+ case spv::Op::OpSRem:
+ case spv::Op::OpIAdd:
+ case spv::Op::OpISub:
+ case spv::Op::OpIMul:
+ case spv::Op::OpBitwiseAnd:
+ case spv::Op::OpBitwiseOr:
+ case spv::Op::OpBitwiseXor:
+ case spv::Op::OpShiftLeftLogical:
+ case spv::Op::OpShiftRightLogical:
+ case spv::Op::OpShiftRightArithmetic:
return true;
default:
break;
@@ -240,12 +310,12 @@
if (deco.size() < 1) {
return false;
}
- switch (deco[0]) {
- case SpvDecorationLocation:
- case SpvDecorationFlat:
- case SpvDecorationNoPerspective:
- case SpvDecorationCentroid:
- case SpvDecorationSample:
+ switch (static_cast<spv::Decoration>(deco[0])) {
+ case spv::Decoration::Location:
+ case spv::Decoration::Flat:
+ case spv::Decoration::NoPerspective:
+ case spv::Decoration::Centroid:
+ case spv::Decoration::Sample:
return true;
default:
break;
@@ -388,16 +458,16 @@
const auto& decorations = deco_mgr_->GetDecorationsFor(id, true);
std::unordered_set<uint32_t> visited;
for (const auto* inst : decorations) {
- if (inst->opcode() != SpvOpDecorate) {
+ if (opcode(inst) != spv::Op::OpDecorate) {
continue;
}
// Example: OpDecorate %struct_id Block
// Example: OpDecorate %array_ty ArrayStride 16
auto decoration_kind = inst->GetSingleWordInOperand(1);
- switch (decoration_kind) {
+ switch (static_cast<spv::Decoration>(decoration_kind)) {
// Restrict and RestrictPointer have no effect in graphics APIs.
- case SpvDecorationRestrict:
- case SpvDecorationRestrictPointer:
+ case spv::Decoration::Restrict:
+ case spv::Decoration::RestrictPointer:
break;
default:
if (visited.emplace(decoration_kind).second) {
@@ -418,15 +488,15 @@
std::unordered_set<uint32_t> visited;
for (const auto* inst : decorations) {
// Example: OpMemberDecorate %struct_id 1 Offset 16
- if ((inst->opcode() != SpvOpMemberDecorate) ||
+ if ((opcode(inst) != spv::Op::OpMemberDecorate) ||
(inst->GetSingleWordInOperand(1) != member_index)) {
continue;
}
auto decoration_kind = inst->GetSingleWordInOperand(2);
- switch (decoration_kind) {
+ switch (static_cast<spv::Decoration>(decoration_kind)) {
// Restrict and RestrictPointer have no effect in graphics APIs.
- case SpvDecorationRestrict:
- case SpvDecorationRestrictPointer:
+ case spv::Decoration::Restrict:
+ case spv::Decoration::RestrictPointer:
break;
default:
if (visited.emplace(decoration_kind).second) {
@@ -458,8 +528,8 @@
Fail() << "malformed SPIR-V decoration: it's empty";
return {};
}
- switch (decoration[0]) {
- case SpvDecorationOffset:
+ switch (static_cast<spv::Decoration>(decoration[0])) {
+ case spv::Decoration::Offset:
if (decoration.size() != 2) {
Fail() << "malformed Offset decoration: expected 1 literal operand, has "
<< decoration.size() - 1 << ": member " << member_index << " of "
@@ -469,24 +539,24 @@
return {
builder_.MemberOffset(Source{}, AInt(decoration[1])),
};
- case SpvDecorationNonReadable:
+ case spv::Decoration::NonReadable:
// WGSL doesn't have a member decoration for this. Silently drop it.
return {};
- case SpvDecorationNonWritable:
+ case spv::Decoration::NonWritable:
// WGSL doesn't have a member decoration for this.
return {};
- case SpvDecorationColMajor:
+ case spv::Decoration::ColMajor:
// WGSL only supports column major matrices.
return {};
- case SpvDecorationRelaxedPrecision:
+ case spv::Decoration::RelaxedPrecision:
// WGSL doesn't support relaxed precision.
return {};
- case SpvDecorationRowMajor:
+ case spv::Decoration::RowMajor:
Fail() << "WGSL does not support row-major matrices: can't "
"translate member "
<< member_index << " of " << ShowType(struct_type_id);
return {};
- case SpvDecorationMatrixStride: {
+ case spv::Decoration::MatrixStride: {
if (decoration.size() != 2) {
Fail() << "malformed MatrixStride decoration: expected 1 literal "
"operand, has "
@@ -589,15 +659,15 @@
[this, &in_op_line_scope, &op_line_source,
&instruction_number](const spvtools::opt::Instruction* inst) {
++instruction_number.line;
- switch (inst->opcode()) {
- case SpvOpLine:
+ switch (opcode(inst)) {
+ case spv::Op::OpLine:
in_op_line_scope = true;
// TODO(dneto): This ignores the File ID (operand 0), since the Tint
// Source concept doesn't represent that.
op_line_source.line = inst->GetSingleWordInOperand(1);
op_line_source.column = inst->GetSingleWordInOperand(2);
break;
- case SpvOpNoLine:
+ case spv::Op::OpNoLine:
in_op_line_scope = false;
break;
default:
@@ -670,12 +740,12 @@
}
bool ParserImpl::IsGlslExtendedInstruction(const spvtools::opt::Instruction& inst) const {
- return (inst.opcode() == SpvOpExtInst) &&
+ return (opcode(inst) == spv::Op::OpExtInst) &&
(glsl_std_450_imports_.count(inst.GetSingleWordInOperand(0)) > 0);
}
bool ParserImpl::IsIgnoredExtendedInstruction(const spvtools::opt::Instruction& inst) const {
- return (inst.opcode() == SpvOpExtInst) &&
+ return (opcode(inst) == spv::Op::OpExtInst) &&
(ignored_imports_.count(inst.GetSingleWordInOperand(0)) > 0);
}
@@ -722,15 +792,15 @@
// Register names from OpName and OpMemberName
for (const auto& inst : module_->debugs2()) {
- switch (inst.opcode()) {
- case SpvOpName: {
+ switch (opcode(inst)) {
+ case spv::Op::OpName: {
const auto name = inst.GetInOperand(1).AsString();
if (!name.empty()) {
namer_.SuggestSanitizedName(inst.GetSingleWordInOperand(0), name);
}
break;
}
- case SpvOpMemberName: {
+ case spv::Op::OpMemberName: {
const auto name = inst.GetInOperand(2).AsString();
if (!name.empty()) {
namer_.SuggestSanitizedMemberName(inst.GetSingleWordInOperand(0),
@@ -745,7 +815,7 @@
// Fill in struct member names, and disambiguate them.
for (const auto* type_inst : module_->GetTypes()) {
- if (type_inst->opcode() == SpvOpTypeStruct) {
+ if (opcode(type_inst) == spv::Op::OpTypeStruct) {
namer_.ResolveMemberNamesForStruct(type_inst->result_id(), type_inst->NumInOperands());
}
}
@@ -779,13 +849,13 @@
bool ParserImpl::RegisterWorkgroupSizeBuiltin() {
WorkgroupSizeInfo& info = workgroup_size_builtin_;
for (const spvtools::opt::Instruction& inst : module_->annotations()) {
- if (inst.opcode() != SpvOpDecorate) {
+ if (opcode(inst) != spv::Op::OpDecorate) {
continue;
}
- if (inst.GetSingleWordInOperand(1) != SpvDecorationBuiltIn) {
+ if (inst.GetSingleWordInOperand(1) != uint32_t(spv::Decoration::BuiltIn)) {
continue;
}
- if (inst.GetSingleWordInOperand(2) != SpvBuiltInWorkgroupSize) {
+ if (inst.GetSingleWordInOperand(2) != uint32_t(spv::BuiltIn::WorkgroupSize)) {
continue;
}
info.id = inst.GetSingleWordInOperand(0);
@@ -806,8 +876,8 @@
// const-expr yet, so avoid supporting OpSpecConstantOp here.
// TODO(dneto): See https://github.com/gpuweb/gpuweb/issues/1272 for WGSL
// const_expr proposals.
- if ((composite_def->opcode() != SpvOpSpecConstantComposite &&
- composite_def->opcode() != SpvOpConstantComposite)) {
+ if ((opcode(composite_def) != spv::Op::OpSpecConstantComposite &&
+ opcode(composite_def) != spv::Op::OpConstantComposite)) {
return Fail() << "Invalid WorkgroupSize builtin. Expected 3-element "
"OpSpecConstantComposite or OpConstantComposite: "
<< composite_def->PrettyPrint();
@@ -822,7 +892,8 @@
int index) -> bool {
const auto id = composite_def->GetSingleWordInOperand(static_cast<uint32_t>(index));
const auto* def = def_use_mgr_->GetDef(id);
- if (!def || (def->opcode() != SpvOpSpecConstant && def->opcode() != SpvOpConstant) ||
+ if (!def ||
+ (opcode(def) != spv::Op::OpSpecConstant && opcode(def) != spv::Op::OpConstant) ||
(def->NumInOperands() != 1)) {
return Fail() << "invalid component " << index << " of workgroupsize "
<< (def ? def->PrettyPrint() : std::string("no definition"));
@@ -842,8 +913,8 @@
// decorations.
std::unordered_map<uint32_t, GridSize> local_size;
for (const spvtools::opt::Instruction& inst : module_->execution_modes()) {
- auto mode = static_cast<SpvExecutionMode>(inst.GetSingleWordInOperand(1));
- if (mode == SpvExecutionModeLocalSize) {
+ auto mode = static_cast<spv::ExecutionMode>(inst.GetSingleWordInOperand(1));
+ if (mode == spv::ExecutionMode::LocalSize) {
if (inst.NumInOperands() != 5) {
// This won't even get past SPIR-V binary parsing.
return Fail() << "invalid LocalSize execution mode: " << inst.PrettyPrint();
@@ -856,7 +927,7 @@
}
for (const spvtools::opt::Instruction& entry_point : module_->entry_points()) {
- const auto stage = SpvExecutionModel(entry_point.GetSingleWordInOperand(0));
+ const auto stage = spv::ExecutionModel(entry_point.GetSingleWordInOperand(0));
const uint32_t function_id = entry_point.GetSingleWordInOperand(1);
const std::string ep_name = entry_point.GetOperand(2).AsString();
@@ -886,11 +957,11 @@
for (unsigned iarg = 3; iarg < entry_point.NumInOperands(); iarg++) {
const uint32_t var_id = entry_point.GetSingleWordInOperand(iarg);
if (const auto* var_inst = def_use_mgr_->GetDef(var_id)) {
- switch (SpvStorageClass(var_inst->GetSingleWordInOperand(0))) {
- case SpvStorageClassInput:
+ switch (spv::StorageClass(var_inst->GetSingleWordInOperand(0))) {
+ case spv::StorageClass::Input:
inputs.Add(var_id);
break;
- case SpvStorageClassOutput:
+ case spv::StorageClass::Output:
outputs.Add(var_id);
break;
default:
@@ -1036,7 +1107,7 @@
*array_stride = 0; // Implicit stride case.
const auto type_id = type_mgr_->GetId(spv_type);
for (auto& decoration : this->GetDecorationsFor(type_id)) {
- if (decoration.size() == 2 && decoration[0] == SpvDecorationArrayStride) {
+ if (decoration.size() == 2 && decoration[0] == uint32_t(spv::Decoration::ArrayStride)) {
const auto stride = decoration[1];
if (stride == 0) {
return Fail() << "invalid array type ID " << type_id << ": ArrayStride can't be 0";
@@ -1057,9 +1128,9 @@
auto struct_decorations = this->GetDecorationsFor(type_id);
if (struct_decorations.size() == 1) {
const auto decoration = struct_decorations[0][0];
- if (decoration == SpvDecorationBufferBlock) {
+ if (decoration == uint32_t(spv::Decoration::BufferBlock)) {
remap_buffer_block_type_.insert(type_id);
- } else if (decoration != SpvDecorationBlock) {
+ } else if (decoration != uint32_t(spv::Decoration::Block)) {
Fail() << "struct with ID " << type_id
<< " has unrecognized decoration: " << int(decoration);
}
@@ -1097,21 +1168,21 @@
Fail() << "malformed SPIR-V decoration: it's empty";
return nullptr;
}
- if ((decoration[0] == SpvDecorationBuiltIn) && (decoration.size() > 1)) {
- switch (decoration[1]) {
- case SpvBuiltInPosition:
+ if ((decoration[0] == uint32_t(spv::Decoration::BuiltIn)) && (decoration.size() > 1)) {
+ switch (static_cast<spv::BuiltIn>(decoration[1])) {
+ case spv::BuiltIn::Position:
// Record this built-in variable specially.
builtin_position_.struct_type_id = type_id;
builtin_position_.position_member_index = member_index;
builtin_position_.position_member_type_id = member_type_id;
create_ast_member = false; // Not part of the WGSL structure.
break;
- case SpvBuiltInPointSize: // not supported in WGSL, but ignore
+ case spv::BuiltIn::PointSize: // not supported in WGSL, but ignore
builtin_position_.pointsize_member_index = member_index;
create_ast_member = false; // Not part of the WGSL structure.
break;
- case SpvBuiltInClipDistance: // not supported in WGSL
- case SpvBuiltInCullDistance: // not supported in WGSL
+ case spv::BuiltIn::ClipDistance: // not supported in WGSL
+ case spv::BuiltIn::CullDistance: // not supported in WGSL
create_ast_member = false; // Not part of the WGSL structure.
break;
default:
@@ -1131,7 +1202,7 @@
if (IsPipelineDecoration(decoration)) {
// IO decorations are handled when emitting the entry point.
continue;
- } else if (decoration[0] == SpvDecorationNonWritable) {
+ } else if (decoration[0] == uint32_t(spv::Decoration::NonWritable)) {
// WGSL doesn't represent individual members as non-writable. Instead,
// apply the ReadOnly access control to the containing struct if all
// the members are non-writable.
@@ -1193,12 +1264,12 @@
const spvtools::opt::analysis::Pointer*) {
const auto* inst = def_use_mgr_->GetDef(type_id);
const auto pointee_type_id = inst->GetSingleWordInOperand(1);
- const auto storage_class = SpvStorageClass(inst->GetSingleWordInOperand(0));
+ const auto storage_class = spv::StorageClass(inst->GetSingleWordInOperand(0));
if (pointee_type_id == builtin_position_.struct_type_id) {
builtin_position_.pointer_type_id = type_id;
// Pipeline IO builtins map to private variables.
- builtin_position_.storage_class = SpvStorageClassPrivate;
+ builtin_position_.storage_class = spv::StorageClass::Private;
return nullptr;
}
auto* ast_elem_ty = ConvertType(pointee_type_id, PtrAs::Ptr);
@@ -1244,21 +1315,21 @@
// in WGSL. In particular, exclude user-defined pipeline IO in a
// block-decorated struct.
for (const auto& type_or_value : module_->types_values()) {
- if (type_or_value.opcode() != SpvOpVariable) {
+ if (opcode(type_or_value) != spv::Op::OpVariable) {
continue;
}
const auto& var = type_or_value;
- const auto spirv_storage_class = SpvStorageClass(var.GetSingleWordInOperand(0));
- if ((spirv_storage_class != SpvStorageClassStorageBuffer) &&
- (spirv_storage_class != SpvStorageClassUniform)) {
+ const auto spirv_storage_class = spv::StorageClass(var.GetSingleWordInOperand(0));
+ if ((spirv_storage_class != spv::StorageClass::StorageBuffer) &&
+ (spirv_storage_class != spv::StorageClass::Uniform)) {
continue;
}
const auto* ptr_type = def_use_mgr_->GetDef(var.type_id());
- if (ptr_type->opcode() != SpvOpTypePointer) {
+ if (opcode(ptr_type) != spv::Op::OpTypePointer) {
return Fail() << "OpVariable type expected to be a pointer: " << var.PrettyPrint();
}
const auto* store_type = def_use_mgr_->GetDef(ptr_type->GetSingleWordInOperand(1));
- if (store_type->opcode() == SpvOpTypeStruct) {
+ if (opcode(store_type) == spv::Op::OpTypeStruct) {
struct_types_for_buffers_.insert(store_type->result_id());
} else {
Fail() << "WGSL does not support arrays of buffers: " << var.PrettyPrint();
@@ -1276,7 +1347,8 @@
// Manufacture a type for the gl_Position variable if we have to.
if ((builtin_position_.struct_type_id != 0) &&
(builtin_position_.position_member_pointer_type_id == 0)) {
- builtin_position_.position_member_pointer_type_id = type_mgr_->FindPointerToType(
+ builtin_position_.position_member_pointer_type_id = three_sided_patch::FindPointerToType(
+ type_mgr_, &spvtools::opt::analysis::TypeManager::FindPointerToType,
builtin_position_.position_member_type_id, builtin_position_.storage_class);
ConvertType(builtin_position_.position_member_pointer_type_id);
}
@@ -1290,13 +1362,13 @@
for (auto& inst : module_->types_values()) {
if (const auto* result_type = type_mgr_->GetType(inst.type_id())) {
if (result_type->AsPointer()) {
- switch (inst.opcode()) {
- case SpvOpVariable:
+ switch (opcode(inst)) {
+ case spv::Op::OpVariable:
// This is the only valid case.
break;
- case SpvOpUndef:
+ case spv::Op::OpUndef:
return Fail() << "undef pointer is not valid: " << inst.PrettyPrint();
- case SpvOpConstantNull:
+ case spv::Op::OpConstantNull:
return Fail() << "null pointer is not valid: " << inst.PrettyPrint();
default:
return Fail()
@@ -1319,15 +1391,15 @@
const Type* ast_type = nullptr;
ast::LiteralExpression* ast_expr = nullptr;
- switch (inst.opcode()) {
- case SpvOpSpecConstantTrue:
- case SpvOpSpecConstantFalse: {
+ switch (opcode(inst)) {
+ case spv::Op::OpSpecConstantTrue:
+ case spv::Op::OpSpecConstantFalse: {
ast_type = ConvertType(inst.type_id());
ast_expr = create<ast::BoolLiteralExpression>(
- Source{}, inst.opcode() == SpvOpSpecConstantTrue);
+ Source{}, opcode(inst) == spv::Op::OpSpecConstantTrue);
break;
}
- case SpvOpSpecConstant: {
+ case spv::Op::OpSpecConstant: {
ast_type = ConvertType(inst.type_id());
const uint32_t literal_value = inst.GetSingleWordInOperand(0);
ast_expr = Switch(
@@ -1362,7 +1434,7 @@
if (ast_type && ast_expr) {
AttributeList spec_id_decos;
for (const auto& deco : GetDecorationsFor(inst.result_id())) {
- if ((deco.size() == 2) && (deco[0] == SpvDecorationSpecId)) {
+ if ((deco.size() == 2) && (deco[0] == uint32_t(spv::Decoration::SpecId))) {
const uint32_t id = deco[1];
if (id > 65535) {
return Fail() << "SpecId too large. WGSL override IDs must be "
@@ -1430,16 +1502,16 @@
return false;
}
for (const auto& type_or_value : module_->types_values()) {
- if (type_or_value.opcode() != SpvOpVariable) {
+ if (opcode(type_or_value) != spv::Op::OpVariable) {
continue;
}
const auto& var = type_or_value;
- const auto spirv_storage_class = SpvStorageClass(var.GetSingleWordInOperand(0));
+ const auto spirv_storage_class = spv::StorageClass(var.GetSingleWordInOperand(0));
uint32_t type_id = var.type_id();
if ((type_id == builtin_position_.pointer_type_id) &&
- ((spirv_storage_class == SpvStorageClassInput) ||
- (spirv_storage_class == SpvStorageClassOutput))) {
+ ((spirv_storage_class == spv::StorageClass::Input) ||
+ (spirv_storage_class == spv::StorageClass::Output))) {
// Skip emitting gl_PerVertex.
builtin_position_.per_vertex_var_id = var.result_id();
builtin_position_.per_vertex_var_init_id =
@@ -1464,7 +1536,7 @@
return false;
}
const Type* ast_type = nullptr;
- if (spirv_storage_class == SpvStorageClassUniformConstant) {
+ if (spirv_storage_class == spv::StorageClass::UniformConstant) {
// These are opaque handles: samplers or textures
ast_type = GetTypeForHandleVar(var);
if (!ast_type) {
@@ -1509,9 +1581,9 @@
if (builtin_position_.per_vertex_var_init_id) {
// The initializer is complex.
const auto* init = def_use_mgr_->GetDef(builtin_position_.per_vertex_var_init_id);
- switch (init->opcode()) {
- case SpvOpConstantComposite:
- case SpvOpSpecConstantComposite:
+ switch (opcode(init)) {
+ case spv::Op::OpConstantComposite:
+ case spv::Op::OpSpecConstantComposite:
ast_initializer =
MakeConstantExpression(
init->GetSingleWordInOperand(builtin_position_.position_member_index))
@@ -1541,15 +1613,15 @@
// @returns the IntConstant for the size of the array, or nullptr
const spvtools::opt::analysis::IntConstant* ParserImpl::GetArraySize(uint32_t var_id) {
auto* var = def_use_mgr_->GetDef(var_id);
- if (!var || var->opcode() != SpvOpVariable) {
+ if (!var || opcode(var) != spv::Op::OpVariable) {
return nullptr;
}
auto* ptr_type = def_use_mgr_->GetDef(var->type_id());
- if (!ptr_type || ptr_type->opcode() != SpvOpTypePointer) {
+ if (!ptr_type || opcode(ptr_type) != spv::Op::OpTypePointer) {
return nullptr;
}
auto* array_type = def_use_mgr_->GetDef(ptr_type->GetSingleWordInOperand(1));
- if (!array_type || array_type->opcode() != SpvOpTypeArray) {
+ if (!array_type || opcode(array_type) != spv::Op::OpTypeArray) {
return nullptr;
}
auto* size = constant_mgr_->FindDeclaredConstant(array_type->GetSingleWordInOperand(1));
@@ -1632,23 +1704,23 @@
if (deco.empty()) {
return Fail() << "malformed decoration on ID " << id << ": it is empty";
}
- if (deco[0] == SpvDecorationBuiltIn) {
+ if (deco[0] == uint32_t(spv::Decoration::BuiltIn)) {
if (deco.size() == 1) {
return Fail() << "malformed BuiltIn decoration on ID " << id << ": has no operand";
}
- const auto spv_builtin = static_cast<SpvBuiltIn>(deco[1]);
+ const auto spv_builtin = static_cast<spv::BuiltIn>(deco[1]);
switch (spv_builtin) {
- case SpvBuiltInPointSize:
+ case spv::BuiltIn::PointSize:
special_builtins_[id] = spv_builtin;
return false; // This is not an error
- case SpvBuiltInSampleId:
- case SpvBuiltInVertexIndex:
- case SpvBuiltInInstanceIndex:
- case SpvBuiltInLocalInvocationId:
- case SpvBuiltInLocalInvocationIndex:
- case SpvBuiltInGlobalInvocationId:
- case SpvBuiltInWorkgroupId:
- case SpvBuiltInNumWorkgroups:
+ case spv::BuiltIn::SampleId:
+ case spv::BuiltIn::VertexIndex:
+ case spv::BuiltIn::InstanceIndex:
+ case spv::BuiltIn::LocalInvocationId:
+ case spv::BuiltIn::LocalInvocationIndex:
+ case spv::BuiltIn::GlobalInvocationId:
+ case spv::BuiltIn::WorkgroupId:
+ case spv::BuiltIn::NumWorkgroups:
// The SPIR-V variable may signed (because GLSL requires signed for
// some of these), but WGSL requires unsigned. Handle specially
// so we always perform the conversion at load and store.
@@ -1660,7 +1732,7 @@
}
}
break;
- case SpvBuiltInSampleMask: {
+ case spv::BuiltIn::SampleMask: {
// In SPIR-V this is used for both input and output variable.
// The SPIR-V variable has store type of array of integer scalar,
// either signed or unsigned.
@@ -1691,14 +1763,14 @@
if (transfer_pipeline_io && IsPipelineDecoration(deco)) {
non_builtin_pipeline_decorations.push_back(deco);
}
- if (deco[0] == SpvDecorationDescriptorSet) {
+ if (deco[0] == uint32_t(spv::Decoration::DescriptorSet)) {
if (deco.size() == 1) {
return Fail() << "malformed DescriptorSet decoration on ID " << id
<< ": has no operand";
}
decorations->Push(builder_.Group(Source{}, AInt(deco[1])));
}
- if (deco[0] == SpvDecorationBinding) {
+ if (deco[0] == uint32_t(spv::Decoration::Binding)) {
if (deco.size() == 1) {
return Fail() << "malformed Binding decoration on ID " << id << ": has no operand";
}
@@ -1756,8 +1828,8 @@
for (const auto& deco : decorations) {
TINT_ASSERT(Reader, deco.size() > 0);
- switch (deco[0]) {
- case SpvDecorationLocation:
+ switch (static_cast<spv::Decoration>(deco[0])) {
+ case spv::Decoration::Location:
if (deco.size() != 2) {
return Fail() << "malformed Location decoration on ID requires one "
"literal operand";
@@ -1768,24 +1840,24 @@
type = ast::InterpolationType::kFlat;
}
break;
- case SpvDecorationFlat:
+ case spv::Decoration::Flat:
type = ast::InterpolationType::kFlat;
break;
- case SpvDecorationNoPerspective:
+ case spv::Decoration::NoPerspective:
if (store_type->IsIntegerScalarOrVector()) {
// This doesn't capture the array or struct case.
return Fail() << "NoPerspective is invalid on integral IO";
}
type = ast::InterpolationType::kLinear;
break;
- case SpvDecorationCentroid:
+ case spv::Decoration::Centroid:
if (store_type->IsIntegerScalarOrVector()) {
// This doesn't capture the array or struct case.
return Fail() << "Centroid interpolation sampling is invalid on integral IO";
}
sampling = ast::InterpolationSampling::kCentroid;
break;
- case SpvDecorationSample:
+ case spv::Decoration::Sample:
if (store_type->IsIntegerScalarOrVector()) {
// This doesn't capture the array or struct case.
return Fail() << "Sample interpolation sampling is invalid on integral IO";
@@ -1801,7 +1873,7 @@
!ast::HasAttribute<ast::LocationAttribute>(*attributes)) {
// WGSL requires that '@interpolate(flat)' needs to be paired with '@location', however
// SPIR-V requires all fragment shader integer Inputs are 'flat'. If the decorations do not
- // contain a SpvDecorationLocation, then make this perspective.
+ // contain a spv::Decoration::Location, then make this perspective.
type = ast::InterpolationType::kPerspective;
}
@@ -1825,7 +1897,7 @@
if (!inst) {
return false;
}
- if (inst->opcode() == SpvOpUndef) {
+ if (opcode(inst) == spv::Op::OpUndef) {
return true;
}
return nullptr != constant_mgr_->FindDeclaredConstant(id);
@@ -1880,13 +1952,13 @@
return {};
}
- switch (inst->opcode()) {
- case SpvOpUndef: // Remap undef to null.
- case SpvOpConstantNull:
+ switch (opcode(inst)) {
+ case spv::Op::OpUndef: // Remap undef to null.
+ case spv::Op::OpConstantNull:
return {original_ast_type, MakeNullValue(original_ast_type)};
- case SpvOpConstantTrue:
- case SpvOpConstantFalse:
- case SpvOpConstant: {
+ case spv::Op::OpConstantTrue:
+ case spv::Op::OpConstantFalse:
+ case spv::Op::OpConstant: {
const auto* spirv_const = constant_mgr_->FindDeclaredConstant(id);
if (spirv_const == nullptr) {
Fail() << "ID " << id << " is not a constant";
@@ -1895,7 +1967,7 @@
return MakeConstantExpressionForScalarSpirvConstant(source, original_ast_type,
spirv_const);
}
- case SpvOpConstantComposite: {
+ case spv::Op::OpConstantComposite: {
// Handle vector, matrix, array, and struct
// Generate a composite from explicit components.
@@ -2047,9 +2119,9 @@
requires_signed = AssumesSignedOperands(extended_opcode);
requires_unsigned = AssumesUnsignedOperands(extended_opcode);
} else {
- const auto opcode = inst.opcode();
- requires_signed = AssumesSignedOperands(opcode);
- requires_unsigned = AssumesUnsignedOperands(opcode);
+ const auto op = opcode(inst);
+ requires_signed = AssumesSignedOperands(op);
+ requires_unsigned = AssumesUnsignedOperands(op);
}
if (!requires_signed && !requires_unsigned) {
// No conversion is required, assuming our tables are complete.
@@ -2087,7 +2159,7 @@
TypedExpression&& second_operand_expr) {
const Type* target_type = first_operand_type->UnwrapRef();
if ((target_type != second_operand_expr.type->UnwrapRef()) &&
- AssumesSecondOperandSignednessMatchesFirstOperand(inst.opcode())) {
+ AssumesSecondOperandSignednessMatchesFirstOperand(opcode(inst))) {
// Conversion is required.
return {target_type, create<ast::BitcastExpression>(Source{}, target_type->Build(builder_),
second_operand_expr.expr)};
@@ -2099,8 +2171,8 @@
const Type* ParserImpl::ForcedResultType(const spvtools::opt::Instruction& inst,
const Type* first_operand_type) {
first_operand_type = first_operand_type->UnwrapRef();
- const auto opcode = inst.opcode();
- if (AssumesResultSignednessMatchesFirstOperand(opcode)) {
+ const auto op = opcode(inst);
+ if (AssumesResultSignednessMatchesFirstOperand(op)) {
return first_operand_type;
}
if (IsGlslExtendedInstruction(inst)) {
@@ -2239,36 +2311,36 @@
if (inst == nullptr) {
return local_fail();
}
- switch (inst->opcode()) {
- case SpvOpFunctionParameter:
- case SpvOpVariable:
+ switch (opcode(inst)) {
+ case spv::Op::OpFunctionParameter:
+ case spv::Op::OpVariable:
// We found the memory object declaration.
// Remember it as the answer for the whole path.
for (auto iter : visited) {
memo_table[iter] = inst;
}
return inst;
- case SpvOpLoad:
+ case spv::Op::OpLoad:
// Follow the pointer being loaded
id = inst->GetSingleWordInOperand(0);
break;
- case SpvOpCopyObject:
+ case spv::Op::OpCopyObject:
// Follow the object being copied.
id = inst->GetSingleWordInOperand(0);
break;
- case SpvOpAccessChain:
- case SpvOpInBoundsAccessChain:
- case SpvOpPtrAccessChain:
- case SpvOpInBoundsPtrAccessChain:
+ case spv::Op::OpAccessChain:
+ case spv::Op::OpInBoundsAccessChain:
+ case spv::Op::OpPtrAccessChain:
+ case spv::Op::OpInBoundsPtrAccessChain:
// Follow the base pointer.
id = inst->GetSingleWordInOperand(0);
break;
- case SpvOpSampledImage:
+ case spv::Op::OpSampledImage:
// Follow the image or the sampler, depending on the follow_image
// parameter.
id = inst->GetSingleWordInOperand(follow_image ? 0 : 1);
break;
- case SpvOpImage:
+ case spv::Op::OpImage:
// Follow the sampled image
id = inst->GetSingleWordInOperand(0);
break;
@@ -2300,7 +2372,7 @@
// Get the SPIR-V handle type.
const auto* ptr_type = def_use_mgr_->GetDef(var.type_id());
- if (!ptr_type || (ptr_type->opcode() != SpvOpTypePointer)) {
+ if (!ptr_type || (opcode(ptr_type) != spv::Op::OpTypePointer)) {
Fail() << "Invalid type for variable or function parameter " << var.PrettyPrint();
return nullptr;
}
@@ -2309,18 +2381,18 @@
Fail() << "Invalid pointer type for variable or function parameter " << var.PrettyPrint();
return nullptr;
}
- switch (raw_handle_type->opcode()) {
- case SpvOpTypeSampler:
- case SpvOpTypeImage:
+ switch (opcode(raw_handle_type)) {
+ case spv::Op::OpTypeSampler:
+ case spv::Op::OpTypeImage:
// The expected cases.
break;
- case SpvOpTypeArray:
- case SpvOpTypeRuntimeArray:
+ case spv::Op::OpTypeArray:
+ case spv::Op::OpTypeRuntimeArray:
Fail() << "arrays of textures or samplers are not supported in WGSL; can't "
"translate variable or function parameter: "
<< var.PrettyPrint();
return nullptr;
- case SpvOpTypeSampledImage:
+ case spv::Op::OpTypeSampledImage:
Fail() << "WGSL does not support combined image-samplers: " << var.PrettyPrint();
return nullptr;
default:
@@ -2358,7 +2430,7 @@
// In SPIR-V you could statically reference a texture or sampler without
// using it in a way that gives us a clue on how to declare it. Look inside
// the store type to infer a usage.
- if (raw_handle_type->opcode() == SpvOpTypeSampler) {
+ if (opcode(raw_handle_type) == spv::Op::OpTypeSampler) {
usage.AddSampler();
} else {
// It's a texture.
@@ -2370,7 +2442,7 @@
const auto sampled_param = raw_handle_type->GetSingleWordInOperand(5);
const auto format_param = raw_handle_type->GetSingleWordInOperand(6);
// Only storage images have a format.
- if ((format_param != SpvImageFormatUnknown) ||
+ if ((format_param != uint32_t(spv::ImageFormat::Unknown)) ||
sampled_param == 2 /* without sampler */) {
// Get NonWritable and NonReadable attributes of the variable.
bool is_nonwritable = false;
@@ -2379,10 +2451,10 @@
if (deco.size() != 1) {
continue;
}
- if (deco[0] == SpvDecorationNonWritable) {
+ if (deco[0] == uint32_t(spv::Decoration::NonWritable)) {
is_nonwritable = true;
}
- if (deco[0] == SpvDecorationNonReadable) {
+ if (deco[0] == uint32_t(spv::Decoration::NonReadable)) {
is_nonreadable = true;
}
}
@@ -2429,9 +2501,9 @@
if (image_type->is_arrayed()) {
// Give a nicer error message here, where we have the offending variable
// in hand, rather than inside the enum converter.
- switch (image_type->dim()) {
- case SpvDim2D:
- case SpvDimCube:
+ switch (static_cast<spv::Dim>(image_type->dim())) {
+ case spv::Dim::Dim2D:
+ case spv::Dim::Cube:
break;
default:
Fail() << "WGSL arrayed textures must be 2d_array or cube_array: "
@@ -2450,7 +2522,7 @@
// WGSL textures are always formatted. Unformatted textures are always
// sampled.
if (usage.IsSampledTexture() || usage.IsStorageReadTexture() ||
- (image_type->format() == SpvImageFormatUnknown)) {
+ (uint32_t(image_type->format()) == uint32_t(spv::ImageFormat::Unknown))) {
// Make a sampled texture type.
auto* ast_sampled_component_type =
ConvertType(raw_handle_type->GetSingleWordInOperand(0));
@@ -2607,34 +2679,34 @@
for (const auto* f : topologically_ordered_functions_) {
for (const auto& bb : *f) {
for (const auto& inst : bb) {
- switch (inst.opcode()) {
+ switch (opcode(inst)) {
// Single texel reads and writes
- case SpvOpImageRead:
+ case spv::Op::OpImageRead:
handle_usage_[get_image(inst)].AddStorageReadTexture();
break;
- case SpvOpImageWrite:
+ case spv::Op::OpImageWrite:
handle_usage_[get_image(inst)].AddStorageWriteTexture();
break;
- case SpvOpImageFetch:
+ case spv::Op::OpImageFetch:
handle_usage_[get_image(inst)].AddSampledTexture();
break;
// Sampling and gathering from a sampled image.
- case SpvOpImageSampleImplicitLod:
- case SpvOpImageSampleExplicitLod:
- case SpvOpImageSampleProjImplicitLod:
- case SpvOpImageSampleProjExplicitLod:
- case SpvOpImageGather:
+ case spv::Op::OpImageSampleImplicitLod:
+ case spv::Op::OpImageSampleExplicitLod:
+ case spv::Op::OpImageSampleProjImplicitLod:
+ case spv::Op::OpImageSampleProjExplicitLod:
+ case spv::Op::OpImageGather:
handle_usage_[get_image(inst)].AddSampledTexture();
handle_usage_[get_sampler(inst)].AddSampler();
break;
- case SpvOpImageSampleDrefImplicitLod:
- case SpvOpImageSampleDrefExplicitLod:
- case SpvOpImageSampleProjDrefImplicitLod:
- case SpvOpImageSampleProjDrefExplicitLod:
- case SpvOpImageDrefGather:
+ case spv::Op::OpImageSampleDrefImplicitLod:
+ case spv::Op::OpImageSampleDrefExplicitLod:
+ case spv::Op::OpImageSampleProjDrefImplicitLod:
+ case spv::Op::OpImageSampleProjDrefExplicitLod:
+ case spv::Op::OpImageDrefGather:
// Depth reference access implies usage as a depth texture, which
// in turn is a sampled texture.
handle_usage_[get_image(inst)].AddDepthTexture();
@@ -2643,29 +2715,29 @@
// Image queries
- case SpvOpImageQuerySizeLod:
+ case spv::Op::OpImageQuerySizeLod:
// Vulkan requires Sampled=1 for this. SPIR-V already requires MS=0.
handle_usage_[get_image(inst)].AddSampledTexture();
break;
- case SpvOpImageQuerySize:
+ case spv::Op::OpImageQuerySize:
// Applies to either MS=1 or Sampled=0 or 2.
// So we can't force it to be multisampled, or storage image.
break;
- case SpvOpImageQueryLod:
+ case spv::Op::OpImageQueryLod:
handle_usage_[get_image(inst)].AddSampledTexture();
handle_usage_[get_sampler(inst)].AddSampler();
break;
- case SpvOpImageQueryLevels:
+ case spv::Op::OpImageQueryLevels:
// We can't tell anything more than that it's an image.
handle_usage_[get_image(inst)].AddTexture();
break;
- case SpvOpImageQuerySamples:
+ case spv::Op::OpImageQuerySamples:
handle_usage_[get_image(inst)].AddMultisampledTexture();
break;
// Function calls
- case SpvOpFunctionCall: {
+ case spv::Op::OpFunctionCall: {
// Propagate handle usages from callee function formal parameters to
// the matching caller parameters. This is where we rely on the
// fact that callees have been processed earlier in the flow.
diff --git a/src/tint/reader/spirv/parser_impl.h b/src/tint/reader/spirv/parser_impl.h
index 9965aa4..6e79ecc 100644
--- a/src/tint/reader/spirv/parser_impl.h
+++ b/src/tint/reader/spirv/parser_impl.h
@@ -66,8 +66,8 @@
/// The binary representation of a SPIR-V decoration enum followed by its
/// operands, if any.
-/// Example: { SpvDecorationBlock }
-/// Example: { SpvDecorationArrayStride, 16 }
+/// Example: { spv::Decoration::Block }
+/// Example: { spv::Decoration::ArrayStride, 16 }
using Decoration = std::vector<uint32_t>;
/// DecorationList is a list of decorations
@@ -338,13 +338,13 @@
/// Returns true when the given instruction is an extended instruction
/// for GLSL.std.450.
/// @param inst a SPIR-V instruction
- /// @returns true if its an SpvOpExtInst for GLSL.std.450
+ /// @returns true if its an spv::Op::ExtInst for GLSL.std.450
bool IsGlslExtendedInstruction(const spvtools::opt::Instruction& inst) const;
/// Returns true when the given instruction is an extended instruction
/// from an ignored extended instruction set.
/// @param inst a SPIR-V instruction
- /// @returns true if its an SpvOpExtInst for an ignored extended instruction
+ /// @returns true if its an spv::Op::ExtInst for an ignored extended instruction
bool IsIgnoredExtendedInstruction(const spvtools::opt::Instruction& inst) const;
/// Registers user names for SPIR-V objects, from OpName, and OpMemberName.
@@ -585,7 +585,7 @@
/// class class.
uint32_t pointer_type_id = 0;
/// The SPIR-V address space.
- SpvStorageClass storage_class = SpvStorageClassOutput;
+ spv::StorageClass storage_class = spv::StorageClass::Output;
/// The ID of the type of a pointer to the Position member.
uint32_t position_member_pointer_type_id = 0;
/// The ID of the gl_PerVertex variable, if it was declared.
@@ -690,7 +690,7 @@
const spvtools::opt::Instruction* GetInstructionForTest(uint32_t id) const;
/// A map of SPIR-V identifiers to builtins
- using BuiltInsMap = std::unordered_map<uint32_t, SpvBuiltIn>;
+ using BuiltInsMap = std::unordered_map<uint32_t, spv::BuiltIn>;
/// @returns a map of builtins that should be handled specially by code
/// generation. Either the builtin does not exist in WGSL, or a type
@@ -699,7 +699,7 @@
/// @param builtin the SPIR-V builtin variable kind
/// @returns the SPIR-V ID for the variable defining the given builtin, or 0
- uint32_t IdForSpecialBuiltIn(SpvBuiltIn builtin) const {
+ uint32_t IdForSpecialBuiltIn(spv::BuiltIn builtin) const {
// Do a linear search.
for (const auto& entry : special_builtins_) {
if (entry.second == builtin) {
diff --git a/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc b/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc
index 767b333..402cc41 100644
--- a/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc
+++ b/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc
@@ -31,7 +31,7 @@
TEST_F(SpvParserTest, ConvertMemberDecoration_OffsetWithoutOperand) {
auto p = parser(std::vector<uint32_t>{});
- auto result = p->ConvertMemberDecoration(12, 13, nullptr, {SpvDecorationOffset});
+ auto result = p->ConvertMemberDecoration(12, 13, nullptr, {uint32_t(spv::Decoration::Offset)});
EXPECT_TRUE(result.IsEmpty());
EXPECT_THAT(p->error(), Eq("malformed Offset decoration: expected 1 literal "
"operand, has 0: member 13 of SPIR-V type 12"));
@@ -40,7 +40,8 @@
TEST_F(SpvParserTest, ConvertMemberDecoration_OffsetWithTooManyOperands) {
auto p = parser(std::vector<uint32_t>{});
- auto result = p->ConvertMemberDecoration(12, 13, nullptr, {SpvDecorationOffset, 3, 4});
+ auto result =
+ p->ConvertMemberDecoration(12, 13, nullptr, {uint32_t(spv::Decoration::Offset), 3, 4});
EXPECT_TRUE(result.IsEmpty());
EXPECT_THAT(p->error(), Eq("malformed Offset decoration: expected 1 literal "
"operand, has 2: member 13 of SPIR-V type 12"));
@@ -49,7 +50,7 @@
TEST_F(SpvParserTest, ConvertMemberDecoration_Offset) {
auto p = parser(std::vector<uint32_t>{});
- auto result = p->ConvertMemberDecoration(1, 1, nullptr, {SpvDecorationOffset, 8});
+ auto result = p->ConvertMemberDecoration(1, 1, nullptr, {uint32_t(spv::Decoration::Offset), 8});
ASSERT_FALSE(result.IsEmpty());
EXPECT_TRUE(result[0]->Is<ast::StructMemberOffsetAttribute>());
auto* offset_deco = result[0]->As<ast::StructMemberOffsetAttribute>();
@@ -64,7 +65,8 @@
spirv::F32 f32;
spirv::Matrix matrix(&f32, 2, 2);
- auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 8});
+ auto result =
+ p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 8});
EXPECT_TRUE(result.IsEmpty());
EXPECT_TRUE(p->error().empty());
}
@@ -74,7 +76,8 @@
spirv::F32 f32;
spirv::Matrix matrix(&f32, 2, 2);
- auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 16});
+ auto result =
+ p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 16});
ASSERT_FALSE(result.IsEmpty());
EXPECT_TRUE(result[0]->Is<ast::StrideAttribute>());
auto* stride_deco = result[0]->As<ast::StrideAttribute>();
@@ -88,7 +91,8 @@
spirv::F32 f32;
spirv::Matrix matrix(&f32, 2, 4);
- auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 16});
+ auto result =
+ p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 16});
EXPECT_TRUE(result.IsEmpty());
EXPECT_TRUE(p->error().empty());
}
@@ -98,7 +102,8 @@
spirv::F32 f32;
spirv::Matrix matrix(&f32, 2, 4);
- auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 64});
+ auto result =
+ p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 64});
ASSERT_FALSE(result.IsEmpty());
EXPECT_TRUE(result[0]->Is<ast::StrideAttribute>());
auto* stride_deco = result[0]->As<ast::StrideAttribute>();
@@ -112,7 +117,8 @@
spirv::F32 f32;
spirv::Matrix matrix(&f32, 2, 3);
- auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 32});
+ auto result =
+ p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 32});
ASSERT_FALSE(result.IsEmpty());
EXPECT_TRUE(result[0]->Is<ast::StrideAttribute>());
auto* stride_deco = result[0]->As<ast::StrideAttribute>();
@@ -127,7 +133,8 @@
// relaxed precision f32.
auto p = parser(std::vector<uint32_t>{});
- auto result = p->ConvertMemberDecoration(1, 1, nullptr, {SpvDecorationRelaxedPrecision});
+ auto result =
+ p->ConvertMemberDecoration(1, 1, nullptr, {uint32_t(spv::Decoration::RelaxedPrecision)});
EXPECT_TRUE(result.IsEmpty());
EXPECT_TRUE(p->error().empty());
}
diff --git a/src/tint/reader/spirv/parser_impl_get_decorations_test.cc b/src/tint/reader/spirv/parser_impl_get_decorations_test.cc
index 929ab1f..cbc384f 100644
--- a/src/tint/reader/spirv/parser_impl_get_decorations_test.cc
+++ b/src/tint/reader/spirv/parser_impl_get_decorations_test.cc
@@ -52,7 +52,7 @@
)"));
EXPECT_TRUE(p->BuildAndParseInternalModule());
auto decorations = p->GetDecorationsFor(10);
- EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationBlock}));
+ EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::Block)}));
EXPECT_TRUE(p->error().empty());
p->SkipDumpingPending(kSkipReason);
}
@@ -66,7 +66,7 @@
)"));
EXPECT_TRUE(p->BuildAndParseInternalModule());
auto decorations = p->GetDecorationsFor(10);
- EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationBlock}));
+ EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::Block)}));
EXPECT_TRUE(p->error().empty());
p->SkipDumpingPending(kSkipReason);
}
@@ -80,8 +80,9 @@
)"));
EXPECT_TRUE(p->BuildAndParseInternalModule());
auto decorations = p->GetDecorationsFor(5);
- EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision},
- Decoration{SpvDecorationLocation, 7}));
+ EXPECT_THAT(decorations,
+ UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)},
+ Decoration{uint32_t(spv::Decoration::Location), 7}));
EXPECT_TRUE(p->error().empty());
p->SkipDumpingPending(kSkipReason);
}
@@ -124,7 +125,8 @@
)"));
EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
auto decorations = p->GetDecorationsForMember(10, 0);
- EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision}));
+ EXPECT_THAT(decorations,
+ UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)}));
EXPECT_TRUE(p->error().empty());
p->SkipDumpingPending(kSkipReason);
}
@@ -138,7 +140,8 @@
)"));
EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
auto decorations = p->GetDecorationsForMember(10, 0);
- EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision}));
+ EXPECT_THAT(decorations,
+ UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)}));
EXPECT_TRUE(p->error().empty());
p->SkipDumpingPending(kSkipReason);
}
@@ -154,7 +157,8 @@
)"));
EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
auto decorations = p->GetDecorationsForMember(10, 1);
- EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationArrayStride, 12}));
+ EXPECT_THAT(decorations,
+ UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::ArrayStride), 12}));
EXPECT_TRUE(p->error().empty());
}
@@ -179,11 +183,11 @@
EXPECT_TRUE(p->GetDecorationsForMember(50, 0).empty());
EXPECT_THAT(p->GetDecorationsForMember(50, 1),
- UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision}));
+ UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)}));
EXPECT_THAT(p->GetDecorationsForMember(50, 2),
- UnorderedElementsAre(Decoration{SpvDecorationColMajor},
- Decoration{SpvDecorationMatrixStride, 8},
- Decoration{SpvDecorationArrayStride, 16}));
+ UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::ColMajor)},
+ Decoration{uint32_t(spv::Decoration::MatrixStride), 8},
+ Decoration{uint32_t(spv::Decoration::ArrayStride), 16}));
EXPECT_TRUE(p->error().empty());
}
