Import Tint changes from Dawn
Changes:
- 4f3f45c00f414e6b57782126b4d98c8c44813d8f Tint/transform: Fix insertBits polyfill for HLSL by Zhaoming Jiang <zhaoming.jiang@intel.com>
- a70e365313d31757bc9b75002fa46b3366d551a8 tint: fix extractBits edge case by Antonio Maiorano <amaiorano@google.com>
- c36417343fe1c2a526dc83cc378d96a6db3ce5e5 tint: const eval of cross by Antonio Maiorano <amaiorano@google.com>
- eb949c87ee5aaa819faae6252dd36bfeda649070 spirv-reader: use spirv.hpp11 by David Neto <dneto@google.com>
- 1bdaded736d61e4bdbf18a90c8ebac9f3274cb2f Implemement const-eval for some unpack routines. by dan sinclair <dsinclair@chromium.org>
- ec4b650adb133672f48ed7d76c12459933cb229c Remove fallthrough from fuzzer AST mutation. by dan sinclair <dsinclair@chromium.org>
- ba200e87c5db20d67e3c0d51bfb4ddf3151766fc tint: Bump stack size for MSVC + debug build by Ben Clayton <bclayton@google.com>
- 5071a54af491e05444c02122b48e9cd5d65f82c7 Make reserved words an error. by dan sinclair <dsinclair@chromium.org>
- 2d706a04360dc35409a2df324b936f6735eb1cd8 Implemement const-eval for some pack routines. by dan sinclair <dsinclair@chromium.org>
- 3b2b5484e239fc558871c04eb061c201d0c15c36 Remove module-scope let deprecation. by dan sinclair <dsinclair@chromium.org>
- c027f33cfd88efd155e88623300dc3356fcd6b95 tint: Fix gdb and lldb pretty printers for recently updat... by Antonio Maiorano <amaiorano@google.com>
- ca98b1b1b9d797bf1fc65340036193a2e7a205b1 tint/resolver: Resolve dependencies of parameter attributes by Ben Clayton <bclayton@google.com>
- 6f799d676d8cdec69153cc6b5ca6003598baf19f Change default acos test value. by dan sinclair <dsinclair@chromium.org>
- f218af04610a66598d44ecbba6c2b17bb1e78433 Remove the `@stage` attribute by dan sinclair <dsinclair@chromium.org>
GitOrigin-RevId: 4f3f45c00f414e6b57782126b4d98c8c44813d8f
Change-Id: Ica3b6cfbfc9b29a0d28600b6f3116001246ad3a2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/109100
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 337b6a1..39fae60 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -581,8 +581,8 @@
"utils/enum_set.h",
"utils/foreach_macro.h",
"utils/hash.h",
- "utils/hashmap_base.h",
"utils/hashmap.h",
+ "utils/hashmap_base.h",
"utils/hashset.h",
"utils/map.h",
"utils/math.h",
@@ -1423,7 +1423,6 @@
"reader/wgsl/parser_impl_multiplicative_expression_test.cc",
"reader/wgsl/parser_impl_param_list_test.cc",
"reader/wgsl/parser_impl_paren_expression_test.cc",
- "reader/wgsl/parser_impl_pipeline_stage_test.cc",
"reader/wgsl/parser_impl_primary_expression_test.cc",
"reader/wgsl/parser_impl_relational_expression_test.cc",
"reader/wgsl/parser_impl_reserved_keyword_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index a22d9fc..1d39243 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -752,6 +752,7 @@
endif()
endif()
+
################################################################################
# Tests
################################################################################
@@ -1054,7 +1055,6 @@
reader/wgsl/parser_impl_multiplicative_expression_test.cc
reader/wgsl/parser_impl_param_list_test.cc
reader/wgsl/parser_impl_paren_expression_test.cc
- reader/wgsl/parser_impl_pipeline_stage_test.cc
reader/wgsl/parser_impl_primary_expression_test.cc
reader/wgsl/parser_impl_relational_expression_test.cc
reader/wgsl/parser_impl_reserved_keyword_test.cc
@@ -1352,7 +1352,16 @@
add_executable(tint_unittests ${TINT_TEST_SRCS})
set_target_properties(${target} PROPERTIES FOLDER "Tests")
- if(NOT MSVC)
+ if(MSVC)
+ # TODO(crbug.com/tint/1749): MSVC debug builds can suffer from stack
+ # overflows when resolving deeply nested expression chains or statements.
+ # Production builds neither use MSVC nor debug, so just bump the stack size
+ # for this build combination.
+ string(TOUPPER "${CMAKE_BUILD_TYPE}" build_type)
+ if ((NOT ${build_type} STREQUAL "RELEASE") AND (NOT ${build_type} STREQUAL "RELWITHDEBINFO"))
+ target_link_options(tint_unittests PRIVATE "/STACK 2097152") # 2MB, default is 1MB
+ endif()
+ else()
target_compile_options(tint_unittests PRIVATE
-Wno-global-constructors
-Wno-weak-vtables
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc
index 05435d0..1ec3ca4 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc
@@ -19,7 +19,6 @@
#include <vector>
#include "src/tint/ast/block_statement.h"
-#include "src/tint/ast/fallthrough_statement.h"
#include "src/tint/ast/for_loop_statement.h"
#include "src/tint/ast/if_statement.h"
#include "src/tint/ast/loop_statement.h"
@@ -151,24 +150,9 @@
}
if (auto* case_statement = statement_node.As<ast::CaseStatement>()) {
- // It is not OK to delete the final case statement in a switch statement if the penultimate
- // case statement falls through to the final case statement.
- auto* switch_statement =
- program.Sem().Get(case_statement)->Parent()->Declaration()->As<ast::SwitchStatement>();
-
- if (switch_statement->body.Length() > 1 &&
- switch_statement->body[switch_statement->body.Length() - 1] == case_statement) {
- // There are at least two case statements, and this is the final case statement.
- auto& penultimate_case_statement_body_statements =
- switch_statement->body[switch_statement->body.Length() - 2]->body->statements;
- if (penultimate_case_statement_body_statements.Length() > 0 &&
- penultimate_case_statement_body_statements
- [penultimate_case_statement_body_statements.Length() - 1]
- ->Is<ast::FallthroughStatement>()) {
- // The penultimate case statement falls through to the final case statement, thus
- // the final case statement cannot be removed.
- return false;
- }
+ // It is not OK to delete the case statement which contains the default selector.
+ if (case_statement->ContainsDefault()) {
+ return false;
}
}
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc
index 925ee80..275a19e 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc
@@ -22,7 +22,6 @@
#include "src/tint/ast/assignment_statement.h"
#include "src/tint/ast/block_statement.h"
#include "src/tint/ast/case_statement.h"
-#include "src/tint/ast/fallthrough_statement.h"
#include "src/tint/ast/for_loop_statement.h"
#include "src/tint/ast/if_statement.h"
#include "src/tint/ast/switch_statement.h"
@@ -151,10 +150,7 @@
switch(1) {
case 0, 1: {
}
- default: {
- fallthrough;
- }
- case 2: {
+ case 2, default: {
}
}
})";
@@ -171,20 +167,14 @@
switch(1) {
case 0, 1: {
}
- default: {
- fallthrough;
- }
- case 2: {
+ case 2, default: {
}
}
})";
auto expected = R"(
fn main() {
switch(1) {
- default: {
- fallthrough;
- }
- case 2: {
+ case 2, default: {
}
}
})";
@@ -199,43 +189,6 @@
CheckStatementDeletionWorks(original, expected, statement_finder);
}
-TEST(DeleteStatementTest, DeleteFallthroughStatement) {
- auto original = R"(
-fn main() {
- switch(1) {
- case 0, 1: {
- }
- default: {
- fallthrough;
- }
- case 2: {
- }
- }
-})";
- auto expected = R"(
-fn main() {
- switch(1) {
- case 0, 1: {
- }
- default: {
- }
- case 2: {
- }
- }
-})";
- auto statement_finder = [](const Program& program) -> const ast::Statement* {
- return program.AST()
- .Functions()[0]
- ->body->statements[0]
- ->As<ast::SwitchStatement>()
- ->body[1]
- ->As<ast::CaseStatement>()
- ->body->statements[0]
- ->As<ast::FallthroughStatement>();
- };
- CheckStatementDeletionWorks(original, expected, statement_finder);
-}
-
TEST(DeleteStatementTest, DeleteElse) {
auto original = R"(
fn main() {
@@ -570,14 +523,11 @@
CheckStatementDeletionNotAllowed(original, statement_finder);
}
-TEST(DeleteStatementTest, DoNotDeleteCaseDueToFallthrough) {
+TEST(DeleteStatementTest, DoNotDeleteCaseDueToDefault) {
auto original = R"(
fn main() {
switch(1) {
- default: {
- fallthrough;
- }
- case 2: {
+ case 2, default: {
}
}
})";
@@ -586,7 +536,7 @@
.Functions()[0]
->body->statements[0]
->As<ast::SwitchStatement>()
- ->body[1]
+ ->body[0]
->As<ast::CaseStatement>();
};
CheckStatementDeletionNotAllowed(original, statement_finder);
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
index ac5743e..8339ac0 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
@@ -386,7 +386,7 @@
}
var<private> a: S;
-let e = 3;
+const e = 3;
@group(1) @binding(1) var<uniform> b: S;
fn f() {
*&a = *&b;
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index f497405..96f6d08 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -406,8 +406,8 @@
// https://gpuweb.github.io/gpuweb/wgsl/#builtin-functions
@const fn abs<T: fia_fiu32_f16>(T) -> T
@const fn abs<N: num, T: fia_fiu32_f16>(vec<N, T>) -> vec<N, T>
-@const fn acos<T: fa_f32_f16>(@test_value(0.87758256189) T) -> T
-@const fn acos<N: num, T: fa_f32_f16>(@test_value(0.87758256189) vec<N, T>) -> vec<N, T>
+@const fn acos<T: fa_f32_f16>(@test_value(0.96891242171) T) -> T
+@const fn acos<N: num, T: fa_f32_f16>(@test_value(0.96891242171) vec<N, T>) -> vec<N, T>
fn acosh<T: f32_f16>(T) -> T
fn acosh<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
@const fn all(bool) -> bool
@@ -439,7 +439,7 @@
@const fn countOneBits<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
@const fn countTrailingZeros<T: iu32>(T) -> T
@const fn countTrailingZeros<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
-fn cross<T: f32_f16>(vec3<T>, vec3<T>) -> vec3<T>
+@const fn cross<T: fa_f32_f16>(vec3<T>, vec3<T>) -> vec3<T>
fn degrees<T: f32_f16>(T) -> T
fn degrees<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
fn determinant<N: num, T: f32_f16>(mat<N, N, T>) -> T
@@ -508,10 +508,10 @@
fn modf<N: num, T: f32_f16>(vec<N, T>) -> __modf_result_vec<N, T>
fn normalize<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
fn pack2x16float(vec2<f32>) -> u32
-fn pack2x16snorm(vec2<f32>) -> u32
-fn pack2x16unorm(vec2<f32>) -> u32
-fn pack4x8snorm(vec4<f32>) -> u32
-fn pack4x8unorm(vec4<f32>) -> u32
+@const fn pack2x16snorm(vec2<f32>) -> u32
+@const fn pack2x16unorm(vec2<f32>) -> u32
+@const fn pack4x8snorm(vec4<f32>) -> u32
+@const fn pack4x8unorm(vec4<f32>) -> u32
fn pow<T: f32_f16>(T, T) -> T
fn pow<N: num, T: f32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
@const fn quantizeToF16(f32) -> f32
@@ -550,10 +550,10 @@
fn trunc<T: f32_f16>(T) -> T
fn trunc<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
fn unpack2x16float(u32) -> vec2<f32>
-fn unpack2x16snorm(u32) -> vec2<f32>
-fn unpack2x16unorm(u32) -> vec2<f32>
-fn unpack4x8snorm(u32) -> vec4<f32>
-fn unpack4x8unorm(u32) -> vec4<f32>
+@const fn unpack2x16snorm(u32) -> vec2<f32>
+@const fn unpack2x16unorm(u32) -> vec2<f32>
+@const fn unpack4x8snorm(u32) -> vec4<f32>
+@const fn unpack4x8unorm(u32) -> vec4<f32>
@stage("compute") fn workgroupBarrier()
fn textureDimensions<T: fiu32>(texture: texture_1d<T>) -> u32
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());
}
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index 976706a..398c123 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -615,12 +615,11 @@
Source source;
if (match(Token::Type::kConst)) {
use = "'const' declaration";
- } else if (match(Token::Type::kLet, &source)) {
- use = "'let' declaration";
- deprecated(source, "module-scope 'let' has been replaced with 'const'");
} else if (match(Token::Type::kOverride)) {
use = "'override' declaration";
is_overridable = true;
+ } else if (match(Token::Type::kLet, &source)) {
+ return add_error(source, "module-scope 'let' is invalid, use 'const'");
} else {
return Failure::kNoMatch;
}
@@ -1613,29 +1612,6 @@
std::move(attrs.value)); // attributes
}
-// pipeline_stage
-// : VERTEX
-// | FRAGMENT
-// | COMPUTE
-//
-// TODO(crbug.com/tint/1503): Remove when deprecation period is over.
-Expect<ast::PipelineStage> ParserImpl::expect_pipeline_stage() {
- auto& t = peek();
- if (t == "vertex") {
- next(); // Consume the peek
- return {ast::PipelineStage::kVertex, t.source()};
- }
- if (t == "fragment") {
- next(); // Consume the peek
- return {ast::PipelineStage::kFragment, t.source()};
- }
- if (t == "compute") {
- next(); // Consume the peek
- return {ast::PipelineStage::kCompute, t.source()};
- }
- return add_error(peek(), "invalid value for stage attribute");
-}
-
// interpolation_sample_name
// : 'center'
// | 'centroid'
@@ -3637,34 +3613,6 @@
});
}
- // TODO(crbug.com/tint/1503): Remove when deprecation period is over.
- if (t == "stage") {
- return expect_paren_block("stage attribute", [&]() -> Result {
- auto stage = expect_pipeline_stage();
- if (stage.errored) {
- return Failure::kErrored;
- }
-
- std::string warning = "remove stage and use @";
- switch (stage.value) {
- case ast::PipelineStage::kVertex:
- warning += "vertex";
- break;
- case ast::PipelineStage::kFragment:
- warning += "fragment";
- break;
- case ast::PipelineStage::kCompute:
- warning += "compute";
- break;
- case ast::PipelineStage::kNone:
- break;
- }
- deprecated(t.source(), warning);
-
- return create<ast::StageAttribute>(t.source(), stage.value);
- });
- }
-
if (t == "vertex") {
return create<ast::StageAttribute>(t.source(), ast::PipelineStage::kVertex);
}
@@ -3826,7 +3774,7 @@
next();
if (is_reserved(t)) {
- deprecated(t.source(), "'" + t.to_str() + "' is a reserved keyword");
+ return add_error(t.source(), "'" + t.to_str() + "' is a reserved keyword");
}
return {t.to_str(), t.source()};
diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
index fbfe6f6..6c4f14a 100644
--- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
@@ -351,48 +351,6 @@
)");
}
-// TODO(crbug.com/tint/1503): Remove this when @stage is removed
-TEST_F(ParserImplErrorTest, FunctionDeclStageMissingLParen) {
- EXPECT("@stage vertex) fn f() {}",
- R"(test.wgsl:1:8 error: expected '(' for stage attribute
-@stage vertex) fn f() {}
- ^^^^^^
-)");
-}
-
-TEST_F(ParserImplErrorTest, FunctionDeclStageMissingRParen) {
- EXPECT(
- "@stage(vertex fn f() {}",
- R"(test.wgsl:1:2 warning: use of deprecated language feature: remove stage and use @vertex
-@stage(vertex fn f() {}
- ^^^^^
-
-test.wgsl:1:15 error: expected ')' for stage attribute
-@stage(vertex fn f() {}
- ^^
-)");
-}
-
-TEST_F(ParserImplErrorTest, FunctionDeclStageInvalid) {
- EXPECT("@stage(x) fn f() {}",
- R"(test.wgsl:1:8 error: invalid value for stage attribute
-@stage(x) fn f() {}
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, FunctionDeclStageTypeInvalid) {
- EXPECT("@shader(vertex) fn main() {}",
- R"(test.wgsl:1:2 error: expected attribute
-@shader(vertex) fn main() {}
- ^^^^^^
-
-test.wgsl:1:8 error: unexpected token
-@shader(vertex) fn main() {}
- ^
-)");
-}
-
TEST_F(ParserImplErrorTest, FunctionDeclWorkgroupSizeXInvalid) {
EXPECT("@workgroup_size() fn f() {}",
R"(test.wgsl:1:17 error: expected workgroup_size x parameter
@@ -595,94 +553,11 @@
)");
}
-TEST_F(ParserImplErrorTest, GlobalDeclLetInvalidIdentifier) {
- EXPECT(
- "let ^ : i32 = 1;",
- R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let ^ : i32 = 1;
+TEST_F(ParserImplErrorTest, GlobalDeclLet) {
+ EXPECT("let a : i32 = 1;",
+ R"(test.wgsl:1:1 error: module-scope 'let' is invalid, use 'const'
+let a : i32 = 1;
^^^
-
-test.wgsl:1:5 error: expected identifier for 'let' declaration
-let ^ : i32 = 1;
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetMissingSemicolon) {
- EXPECT(
- "let i : i32 = 1",
- R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : i32 = 1
-^^^
-
-test.wgsl:1:16 error: expected ';' for 'const' declaration
-let i : i32 = 1
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetMissingLParen) {
- EXPECT(
- "let i : vec2<i32> = vec2<i32>;",
- R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>;
-^^^
-
-test.wgsl:1:30 error: expected '(' for type initializer
-let i : vec2<i32> = vec2<i32>;
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetMissingRParen) {
- EXPECT(
- "let i : vec2<i32> = vec2<i32>(1., 2.;",
- R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>(1., 2.;
-^^^
-
-test.wgsl:1:37 error: expected ')' for type initializer
-let i : vec2<i32> = vec2<i32>(1., 2.;
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetBadConstLiteral) {
- EXPECT(
- "let i : vec2<i32> = vec2<i32>(!);",
- R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>(!);
-^^^
-
-test.wgsl:1:32 error: unable to parse right side of ! expression
-let i : vec2<i32> = vec2<i32>(!);
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetExprMissingLParen) {
- EXPECT(
- "let i : vec2<i32> = vec2<i32> 1, 2);",
- R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32> 1, 2);
-^^^
-
-test.wgsl:1:31 error: expected '(' for type initializer
-let i : vec2<i32> = vec2<i32> 1, 2);
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetExprMissingRParen) {
- EXPECT(
- "let i : vec2<i32> = vec2<i32>(1, 2;",
- R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>(1, 2;
-^^^
-
-test.wgsl:1:35 error: expected ')' for type initializer
-let i : vec2<i32> = vec2<i32>(1, 2;
- ^
)");
}
diff --git a/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc b/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
index 28b84a9..785d23c 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
@@ -18,33 +18,6 @@
namespace tint::reader::wgsl {
namespace {
-// TODO(crbug.com/tint/1503): Remove this when @stage is removed
-TEST_F(ParserImplTest, AttributeList_Parses_Stage) {
- auto p = parser("@workgroup_size(2) @stage(compute)");
- auto attrs = p->attribute_list();
- EXPECT_FALSE(p->has_error()) << p->error();
- EXPECT_FALSE(attrs.errored);
- EXPECT_TRUE(attrs.matched);
- ASSERT_EQ(attrs.value.Length(), 2u);
-
- auto* attr_0 = attrs.value[0]->As<ast::Attribute>();
- auto* attr_1 = attrs.value[1]->As<ast::Attribute>();
- ASSERT_NE(attr_0, nullptr);
- ASSERT_NE(attr_1, nullptr);
-
- ASSERT_TRUE(attr_0->Is<ast::WorkgroupAttribute>());
- const ast::Expression* x = attr_0->As<ast::WorkgroupAttribute>()->x;
- ASSERT_NE(x, nullptr);
- auto* x_literal = x->As<ast::LiteralExpression>();
- ASSERT_NE(x_literal, nullptr);
- ASSERT_TRUE(x_literal->Is<ast::IntLiteralExpression>());
- EXPECT_EQ(x_literal->As<ast::IntLiteralExpression>()->value, 2);
- EXPECT_EQ(x_literal->As<ast::IntLiteralExpression>()->suffix,
- ast::IntLiteralExpression::Suffix::kNone);
- ASSERT_TRUE(attr_1->Is<ast::StageAttribute>());
- EXPECT_EQ(attr_1->As<ast::StageAttribute>()->stage, ast::PipelineStage::kCompute);
-}
-
TEST_F(ParserImplTest, AttributeList_Parses) {
auto p = parser("@workgroup_size(2) @compute");
auto attrs = p->attribute_list();
@@ -81,14 +54,5 @@
EXPECT_EQ(p->error(), "1:2: expected attribute");
}
-TEST_F(ParserImplTest, AttributeList_BadAttribute) {
- auto p = parser("@stage()");
- auto attrs = p->attribute_list();
- EXPECT_TRUE(p->has_error());
- EXPECT_TRUE(attrs.errored);
- EXPECT_FALSE(attrs.matched);
- EXPECT_EQ(p->error(), "1:8: invalid value for stage attribute");
-}
-
} // namespace
} // namespace tint::reader::wgsl
diff --git a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
index 60e8f86..186c86a 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
@@ -397,61 +397,6 @@
EXPECT_EQ(p->error(), "1:21: expected ')' for workgroup_size attribute");
}
-// TODO(crbug.com/tint/1503): Remove when @stage is removed
-TEST_F(ParserImplTest, Attribute_Stage) {
- auto p = parser("stage(compute)");
- auto attr = p->attribute();
- EXPECT_TRUE(attr.matched);
- EXPECT_FALSE(attr.errored);
- ASSERT_NE(attr.value, nullptr) << p->error();
- ASSERT_FALSE(p->has_error());
- auto* func_attr = attr.value->As<ast::Attribute>();
- ASSERT_NE(func_attr, nullptr);
- ASSERT_TRUE(func_attr->Is<ast::StageAttribute>());
- EXPECT_EQ(func_attr->As<ast::StageAttribute>()->stage, ast::PipelineStage::kCompute);
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingValue) {
- auto p = parser("stage()");
- auto attr = p->attribute();
- EXPECT_FALSE(attr.matched);
- EXPECT_TRUE(attr.errored);
- EXPECT_EQ(attr.value, nullptr);
- EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:7: invalid value for stage attribute");
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingInvalid) {
- auto p = parser("stage(nan)");
- auto attr = p->attribute();
- EXPECT_FALSE(attr.matched);
- EXPECT_TRUE(attr.errored);
- EXPECT_EQ(attr.value, nullptr);
- EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:7: invalid value for stage attribute");
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingLeftParen) {
- auto p = parser("stage compute)");
- auto attr = p->attribute();
- EXPECT_FALSE(attr.matched);
- EXPECT_TRUE(attr.errored);
- EXPECT_EQ(attr.value, nullptr);
- EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:7: expected '(' for stage attribute");
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingRightParen) {
- auto p = parser("stage(compute");
- auto attr = p->attribute();
- EXPECT_FALSE(attr.matched);
- EXPECT_TRUE(attr.errored);
- EXPECT_EQ(attr.value, nullptr);
- EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), R"(1:1: use of deprecated language feature: remove stage and use @compute
-1:14: expected ')' for stage attribute)");
-}
-
TEST_F(ParserImplTest, Attribute_Compute) {
auto p = parser("compute");
auto attr = p->attribute();
diff --git a/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
index 86e4383..c465e4b 100644
--- a/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
@@ -24,79 +24,10 @@
EXPECT_FALSE(attrs.errored);
EXPECT_FALSE(attrs.matched);
auto e = p->global_constant_decl(attrs.value);
- EXPECT_FALSE(p->has_error()) << p->error();
- EXPECT_TRUE(e.matched);
- EXPECT_FALSE(e.errored);
- auto* const_ = e.value->As<ast::Const>();
- ASSERT_NE(const_, nullptr);
-
- EXPECT_EQ(const_->symbol, p->builder().Symbols().Get("a"));
- ASSERT_NE(const_->type, nullptr);
- EXPECT_TRUE(const_->type->Is<ast::F32>());
-
- EXPECT_EQ(const_->source.range.begin.line, 1u);
- EXPECT_EQ(const_->source.range.begin.column, 5u);
- EXPECT_EQ(const_->source.range.end.line, 1u);
- EXPECT_EQ(const_->source.range.end.column, 6u);
-
- ASSERT_NE(const_->initializer, nullptr);
- EXPECT_TRUE(const_->initializer->Is<ast::LiteralExpression>());
-}
-
-TEST_F(ParserImplTest, GlobalLetDecl_Inferred) {
- auto p = parser("let a = 1.");
- auto attrs = p->attribute_list();
- EXPECT_FALSE(attrs.errored);
- EXPECT_FALSE(attrs.matched);
- auto e = p->global_constant_decl(attrs.value);
- EXPECT_FALSE(p->has_error()) << p->error();
- EXPECT_TRUE(e.matched);
- EXPECT_FALSE(e.errored);
- auto* const_ = e.value->As<ast::Const>();
- ASSERT_NE(const_, nullptr);
-
- EXPECT_EQ(const_->symbol, p->builder().Symbols().Get("a"));
- EXPECT_EQ(const_->type, nullptr);
-
- EXPECT_EQ(const_->source.range.begin.line, 1u);
- EXPECT_EQ(const_->source.range.begin.column, 5u);
- EXPECT_EQ(const_->source.range.end.line, 1u);
- EXPECT_EQ(const_->source.range.end.column, 6u);
-
- ASSERT_NE(const_->initializer, nullptr);
- EXPECT_TRUE(const_->initializer->Is<ast::LiteralExpression>());
-}
-
-TEST_F(ParserImplTest, GlobalLetDecl_InvalidExpression) {
- auto p = parser("let a : f32 = if (a) {}");
- auto attrs = p->attribute_list();
- EXPECT_FALSE(attrs.errored);
- EXPECT_FALSE(attrs.matched);
- auto e = p->global_constant_decl(attrs.value);
EXPECT_TRUE(p->has_error());
- EXPECT_TRUE(e.errored);
EXPECT_FALSE(e.matched);
- EXPECT_EQ(e.value, nullptr);
- EXPECT_EQ(
- p->error(),
- R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:15: missing initializer for 'let' declaration)");
-}
-
-TEST_F(ParserImplTest, GlobalLetDecl_MissingExpression) {
- auto p = parser("let a : f32 =");
- auto attrs = p->attribute_list();
- EXPECT_FALSE(attrs.errored);
- EXPECT_FALSE(attrs.matched);
- auto e = p->global_constant_decl(attrs.value);
- EXPECT_TRUE(p->has_error());
EXPECT_TRUE(e.errored);
- EXPECT_FALSE(e.matched);
- EXPECT_EQ(e.value, nullptr);
- EXPECT_EQ(
- p->error(),
- R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:14: missing initializer for 'let' declaration)");
+ EXPECT_EQ(p->error(), "1:1: module-scope 'let' is invalid, use 'const'");
}
TEST_F(ParserImplTest, GlobalConstDecl) {
diff --git a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
index 6c943c8..2f124d5 100644
--- a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
@@ -58,44 +58,11 @@
TEST_F(ParserImplTest, GlobalDecl_GlobalLet) {
auto p = parser("let a : i32 = 2;");
- p->global_decl();
- ASSERT_FALSE(p->has_error()) << p->error();
-
- auto program = p->program();
- ASSERT_EQ(program.AST().GlobalVariables().Length(), 1u);
-
- auto* v = program.AST().GlobalVariables()[0];
- EXPECT_EQ(v->symbol, program.Symbols().Get("a"));
-}
-
-TEST_F(ParserImplTest, GlobalDecl_GlobalLet_MissingInitializer) {
- auto p = parser("let a : vec2<i32>;");
- p->global_decl();
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(
- p->error(),
- R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:18: expected '=' for 'let' declaration)");
-}
-
-TEST_F(ParserImplTest, GlobalDecl_GlobalLet_Invalid) {
- auto p = parser("let a : vec2<i32> 1.0;");
- p->global_decl();
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(
- p->error(),
- R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:19: expected '=' for 'let' declaration)");
-}
-
-TEST_F(ParserImplTest, GlobalDecl_GlobalLet_MissingSemicolon) {
- auto p = parser("let a : vec2<i32> = vec2<i32>(1, 2)");
- p->global_decl();
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(
- p->error(),
- R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:36: expected ';' for 'const' declaration)");
+ auto e = p->global_decl();
+ EXPECT_TRUE(p->has_error());
+ EXPECT_FALSE(e.matched);
+ EXPECT_TRUE(e.errored);
+ EXPECT_EQ(p->error(), "1:1: module-scope 'let' is invalid, use 'const'");
}
TEST_F(ParserImplTest, GlobalDecl_GlobalConst) {
diff --git a/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc b/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc
deleted file mode 100644
index 3c730e2..0000000
--- a/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc
+++ /dev/null
@@ -1,62 +0,0 @@
-// Copyright 2020 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/reader/wgsl/parser_impl_test_helper.h"
-
-namespace tint::reader::wgsl {
-namespace {
-
-struct PipelineStageData {
- std::string input;
- ast::PipelineStage result;
-};
-inline std::ostream& operator<<(std::ostream& out, PipelineStageData data) {
- return out << data.input;
-}
-
-class PipelineStageTest : public ParserImplTestWithParam<PipelineStageData> {};
-
-TEST_P(PipelineStageTest, Parses) {
- auto params = GetParam();
- auto p = parser(params.input);
-
- auto stage = p->expect_pipeline_stage();
- ASSERT_FALSE(p->has_error()) << p->error();
- ASSERT_FALSE(stage.errored);
- EXPECT_EQ(stage.value, params.result);
- EXPECT_EQ(stage.source.range.begin.line, 1u);
- EXPECT_EQ(stage.source.range.begin.column, 1u);
- EXPECT_EQ(stage.source.range.end.line, 1u);
- EXPECT_EQ(stage.source.range.end.column, 1u + params.input.size());
-
- auto& t = p->next();
- EXPECT_TRUE(t.IsEof());
-}
-INSTANTIATE_TEST_SUITE_P(
- ParserImplTest,
- PipelineStageTest,
- testing::Values(PipelineStageData{"vertex", ast::PipelineStage::kVertex},
- PipelineStageData{"fragment", ast::PipelineStage::kFragment},
- PipelineStageData{"compute", ast::PipelineStage::kCompute}));
-
-TEST_F(ParserImplTest, PipelineStage_NoMatch) {
- auto p = parser("not-a-stage");
- auto stage = p->expect_pipeline_stage();
- ASSERT_TRUE(p->has_error());
- ASSERT_TRUE(stage.errored);
- ASSERT_EQ(p->error(), "1:1: invalid value for stage attribute");
-}
-
-} // namespace
-} // namespace tint::reader::wgsl
diff --git a/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc b/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc
index 0ad3cc1..8e178ab 100644
--- a/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc
@@ -21,74 +21,67 @@
TEST_P(ParserImplReservedKeywordTest, Function) {
auto name = GetParam();
auto p = parser("fn " + name + "() {}");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:4: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:4: '" + name + "' is a reserved keyword");
}
TEST_P(ParserImplReservedKeywordTest, ModuleConst) {
auto name = GetParam();
auto p = parser("const " + name + " : i32 = 1;");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:7: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:7: '" + name + "' is a reserved keyword");
}
TEST_P(ParserImplReservedKeywordTest, ModuleVar) {
auto name = GetParam();
auto p = parser("var " + name + " : i32 = 1;");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:5: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:5: '" + name + "' is a reserved keyword");
}
TEST_P(ParserImplReservedKeywordTest, FunctionLet) {
auto name = GetParam();
auto p = parser("fn f() { let " + name + " : i32 = 1; }");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:14: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:14: '" + name + "' is a reserved keyword");
}
TEST_P(ParserImplReservedKeywordTest, FunctionVar) {
auto name = GetParam();
auto p = parser("fn f() { var " + name + " : i32 = 1; }");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:14: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:14: '" + name + "' is a reserved keyword");
}
TEST_P(ParserImplReservedKeywordTest, FunctionParam) {
auto name = GetParam();
auto p = parser("fn f(" + name + " : i32) {}");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:6: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:6: '" + name + "' is a reserved keyword");
}
TEST_P(ParserImplReservedKeywordTest, Struct) {
auto name = GetParam();
auto p = parser("struct " + name + " {};");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:8: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), R"(1:8: ')" + name + R"(' is a reserved keyword
+1:)" + std::to_string(9 + name.length()) +
+ R"(: statement found outside of function body)");
}
TEST_P(ParserImplReservedKeywordTest, StructMember) {
auto name = GetParam();
auto p = parser("struct S { " + name + " : i32, };");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:12: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:12: '" + name + "' is a reserved keyword");
}
TEST_P(ParserImplReservedKeywordTest, Alias) {
auto name = GetParam();
auto p = parser("type " + name + " = i32;");
- EXPECT_TRUE(p->Parse());
- EXPECT_FALSE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:6: use of deprecated language feature: '" + name + "' is a reserved keyword");
+ EXPECT_FALSE(p->Parse());
+ EXPECT_TRUE(p->has_error());
+ EXPECT_EQ(p->error(), "1:6: '" + name + "' is a reserved keyword");
}
INSTANTIATE_TEST_SUITE_P(ParserImplReservedKeywordTest,
ParserImplReservedKeywordTest,
diff --git a/src/tint/resolver/builtin_test.cc b/src/tint/resolver/builtin_test.cc
index c3911fb..3d8a4db 100644
--- a/src/tint/resolver/builtin_test.cc
+++ b/src/tint/resolver/builtin_test.cc
@@ -748,7 +748,7 @@
EXPECT_EQ(r()->error(), R"(error: no matching call to cross()
1 candidate function:
- cross(vec3<T>, vec3<T>) -> vec3<T> where: T is f32 or f16
+ cross(vec3<T>, vec3<T>) -> vec3<T> where: T is abstract-float, f32 or f16
)");
}
@@ -761,7 +761,7 @@
EXPECT_EQ(r()->error(), R"(error: no matching call to cross(f32, f32)
1 candidate function:
- cross(vec3<T>, vec3<T>) -> vec3<T> where: T is f32 or f16
+ cross(vec3<T>, vec3<T>) -> vec3<T> where: T is abstract-float, f32 or f16
)");
}
@@ -775,7 +775,7 @@
R"(error: no matching call to cross(vec3<i32>, vec3<i32>)
1 candidate function:
- cross(vec3<T>, vec3<T>) -> vec3<T> where: T is f32 or f16
+ cross(vec3<T>, vec3<T>) -> vec3<T> where: T is abstract-float, f32 or f16
)");
}
@@ -790,7 +790,7 @@
R"(error: no matching call to cross(vec4<f32>, vec4<f32>)
1 candidate function:
- cross(vec3<T>, vec3<T>) -> vec3<T> where: T is f32 or f16
+ cross(vec3<T>, vec3<T>) -> vec3<T> where: T is abstract-float, f32 or f16
)");
}
@@ -806,7 +806,7 @@
R"(error: no matching call to cross(vec3<f32>, vec3<f32>, vec3<f32>)
1 candidate function:
- cross(vec3<T>, vec3<T>) -> vec3<T> where: T is f32 or f16
+ cross(vec3<T>, vec3<T>) -> vec3<T> where: T is abstract-float, f32 or f16
)");
}
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index afb9f5d..2c012a1 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -37,6 +37,7 @@
#include "src/tint/sem/type_initializer.h"
#include "src/tint/sem/u32.h"
#include "src/tint/sem/vector.h"
+#include "src/tint/utils/bitcast.h"
#include "src/tint/utils/compiler_macros.h"
#include "src/tint/utils/map.h"
#include "src/tint/utils/scoped_assignment.h"
@@ -680,6 +681,33 @@
}
template <typename NumberT>
+utils::Result<NumberT> ConstEval::Sub(NumberT a, NumberT b) {
+ NumberT result;
+ if constexpr (IsAbstract<NumberT>) {
+ // Check for over/underflow for abstract values
+ if (auto r = CheckedSub(a, b)) {
+ result = r->value;
+ } else {
+ AddError(OverflowErrorMessage(a, "-", b), *current_source);
+ return utils::Failure;
+ }
+ } else {
+ using T = UnwrapNumber<NumberT>;
+ auto sub_values = [](T lhs, T rhs) {
+ if constexpr (std::is_integral_v<T> && std::is_signed_v<T>) {
+ // Ensure no UB for signed overflow
+ using UT = std::make_unsigned_t<T>;
+ return static_cast<T>(static_cast<UT>(lhs) - static_cast<UT>(rhs));
+ } else {
+ return lhs - rhs;
+ }
+ };
+ result = sub_values(a.value, b.value);
+ }
+ return result;
+}
+
+template <typename NumberT>
utils::Result<NumberT> ConstEval::Mul(NumberT a, NumberT b) {
using T = UnwrapNumber<NumberT>;
NumberT result;
@@ -793,6 +821,37 @@
return r;
}
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Det2(NumberT a1, NumberT a2, NumberT b1, NumberT b2) {
+ auto r1 = Mul(a1, b2);
+ if (!r1) {
+ return utils::Failure;
+ }
+ auto r2 = Mul(b1, a2);
+ if (!r2) {
+ return utils::Failure;
+ }
+ auto r = Sub(r1.Get(), r2.Get());
+ if (!r) {
+ return utils::Failure;
+ }
+ return r;
+}
+
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Clamp(NumberT e, NumberT low, NumberT high) {
+ return NumberT{std::min(std::max(e, low), high)};
+}
+
+auto ConstEval::ClampFunc(const sem::Type* elem_ty) {
+ return [=](auto e, auto low, auto high) -> ImplResult {
+ if (auto r = Clamp(e, low, high)) {
+ return CreateElement(builder, elem_ty, r.Get());
+ }
+ return utils::Failure;
+ };
+}
+
auto ConstEval::AddFunc(const sem::Type* elem_ty) {
return [=](auto a1, auto a2) -> ImplResult {
if (auto r = Add(a1, a2)) {
@@ -802,6 +861,15 @@
};
}
+auto ConstEval::SubFunc(const sem::Type* elem_ty) {
+ return [=](auto a1, auto a2) -> ImplResult {
+ if (auto r = Sub(a1, a2)) {
+ return CreateElement(builder, elem_ty, r.Get());
+ }
+ return utils::Failure;
+ };
+}
+
auto ConstEval::MulFunc(const sem::Type* elem_ty) {
return [=](auto a1, auto a2) -> ImplResult {
if (auto r = Mul(a1, a2)) {
@@ -839,6 +907,15 @@
};
}
+auto ConstEval::Det2Func(const sem::Type* elem_ty) {
+ return [=](auto a, auto b, auto c, auto d) -> ImplResult {
+ if (auto r = Det2(a, b, c, d)) {
+ return CreateElement(builder, elem_ty, r.Get());
+ }
+ return utils::Failure;
+ };
+}
+
ConstEval::Result ConstEval::Literal(const sem::Type* ty, const ast::LiteralExpression* literal) {
return Switch(
literal,
@@ -1100,34 +1177,9 @@
ConstEval::Result ConstEval::OpMinus(const sem::Type* ty,
utils::VectorRef<const sem::Constant*> args,
const Source& source) {
+ TINT_SCOPED_ASSIGNMENT(current_source, &source);
auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
- auto create = [&](auto i, auto j) -> ImplResult {
- using NumberT = decltype(i);
- NumberT result;
- if constexpr (IsAbstract<NumberT>) {
- // Check for over/underflow for abstract values
- if (auto r = CheckedSub(i, j)) {
- result = r->value;
- } else {
- AddError(OverflowErrorMessage(i, "-", j), source);
- return utils::Failure;
- }
- } else {
- using T = UnwrapNumber<NumberT>;
- auto subtract_values = [](T lhs, T rhs) {
- if constexpr (std::is_integral_v<T> && std::is_signed_v<T>) {
- // Ensure no UB for signed underflow
- using UT = std::make_unsigned_t<T>;
- return static_cast<T>(static_cast<UT>(lhs) - static_cast<UT>(rhs));
- } else {
- return lhs - rhs;
- }
- };
- result = subtract_values(i.value, j.value);
- }
- return CreateElement(builder, c0->Type(), result);
- };
- return Dispatch_fia_fiu32_f16(create, c0, c1);
+ return Dispatch_fia_fiu32_f16(SubFunc(c0->Type()), c0, c1);
};
return TransformBinaryElements(builder, ty, transform, args[0], args[1]);
@@ -1612,7 +1664,8 @@
auto create = [&](auto i) -> ImplResult {
using NumberT = decltype(i);
if (i < NumberT(-1.0) || i > NumberT(1.0)) {
- AddError("acos must be called with a value in the range [-1 .. 1] (inclusive)", source);
+ AddError("acos must be called with a value in the range [-1 .. 1] (inclusive)",
+ source);
return utils::Failure;
}
return CreateElement(builder, c0->Type(), NumberT(std::acos(i.value)));
@@ -1641,7 +1694,8 @@
auto create = [&](auto i) -> ImplResult {
using NumberT = decltype(i);
if (i < NumberT(-1.0) || i > NumberT(1.0)) {
- AddError("asin must be called with a value in the range [-1 .. 1] (inclusive)", source);
+ AddError("asin must be called with a value in the range [-1 .. 1] (inclusive)",
+ source);
return utils::Failure;
}
return CreateElement(builder, c0->Type(), NumberT(std::asin(i.value)));
@@ -1687,7 +1741,8 @@
auto create = [&](auto i) -> ImplResult {
using NumberT = decltype(i);
if (i <= NumberT(-1.0) || i >= NumberT(1.0)) {
- AddError("atanh must be called with a value in the range (-1 .. 1) (exclusive)", source);
+ AddError("atanh must be called with a value in the range (-1 .. 1) (exclusive)",
+ source);
return utils::Failure;
}
return CreateElement(builder, c0->Type(), NumberT(std::atanh(i.value)));
@@ -1727,11 +1782,7 @@
const Source&) {
auto transform = [&](const sem::Constant* c0, const sem::Constant* c1,
const sem::Constant* c2) {
- auto create = [&](auto e, auto low, auto high) {
- return CreateElement(builder, c0->Type(),
- decltype(e)(std::min(std::max(e, low), high)));
- };
- return Dispatch_fia_fiu32_f16(create, c0, c1, c2);
+ return Dispatch_fia_fiu32_f16(ClampFunc(c0->Type()), c0, c1, c2);
};
return TransformElements(builder, ty, transform, args[0], args[1], args[2]);
}
@@ -1790,6 +1841,52 @@
return TransformElements(builder, ty, transform, args[0]);
}
+ConstEval::Result ConstEval::cross(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source) {
+ TINT_SCOPED_ASSIGNMENT(current_source, &source);
+ auto* u = args[0];
+ auto* v = args[1];
+ auto* elem_ty = u->Type()->As<sem::Vector>()->type();
+
+ // cross product of a v3 is the determinant of the 3x3 matrix:
+ //
+ // |i j k |
+ // |u0 u1 u2|
+ // |v0 v1 v2|
+ //
+ // |u1 u2|i - |u0 u2|j + |u0 u1|k
+ // |v1 v2| |v0 v2| |v0 v1|
+ //
+ // |u1 u2|i + |v0 v2|j + |u0 u1|k
+ // |v1 v2| |u0 u2| |v0 v1|
+
+ auto* u0 = u->Index(0);
+ auto* u1 = u->Index(1);
+ auto* u2 = u->Index(2);
+ auto* v0 = v->Index(0);
+ auto* v1 = v->Index(1);
+ auto* v2 = v->Index(2);
+
+ // auto x = Dispatch_fa_f32_f16(ab_minus_cd_func(elem_ty), u->Index(1), v->Index(2),
+ // v->Index(1), u->Index(2));
+ auto x = Dispatch_fa_f32_f16(Det2Func(elem_ty), u1, u2, v1, v2);
+ if (!x) {
+ return utils::Failure;
+ }
+ auto y = Dispatch_fa_f32_f16(Det2Func(elem_ty), v0, v2, u0, u2);
+ if (!y) {
+ return utils::Failure;
+ }
+ auto z = Dispatch_fa_f32_f16(Det2Func(elem_ty), u0, u1, v0, v1);
+ if (!z) {
+ return utils::Failure;
+ }
+
+ return CreateComposite(builder, ty,
+ utils::Vector<const sem::Constant*, 3>{x.Get(), y.Get(), z.Get()});
+}
+
ConstEval::Result ConstEval::extractBits(const sem::Type* ty,
utils::VectorRef<const sem::Constant*> args,
const Source& source) {
@@ -1804,18 +1901,18 @@
NumberUT in_offset = args[1]->As<NumberUT>();
NumberUT in_count = args[2]->As<NumberUT>();
- constexpr UT w = sizeof(UT) * 8;
- if ((in_offset + in_count) > w) {
- AddError("'offset + 'count' must be less than or equal to the bit width of 'e'",
- source);
- return utils::Failure;
- }
-
// Cast all to unsigned
UT e = static_cast<UT>(in_e);
UT o = static_cast<UT>(in_offset);
UT c = static_cast<UT>(in_count);
+ constexpr UT w = sizeof(UT) * 8;
+ if (o > w || c > w || (o + c) > w) {
+ AddError("'offset + 'count' must be less than or equal to the bit width of 'e'",
+ source);
+ return utils::Failure;
+ }
+
NumberT result;
if (c == UT{0}) {
// The result is 0 if c is 0
@@ -1979,6 +2076,78 @@
return TransformElements(builder, ty, transform, args[0], args[1]);
}
+ConstEval::Result ConstEval::pack2x16snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto calc = [&](f32 val) -> u32 {
+ auto clamped = Clamp(val, f32(-1.0f), f32(1.0f)).Get();
+ return u32(utils::Bitcast<uint16_t>(
+ static_cast<int16_t>(std::floor(0.5f + (32767.0f * clamped)))));
+ };
+
+ auto* e = args[0];
+ auto e0 = calc(e->Index(0)->As<f32>());
+ auto e1 = calc(e->Index(1)->As<f32>());
+
+ u32 ret = u32((e0 & 0x0000'ffff) | (e1 << 16));
+ return CreateElement(builder, ty, ret);
+}
+
+ConstEval::Result ConstEval::pack2x16unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto calc = [&](f32 val) -> u32 {
+ auto clamped = Clamp(val, f32(0.0f), f32(1.0f)).Get();
+ return u32{std::floor(0.5f + (65535.0f * clamped))};
+ };
+
+ auto* e = args[0];
+ auto e0 = calc(e->Index(0)->As<f32>());
+ auto e1 = calc(e->Index(1)->As<f32>());
+
+ u32 ret = u32((e0 & 0x0000'ffff) | (e1 << 16));
+ return CreateElement(builder, ty, ret);
+}
+
+ConstEval::Result ConstEval::pack4x8snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto calc = [&](f32 val) -> u32 {
+ auto clamped = Clamp(val, f32(-1.0f), f32(1.0f)).Get();
+ return u32(
+ utils::Bitcast<uint8_t>(static_cast<int8_t>(std::floor(0.5f + (127.0f * clamped)))));
+ };
+
+ auto* e = args[0];
+ auto e0 = calc(e->Index(0)->As<f32>());
+ auto e1 = calc(e->Index(1)->As<f32>());
+ auto e2 = calc(e->Index(2)->As<f32>());
+ auto e3 = calc(e->Index(3)->As<f32>());
+
+ uint32_t mask = 0x0000'00ff;
+ u32 ret = u32((e0 & mask) | ((e1 & mask) << 8) | ((e2 & mask) << 16) | ((e3 & mask) << 24));
+ return CreateElement(builder, ty, ret);
+}
+
+ConstEval::Result ConstEval::pack4x8unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto calc = [&](f32 val) -> u32 {
+ auto clamped = Clamp(val, f32(0.0f), f32(1.0f)).Get();
+ return u32{std::floor(0.5f + (255.0f * clamped))};
+ };
+
+ auto* e = args[0];
+ auto e0 = calc(e->Index(0)->As<f32>());
+ auto e1 = calc(e->Index(1)->As<f32>());
+ auto e2 = calc(e->Index(2)->As<f32>());
+ auto e3 = calc(e->Index(3)->As<f32>());
+
+ uint32_t mask = 0x0000'00ff;
+ u32 ret = u32((e0 & mask) | ((e1 & mask) << 8) | ((e2 & mask) << 16) | ((e3 & mask) << 24));
+ return CreateElement(builder, ty, ret);
+}
+
ConstEval::Result ConstEval::reverseBits(const sem::Type* ty,
utils::VectorRef<const sem::Constant*> args,
const Source&) {
@@ -2085,6 +2254,68 @@
return TransformElements(builder, ty, transform, args[0], args[1]);
}
+ConstEval::Result ConstEval::unpack2x16snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto* inner_ty = sem::Type::DeepestElementOf(ty);
+ auto e = args[0]->As<u32>().value;
+
+ utils::Vector<const sem::Constant*, 2> els;
+ els.Reserve(2);
+ for (size_t i = 0; i < 2; ++i) {
+ auto val = f32(
+ std::max(static_cast<float>(int16_t((e >> (16 * i)) & 0x0000'ffff)) / 32767.f, -1.f));
+ els.Push(CreateElement(builder, inner_ty, val));
+ }
+ return CreateComposite(builder, ty, std::move(els));
+}
+
+ConstEval::Result ConstEval::unpack2x16unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto* inner_ty = sem::Type::DeepestElementOf(ty);
+ auto e = args[0]->As<u32>().value;
+
+ utils::Vector<const sem::Constant*, 2> els;
+ els.Reserve(2);
+ for (size_t i = 0; i < 2; ++i) {
+ auto val = f32(static_cast<float>(uint16_t((e >> (16 * i)) & 0x0000'ffff)) / 65535.f);
+ els.Push(CreateElement(builder, inner_ty, val));
+ }
+ return CreateComposite(builder, ty, std::move(els));
+}
+
+ConstEval::Result ConstEval::unpack4x8snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto* inner_ty = sem::Type::DeepestElementOf(ty);
+ auto e = args[0]->As<u32>().value;
+
+ utils::Vector<const sem::Constant*, 4> els;
+ els.Reserve(4);
+ for (size_t i = 0; i < 4; ++i) {
+ auto val =
+ f32(std::max(static_cast<float>(int8_t((e >> (8 * i)) & 0x0000'00ff)) / 127.f, -1.f));
+ els.Push(CreateElement(builder, inner_ty, val));
+ }
+ return CreateComposite(builder, ty, std::move(els));
+}
+
+ConstEval::Result ConstEval::unpack4x8unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source&) {
+ auto* inner_ty = sem::Type::DeepestElementOf(ty);
+ auto e = args[0]->As<u32>().value;
+
+ utils::Vector<const sem::Constant*, 4> els;
+ els.Reserve(4);
+ for (size_t i = 0; i < 4; ++i) {
+ auto val = f32(static_cast<float>(uint8_t((e >> (8 * i)) & 0x0000'00ff)) / 255.f);
+ els.Push(CreateElement(builder, inner_ty, val));
+ }
+ return CreateComposite(builder, ty, std::move(els));
+}
+
ConstEval::Result ConstEval::quantizeToF16(const sem::Type* ty,
utils::VectorRef<const sem::Constant*> args,
const Source&) {
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index f309150..9907b36 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -503,6 +503,15 @@
utils::VectorRef<const sem::Constant*> args,
const Source& source);
+ /// cross builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result cross(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
/// extractBits builtin
/// @param ty the expression type
/// @param args the input arguments
@@ -548,6 +557,42 @@
utils::VectorRef<const sem::Constant*> args,
const Source& source);
+ /// pack2x16snorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result pack2x16snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
+ /// pack2x16unorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result pack2x16unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
+ /// pack4x8snorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result pack4x8snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
+ /// pack4x8unorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result pack4x8unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
/// reverseBits builtin
/// @param ty the expression type
/// @param args the input arguments
@@ -602,6 +647,42 @@
utils::VectorRef<const sem::Constant*> args,
const Source& source);
+ /// unpack2x16snorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result unpack2x16snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
+ /// unpack2x16unorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result unpack2x16unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
+ /// unpack4x8snorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result unpack4x8snorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
+ /// unpack4x8unorm builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location of the conversion
+ /// @return the result value, or null if the value cannot be calculated
+ Result unpack4x8unorm(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
/// quantizeToF16 builtin
/// @param ty the expression type
/// @param args the input arguments
@@ -625,6 +706,13 @@
template <typename NumberT>
utils::Result<NumberT> Add(NumberT a, NumberT b);
+ /// Subtracts two Number<T>s
+ /// @param a the lhs number
+ /// @param b the rhs number
+ /// @returns the result number on success, or logs an error and returns Failure
+ template <typename NumberT>
+ utils::Result<NumberT> Sub(NumberT a, NumberT b);
+
/// Multiplies two Number<T>s
/// @param a the lhs number
/// @param b the rhs number
@@ -677,12 +765,34 @@
NumberT b3,
NumberT b4);
+ /// Returns the determinant of the 2x2 matrix [(a1, a2), (b1, b2)]
+ /// @param a1 component 1 of the first column vector
+ /// @param a2 component 2 of the first column vector
+ /// @param b1 component 1 of the second column vector
+ /// @param b2 component 2 of the second column vector
+ template <typename NumberT>
+ utils::Result<NumberT> Det2(NumberT a1, NumberT a2, NumberT b1, NumberT b2);
+
+ /// Clamps e between low and high
+ /// @param e the number to clamp
+ /// @param low the lower bound
+ /// @param high the upper bound
+ /// @returns the result number on success, or logs an error and returns Failure
+ template <typename NumberT>
+ utils::Result<NumberT> Clamp(NumberT e, NumberT low, NumberT high);
+
/// Returns a callable that calls Add, and creates a Constant with its result of type `elem_ty`
/// if successful, or returns Failure otherwise.
/// @param elem_ty the element type of the Constant to create on success
/// @returns the callable function
auto AddFunc(const sem::Type* elem_ty);
+ /// Returns a callable that calls Sub, and creates a Constant with its result of type `elem_ty`
+ /// if successful, or returns Failure otherwise.
+ /// @param elem_ty the element type of the Constant to create on success
+ /// @returns the callable function
+ auto SubFunc(const sem::Type* elem_ty);
+
/// Returns a callable that calls Mul, and creates a Constant with its result of type `elem_ty`
/// if successful, or returns Failure otherwise.
/// @param elem_ty the element type of the Constant to create on success
@@ -707,6 +817,18 @@
/// @returns the callable function
auto Dot4Func(const sem::Type* elem_ty);
+ /// Returns a callable that calls Det2, and creates a Constant with its result of type `elem_ty`
+ /// if successful, or returns Failure otherwise.
+ /// @param elem_ty the element type of the Constant to create on success
+ /// @returns the callable function
+ auto Det2Func(const sem::Type* elem_ty);
+
+ /// Returns a callable that calls Clamp, and creates a Constant with its result of type
+ /// `elem_ty` if successful, or returns Failure otherwise.
+ /// @param elem_ty the element type of the Constant to create on success
+ /// @returns the callable function
+ auto ClampFunc(const sem::Type* elem_ty);
+
ProgramBuilder& builder;
const Source* current_source = nullptr;
};
diff --git a/src/tint/resolver/const_eval_builtin_test.cc b/src/tint/resolver/const_eval_builtin_test.cc
index 3dc395b..fddca38 100644
--- a/src/tint/resolver/const_eval_builtin_test.cc
+++ b/src/tint/resolver/const_eval_builtin_test.cc
@@ -9,7 +9,7 @@
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
+// See the License for the empecific language governing permissions and
// limitations under the License.
#include "src/tint/resolver/const_eval_test.h"
@@ -17,6 +17,7 @@
#include "src/tint/utils/result.h"
using namespace tint::number_suffixes; // NOLINT
+using ::testing::HasSubstr;
namespace tint::resolver {
namespace {
@@ -28,8 +29,8 @@
Case(utils::VectorRef<Types> in_args, Types expected_value)
: args(std::move(in_args)), expected(Success{std::move(expected_value), false, false}) {}
- Case(utils::VectorRef<Types> in_args, const char* expected_err)
- : args(std::move(in_args)), expected(Failure{expected_err}) {}
+ Case(utils::VectorRef<Types> in_args, std::string expected_err)
+ : args(std::move(in_args)), expected(Failure{std::move(expected_err)}) {}
/// Expected value may be positive or negative
Case& PosOrNeg() {
@@ -53,7 +54,7 @@
bool float_compare = false;
};
struct Failure {
- const char* error = nullptr;
+ std::string error = nullptr;
};
utils::Vector<Types, 8> args;
@@ -94,17 +95,27 @@
}
/// Creates a Case with Values for args and expected error
-static Case E(std::initializer_list<Types> args, const char* err) {
- return Case{utils::Vector<Types, 8>{args}, err};
+static Case E(std::initializer_list<Types> args, std::string err) {
+ return Case{utils::Vector<Types, 8>{args}, std::move(err)};
}
/// Convenience overload that creates an expected-error Case with just scalars
-static Case E(std::initializer_list<ScalarTypes> sargs, const char* err) {
+static Case E(std::initializer_list<ScalarTypes> sargs, std::string err) {
utils::Vector<Types, 8> args;
for (auto& sa : sargs) {
std::visit([&](auto&& v) { return args.Push(Val(v)); }, sa);
}
- return Case{std::move(args), err};
+ return Case{std::move(args), std::move(err)};
+}
+
+/// Returns the overflow error message for binary ops
+template <typename NumberT>
+std::string OverflowErrorMessage(NumberT lhs, const char* op, NumberT rhs) {
+ std::stringstream ss;
+ ss << std::setprecision(20);
+ ss << "'" << lhs.value << " " << op << " " << rhs.value << "' cannot be represented as '"
+ << FriendlyName<NumberT>() << "'";
+ return ss.str();
}
using ResolverConstEvalBuiltinTest = ResolverTestWithParam<std::tuple<sem::BuiltinType, Case>>;
@@ -776,6 +787,130 @@
testing::ValuesIn(Concat(CountOneBitsCases<i32>(), //
CountOneBitsCases<u32>()))));
+template <typename T, bool finite_only>
+std::vector<Case> CrossCases() {
+ constexpr auto vec_x = [](T v) { return Vec(T(v), T(0), T(0)); };
+ constexpr auto vec_y = [](T v) { return Vec(T(0), T(v), T(0)); };
+ constexpr auto vec_z = [](T v) { return Vec(T(0), T(0), T(v)); };
+
+ const auto zero = Vec(T(0), T(0), T(0));
+ const auto unit_x = vec_x(T(1));
+ const auto unit_y = vec_y(T(1));
+ const auto unit_z = vec_z(T(1));
+ const auto neg_unit_x = vec_x(-T(1));
+ const auto neg_unit_y = vec_y(-T(1));
+ const auto neg_unit_z = vec_z(-T(1));
+ const auto highest_x = vec_x(T::Highest());
+ const auto highest_y = vec_y(T::Highest());
+ const auto highest_z = vec_z(T::Highest());
+ const auto smallest_x = vec_x(T::Smallest());
+ const auto smallest_y = vec_y(T::Smallest());
+ const auto smallest_z = vec_z(T::Smallest());
+ const auto lowest_x = vec_x(T::Lowest());
+ const auto lowest_y = vec_y(T::Lowest());
+ const auto lowest_z = vec_z(T::Lowest());
+ const auto inf_x = vec_x(T::Inf());
+ const auto inf_y = vec_y(T::Inf());
+ const auto inf_z = vec_z(T::Inf());
+ const auto neg_inf_x = vec_x(-T::Inf());
+ const auto neg_inf_y = vec_y(-T::Inf());
+ const auto neg_inf_z = vec_z(-T::Inf());
+
+ std::vector<Case> r = {
+ C({zero, zero}, zero),
+
+ C({unit_x, unit_x}, zero),
+ C({unit_y, unit_y}, zero),
+ C({unit_z, unit_z}, zero),
+
+ C({smallest_x, smallest_x}, zero),
+ C({smallest_y, smallest_y}, zero),
+ C({smallest_z, smallest_z}, zero),
+
+ C({lowest_x, lowest_x}, zero),
+ C({lowest_y, lowest_y}, zero),
+ C({lowest_z, lowest_z}, zero),
+
+ C({highest_x, highest_x}, zero),
+ C({highest_y, highest_y}, zero),
+ C({highest_z, highest_z}, zero),
+
+ C({smallest_x, highest_x}, zero),
+ C({smallest_y, highest_y}, zero),
+ C({smallest_z, highest_z}, zero),
+
+ C({unit_x, neg_unit_x}, zero).PosOrNeg(),
+ C({unit_y, neg_unit_y}, zero).PosOrNeg(),
+ C({unit_z, neg_unit_z}, zero).PosOrNeg(),
+
+ C({unit_x, unit_y}, unit_z),
+ C({unit_y, unit_x}, neg_unit_z),
+
+ C({unit_z, unit_x}, unit_y),
+ C({unit_x, unit_z}, neg_unit_y),
+
+ C({unit_y, unit_z}, unit_x),
+ C({unit_z, unit_y}, neg_unit_x),
+
+ C({vec_x(T(1)), vec_y(T(2))}, vec_z(T(2))),
+ C({vec_y(T(1)), vec_x(T(2))}, vec_z(-T(2))),
+ C({vec_x(T(2)), vec_y(T(3))}, vec_z(T(6))),
+ C({vec_y(T(2)), vec_x(T(3))}, vec_z(-T(6))),
+
+ C({Vec(T(1), T(2), T(3)), Vec(T(1), T(5), T(7))}, Vec(T(-1), T(-4), T(3))),
+ C({Vec(T(33), T(44), T(55)), Vec(T(13), T(42), T(39))}, Vec(T(-594), T(-572), T(814))),
+ C({Vec(T(3.5), T(4), T(5.5)), Vec(T(1), T(4.5), T(3.5))},
+ Vec(T(-10.75), T(-6.75), T(11.75))),
+ };
+
+ ConcatIntoIf<!finite_only>( //
+ r, std::vector<Case>{
+ C({highest_x, highest_y}, inf_z).PosOrNeg(), //
+ C({highest_y, highest_x}, inf_z).PosOrNeg(), //
+ C({highest_z, highest_x}, inf_y).PosOrNeg(), //
+ C({highest_x, highest_z}, inf_y).PosOrNeg(), //
+ C({highest_y, highest_z}, inf_x).PosOrNeg(), //
+ C({highest_z, highest_y}, inf_x).PosOrNeg(), //
+ C({lowest_x, lowest_y}, inf_z).PosOrNeg(), //
+ C({lowest_y, lowest_x}, inf_z).PosOrNeg(), //
+ C({lowest_z, lowest_x}, inf_y).PosOrNeg(), //
+ C({lowest_x, lowest_z}, inf_y).PosOrNeg(), //
+ C({lowest_y, lowest_z}, inf_x).PosOrNeg(), //
+ C({lowest_z, lowest_y}, inf_x).PosOrNeg(),
+ });
+
+ std::string pos_error_msg =
+ "12:34 error: " + OverflowErrorMessage(T::Highest(), "*", T::Highest());
+ std::string neg_error_msg =
+ "12:34 error: " + OverflowErrorMessage(T::Lowest(), "*", T::Lowest());
+
+ ConcatIntoIf<finite_only>( //
+ r, std::vector<Case>{
+ E({highest_x, highest_y}, pos_error_msg),
+ E({highest_y, highest_x}, pos_error_msg),
+ E({highest_z, highest_x}, pos_error_msg),
+ E({highest_x, highest_z}, pos_error_msg),
+ E({highest_y, highest_z}, pos_error_msg),
+ E({highest_z, highest_y}, pos_error_msg),
+ E({lowest_x, lowest_y}, neg_error_msg),
+ E({lowest_y, lowest_x}, neg_error_msg),
+ E({lowest_z, lowest_x}, neg_error_msg),
+ E({lowest_x, lowest_z}, neg_error_msg),
+ E({lowest_y, lowest_z}, neg_error_msg),
+ E({lowest_z, lowest_y}, neg_error_msg),
+ });
+
+ return r;
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Cross,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kCross),
+ testing::ValuesIn(Concat(CrossCases<AFloat, true>(), //
+ CrossCases<f32, false>(),
+ CrossCases<f32, false>(), //
+ CrossCases<f16, false>()))));
+
template <typename T>
std::vector<Case> FirstLeadingBitCases() {
using B = BitValues<T>;
@@ -1091,8 +1226,77 @@
std::make_tuple(33, 33), //
std::make_tuple(34, 34), //
std::make_tuple(1000, 1000), //
+ std::make_tuple(u32::Highest(), 1), //
+ std::make_tuple(1, u32::Highest()), //
std::make_tuple(u32::Highest(), u32::Highest())));
+std::vector<Case> Pack4x8snormCases() {
+ return {
+ C({Vec(f32(0), f32(0), f32(0), f32(0))}, Val(u32(0x0000'0000))),
+ C({Vec(f32(0), f32(0), f32(0), f32(-1))}, Val(u32(0x8100'0000))),
+ C({Vec(f32(0), f32(0), f32(0), f32(1))}, Val(u32(0x7f00'0000))),
+ C({Vec(f32(0), f32(0), f32(-1), f32(0))}, Val(u32(0x0081'0000))),
+ C({Vec(f32(0), f32(1), f32(0), f32(0))}, Val(u32(0x0000'7f00))),
+ C({Vec(f32(-1), f32(0), f32(0), f32(0))}, Val(u32(0x0000'0081))),
+ C({Vec(f32(1), f32(-1), f32(1), f32(-1))}, Val(u32(0x817f'817f))),
+ C({Vec(f32::Highest(), f32(-0.5), f32(0.5), f32::Lowest())}, Val(u32(0x8140'c17f))),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Pack4x8snorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kPack4X8Snorm),
+ testing::ValuesIn(Pack4x8snormCases())));
+
+std::vector<Case> Pack4x8unormCases() {
+ return {
+ C({Vec(f32(0), f32(0), f32(0), f32(0))}, Val(u32(0x0000'0000))),
+ C({Vec(f32(0), f32(0), f32(0), f32(1))}, Val(u32(0xff00'0000))),
+ C({Vec(f32(0), f32(0), f32(1), f32(0))}, Val(u32(0x00ff'0000))),
+ C({Vec(f32(0), f32(1), f32(0), f32(0))}, Val(u32(0x0000'ff00))),
+ C({Vec(f32(1), f32(0), f32(0), f32(0))}, Val(u32(0x0000'00ff))),
+ C({Vec(f32(1), f32(0), f32(1), f32(0))}, Val(u32(0x00ff'00ff))),
+ C({Vec(f32::Highest(), f32(0), f32(0.5), f32::Lowest())}, Val(u32(0x0080'00ff))),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Pack4x8unorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kPack4X8Unorm),
+ testing::ValuesIn(Pack4x8unormCases())));
+
+std::vector<Case> Pack2x16snormCases() {
+ return {
+ C({Vec(f32(0), f32(0))}, Val(u32(0x0000'0000))),
+ C({Vec(f32(0), f32(-1))}, Val(u32(0x8001'0000))),
+ C({Vec(f32(0), f32(1))}, Val(u32(0x7fff'0000))),
+ C({Vec(f32(-1), f32(0))}, Val(u32(0x0000'8001))),
+ C({Vec(f32(1), f32(0))}, Val(u32(0x0000'7fff))),
+ C({Vec(f32(1), f32(-1))}, Val(u32(0x8001'7fff))),
+ C({Vec(f32::Highest(), f32::Lowest())}, Val(u32(0x8001'7fff))),
+ C({Vec(f32(-0.5), f32(0.5))}, Val(u32(0x4000'c001))),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Pack2x16snorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kPack2X16Snorm),
+ testing::ValuesIn(Pack2x16snormCases())));
+
+std::vector<Case> Pack2x16unormCases() {
+ return {
+ C({Vec(f32(0), f32(1))}, Val(u32(0xffff'0000))),
+ C({Vec(f32(1), f32(0))}, Val(u32(0x0000'ffff))),
+ C({Vec(f32(0.5), f32(0))}, Val(u32(0x0000'8000))),
+ C({Vec(f32::Highest(), f32::Lowest())}, Val(u32(0x0000'ffff))),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Pack2x16unorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kPack2X16Unorm),
+ testing::ValuesIn(Pack2x16unormCases())));
+
template <typename T>
std::vector<Case> ReverseBitsCases() {
using B = BitValues<T>;
@@ -1268,6 +1472,74 @@
StepCases<f32>(),
StepCases<f16>()))));
+std::vector<Case> Unpack4x8snormCases() {
+ return {
+ C({Val(u32(0x0000'0000))}, Vec(f32(0), f32(0), f32(0), f32(0))),
+ C({Val(u32(0x8100'0000))}, Vec(f32(0), f32(0), f32(0), f32(-1))),
+ C({Val(u32(0x7f00'0000))}, Vec(f32(0), f32(0), f32(0), f32(1))),
+ C({Val(u32(0x0081'0000))}, Vec(f32(0), f32(0), f32(-1), f32(0))),
+ C({Val(u32(0x0000'7f00))}, Vec(f32(0), f32(1), f32(0), f32(0))),
+ C({Val(u32(0x0000'0081))}, Vec(f32(-1), f32(0), f32(0), f32(0))),
+ C({Val(u32(0x817f'817f))}, Vec(f32(1), f32(-1), f32(1), f32(-1))),
+ C({Val(u32(0x816d'937f))},
+ Vec(f32(1), f32(-0.8582677165354), f32(0.8582677165354), f32(-1))),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Unpack4x8snorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kUnpack4X8Snorm),
+ testing::ValuesIn(Unpack4x8snormCases())));
+
+std::vector<Case> Unpack4x8unormCases() {
+ return {
+ C({Val(u32(0x0000'0000))}, Vec(f32(0), f32(0), f32(0), f32(0))),
+ C({Val(u32(0xff00'0000))}, Vec(f32(0), f32(0), f32(0), f32(1))),
+ C({Val(u32(0x00ff'0000))}, Vec(f32(0), f32(0), f32(1), f32(0))),
+ C({Val(u32(0x0000'ff00))}, Vec(f32(0), f32(1), f32(0), f32(0))),
+ C({Val(u32(0x0000'00ff))}, Vec(f32(1), f32(0), f32(0), f32(0))),
+ C({Val(u32(0x00ff'00ff))}, Vec(f32(1), f32(0), f32(1), f32(0))),
+ C({Val(u32(0x0066'00ff))}, Vec(f32(1), f32(0), f32(0.4), f32(0))),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Unpack4x8unorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kUnpack4X8Unorm),
+ testing::ValuesIn(Unpack4x8unormCases())));
+
+std::vector<Case> Unpack2x16snormCases() {
+ return {
+ C({Val(u32(0x0000'0000))}, Vec(f32(0), f32(0))),
+ C({Val(u32(0x8001'0000))}, Vec(f32(0), f32(-1))),
+ C({Val(u32(0x7fff'0000))}, Vec(f32(0), f32(1))),
+ C({Val(u32(0x0000'8001))}, Vec(f32(-1), f32(0))),
+ C({Val(u32(0x0000'7fff))}, Vec(f32(1), f32(0))),
+ C({Val(u32(0x8001'7fff))}, Vec(f32(1), f32(-1))),
+ C({Val(u32(0x8001'7fff))}, Vec(f32(1), f32(-1))),
+ C({Val(u32(0x4000'999a))}, Vec(f32(-0.80001220740379), f32(0.500015259254737))).FloatComp(),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Unpack2x16snorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kUnpack2X16Snorm),
+ testing::ValuesIn(Unpack2x16snormCases())));
+
+std::vector<Case> Unpack2x16unormCases() {
+ return {
+ C({Val(u32(0xffff'0000))}, Vec(f32(0), f32(1))),
+ C({Val(u32(0x0000'ffff))}, Vec(f32(1), f32(0))),
+ C({Val(u32(0x0000'6666))}, Vec(f32(0.4), f32(0))),
+ C({Val(u32(0x0000'ffff))}, Vec(f32(1), f32(0))),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Unpack2x16unorm,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kUnpack2X16Unorm),
+ testing::ValuesIn(Unpack2x16unormCases())));
+
std::vector<Case> QuantizeToF16Cases() {
(void)E({Vec(0_f, 0_f)}, ""); // Currently unused, but will be soon.
return {
diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc
index 5a21c55..3238c8b 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -193,8 +193,6 @@
},
[&](const ast::Function* func) {
Declare(func->symbol, func);
- TraverseAttributes(func->attributes);
- TraverseAttributes(func->return_type_attributes);
TraverseFunction(func);
},
[&](const ast::Variable* var) {
@@ -216,10 +214,13 @@
/// Traverses the function, performing symbol resolution and determining
/// global dependencies.
void TraverseFunction(const ast::Function* func) {
+ TraverseAttributes(func->attributes);
+ TraverseAttributes(func->return_type_attributes);
// Perform symbol resolution on all the parameter types before registering
// the parameters themselves. This allows the case of declaring a parameter
// with the same identifier as its type.
for (auto* param : func->params) {
+ TraverseAttributes(param->attributes);
TraverseType(param->type);
}
// Resolve the return type
diff --git a/src/tint/resolver/dependency_graph_test.cc b/src/tint/resolver/dependency_graph_test.cc
index 311c9cc..272357b 100644
--- a/src/tint/resolver/dependency_graph_test.cc
+++ b/src/tint/resolver/dependency_graph_test.cc
@@ -1237,9 +1237,14 @@
})});
GlobalVar(Sym(), T, V);
GlobalConst(Sym(), T, V);
- Func(Sym(), //
- utils::Vector{Param(Sym(), T)}, //
- T, // Return type
+ Func(Sym(),
+ utils::Vector{
+ Param(Sym(), T,
+ utils::Vector{
+ Location(V), // Parameter attributes
+ }),
+ },
+ T, // Return type
utils::Vector{
Decl(Var(Sym(), T, V)), //
Decl(Let(Sym(), T, V)), //
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index ddce1ba..dcba75b 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -13766,7 +13766,7 @@
/* parameters */ &kParameters[841],
/* return matcher indices */ &kMatcherIndices[102],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::unpack4x8unorm,
},
{
/* [453] */
@@ -13778,7 +13778,7 @@
/* parameters */ &kParameters[842],
/* return matcher indices */ &kMatcherIndices[102],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::unpack4x8snorm,
},
{
/* [454] */
@@ -13790,7 +13790,7 @@
/* parameters */ &kParameters[843],
/* return matcher indices */ &kMatcherIndices[132],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::unpack2x16unorm,
},
{
/* [455] */
@@ -13802,7 +13802,7 @@
/* parameters */ &kParameters[844],
/* return matcher indices */ &kMatcherIndices[132],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::unpack2x16snorm,
},
{
/* [456] */
@@ -13862,7 +13862,7 @@
/* parameters */ &kParameters[881],
/* return matcher indices */ &kMatcherIndices[95],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::pack4x8unorm,
},
{
/* [461] */
@@ -13874,7 +13874,7 @@
/* parameters */ &kParameters[882],
/* return matcher indices */ &kMatcherIndices[95],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::pack4x8snorm,
},
{
/* [462] */
@@ -13886,7 +13886,7 @@
/* parameters */ &kParameters[884],
/* return matcher indices */ &kMatcherIndices[95],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::pack2x16snorm,
},
{
/* [463] */
@@ -13917,12 +13917,12 @@
/* num parameters */ 2,
/* num template types */ 1,
/* num template numbers */ 0,
- /* template types */ &kTemplateTypes[22],
+ /* template types */ &kTemplateTypes[23],
/* template numbers */ &kTemplateNumbers[10],
/* parameters */ &kParameters[777],
/* return matcher indices */ &kMatcherIndices[104],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::cross,
},
{
/* [466] */
@@ -14006,7 +14006,7 @@
/* parameters */ &kParameters[883],
/* return matcher indices */ &kMatcherIndices[95],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::pack2x16unorm,
},
};
@@ -14020,8 +14020,8 @@
},
{
/* [1] */
- /* fn acos<T : fa_f32_f16>(@test_value(0.87758256189) T) -> T */
- /* fn acos<N : num, T : fa_f32_f16>(@test_value(0.87758256189) vec<N, T>) -> vec<N, T> */
+ /* fn acos<T : fa_f32_f16>(@test_value(0.96891242171) T) -> T */
+ /* fn acos<N : num, T : fa_f32_f16>(@test_value(0.96891242171) vec<N, T>) -> vec<N, T> */
/* num overloads */ 2,
/* overloads */ &kOverloads[428],
},
@@ -14138,7 +14138,7 @@
},
{
/* [18] */
- /* fn cross<T : f32_f16>(vec3<T>, vec3<T>) -> vec3<T> */
+ /* fn cross<T : fa_f32_f16>(vec3<T>, vec3<T>) -> vec3<T> */
/* num overloads */ 1,
/* overloads */ &kOverloads[465],
},
diff --git a/src/tint/tint_gdb.py b/src/tint/tint_gdb.py
index 2dce446..88c614d 100644
--- a/src/tint/tint_gdb.py
+++ b/src/tint/tint_gdb.py
@@ -165,11 +165,11 @@
'UtilsVector', '^tint::utils::VectorRef<.*>$', UtilsVectorRefPrinter)
-class UtilsHashsetPrinter(Printer):
- '''Printer for Hashset<T, N, HASH, EQUAL>'''
+class UtilsHashmapBasePrinter(Printer):
+ '''Base Printer for HashmapBase-derived types'''
def __init__(self, val):
- super(UtilsHashsetPrinter, self).__init__(val)
+ super(UtilsHashmapBasePrinter, self).__init__(val)
self.slice = UtilsVectorPrinter(self.val['slots_']).slice_printer()
self.try_read_std_optional_func = self.try_read_std_optional
@@ -185,26 +185,32 @@
for slot in range(0, self.slice.length()):
v = self.slice.value_at(slot)
if v['hash'] != 0:
- value = v['value']
+ entry = v['entry']
- # value is a std::optional, let's try to extract its value for display
- kvp = self.try_read_std_optional_func(slot, value)
+ # entry is a std::optional, let's try to extract its value for display
+ kvp = self.try_read_std_optional_func(slot, entry)
if kvp is None:
- # If we failed, just output the slot and value as is, which will use
+ # If we failed, just output the slot and entry as is, which will use
# the default visualizer for each.
- kvp = slot, value
+ kvp = slot, entry
yield str(kvp[0]), kvp[1]
def display_hint(self):
return 'array'
- def try_read_std_optional(self, slot, value):
+ def try_read_std_optional(self, slot, entry):
+ return None
+
+
+class UtilsHashsetPrinter(UtilsHashmapBasePrinter):
+ '''Printer for Hashset<T, N, HASH, EQUAL>'''
+
+ def try_read_std_optional(self, slot, entry):
try:
# libstdc++
- v = value['_M_payload']['_M_payload']['_M_value']
+ v = entry['_M_payload']['_M_payload']['_M_value']
return slot, v
- # return str(kvp['key']), kvp['value']
except:
return None
@@ -213,33 +219,16 @@
'UtilsHashset', '^tint::utils::Hashset<.*>$', UtilsHashsetPrinter)
-class UtilsHashmapPrinter(Printer):
+class UtilsHashmapPrinter(UtilsHashmapBasePrinter):
'''Printer for Hashmap<K, V, N, HASH, EQUAL>'''
- def __init__(self, val):
- super(UtilsHashmapPrinter, self).__init__(val)
- self.hash_set = UtilsHashsetPrinter(self.val['set_'])
- # Replace the lookup function so we can extract the key and value out of the std::optionals in the Hashset
- self.hash_set.try_read_std_optional_func = self.try_read_std_optional
-
- def to_string(self):
- return self.hash_set.to_string()
-
- def children(self):
- return self.hash_set.children()
-
- def display_hint(self):
- return 'array'
-
- def try_read_std_optional(self, slot, value):
+ def try_read_std_optional(self, slot, entry):
try:
# libstdc++
- kvp = value['_M_payload']['_M_payload']['_M_value']
+ kvp = entry['_M_payload']['_M_payload']['_M_value']
return str(kvp['key']), kvp['value']
except:
- pass
- # Failed, fall back on hash_set
- return self.hash_set.try_read_std_optional(slot, value)
+ return None
pp_set.add_printer(
diff --git a/src/tint/tint_lldb.py b/src/tint/tint_lldb.py
index 85a93ff..84e6ad3 100644
--- a/src/tint/tint_lldb.py
+++ b/src/tint/tint_lldb.py
@@ -297,8 +297,8 @@
return self.slice_printer.get_child_at_index(index)
-class UtilsHashsetPrinter(Printer):
- '''Printer for Hashset<T, N, HASH, EQUAL>'''
+class UtilsHashmapBasePrinter(Printer):
+ '''Base Printer for HashmapBase-derived types'''
def initialize(self):
self.slice = UtilsVectorPrinter(
@@ -326,26 +326,33 @@
def get_child_at_index(self, index):
slot = self.valid_slots[index]
v = self.slice.value_at(slot)
- value = member(v, 'value')
+ entry = member(v, 'entry')
- # value is a std::optional, let's try to extract its value for display
- kvp = self.try_read_std_optional_func(slot, value)
+ # entry is a std::optional, let's try to extract its value for display
+ kvp = self.try_read_std_optional_func(slot, entry)
if kvp is None:
- # If we failed, just output the slot and value as is, which will use
+ # If we failed, just output the slot and entry as is, which will use
# the default printer for std::optional.
- kvp = slot, value
+ kvp = slot, entry
return kvp[1].CreateChildAtOffset('[{}]'.format(kvp[0]), 0, kvp[1].GetType())
- def try_read_std_optional(self, slot, value):
+ def try_read_std_optional(self, slot, entry):
+ return None
+
+
+class UtilsHashsetPrinter(UtilsHashmapBasePrinter):
+ '''Printer for Hashset<T, N, HASH, EQUAL>'''
+
+ def try_read_std_optional(self, slot, entry):
try:
# libc++
- v = value.EvaluateExpression('__val_')
+ v = entry.EvaluateExpression('__val_')
if v.name is not None:
return slot, v
# libstdc++
- v = value.EvaluateExpression('_M_payload._M_payload._M_value')
+ v = entry.EvaluateExpression('_M_payload._M_payload._M_value')
if v.name is not None:
return slot, v
return None
@@ -353,45 +360,24 @@
return None
-class UtilsHashmapPrinter(Printer):
+class UtilsHashmapPrinter(UtilsHashsetPrinter):
'''Printer for Hashmap<K, V, N, HASH, EQUAL>'''
- def initialize(self):
- self.hash_set = UtilsHashsetPrinter(self.member('set_'))
- # Replace the lookup function so we can extract the key and value out of the std::optionals in the Hashset
- self.hash_set.try_read_std_optional_func = self.try_read_std_optional
-
- def update(self):
- self.hash_set.update()
-
- def get_summary(self):
- return self.hash_set.get_summary()
-
- def num_children(self):
- return self.hash_set.num_children()
-
- def has_children(self):
- return self.hash_set.has_children()
-
- def get_child_at_index(self, index):
- return self.hash_set.get_child_at_index(index)
-
- def try_read_std_optional(self, slot, value):
+ def try_read_std_optional(self, slot, entry):
try:
# libc++
- val = value.EvaluateExpression('__val_')
+ val = entry.EvaluateExpression('__val_')
k = val.EvaluateExpression('key')
v = val.EvaluateExpression('value')
if k.name is not None and v.name is not None:
return k.GetValue(), v
# libstdc++
- val = value.EvaluateExpression('_M_payload._M_payload._M_value')
+ val = entry.EvaluateExpression('_M_payload._M_payload._M_value')
k = val.EvaluateExpression('key')
v = val.EvaluateExpression('value')
if k.name is not None and v.name is not None:
return k.GetValue(), v
+ return None
except:
- pass
- # Failed, fall back on hash_set
- return self.hash_set.try_read_std_optional(slot, value)
+ return None
diff --git a/src/tint/transform/builtin_polyfill.cc b/src/tint/transform/builtin_polyfill.cc
index e80436d..17fcc20 100644
--- a/src/tint/transform/builtin_polyfill.cc
+++ b/src/tint/transform/builtin_polyfill.cc
@@ -462,6 +462,14 @@
auto name = b.Symbols().New("tint_insert_bits");
uint32_t width = WidthOf(ty);
+ // Currently in WGSL parameters of insertBits must be i32, u32, vecN<i32> or vecN<u32>
+ if (!sem::Type::DeepestElementOf(ty)->IsAnyOf<sem::I32, sem::U32>()) {
+ TINT_ICE(Transform, b.Diagnostics())
+ << "insertBits polyfill only support i32, u32, and vector of i32 or u32, got "
+ << b.FriendlyName(ty);
+ return {};
+ }
+
constexpr uint32_t W = 32u; // 32-bit
auto V = [&](auto value) -> const ast::Expression* {
@@ -481,21 +489,60 @@
return b.vec(b.ty.u32(), width, value);
};
- utils::Vector<const ast::Statement*, 8> body = {
- b.Decl(b.Let("s", b.Call("min", "offset", u32(W)))),
- b.Decl(b.Let("e", b.Call("min", u32(W), b.Add("s", "count")))),
- };
+ // Polyfill algorithm:
+ // s = min(offset, 32u);
+ // e = min(32u, (s + count));
+ // mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+ // return (((n << s) & mask) | (v & ~(mask)));
+ // Note that the algorithm above use the left-shifting in C++ manner, but in WGSL, HLSL, MSL
+ // the rhs are modulo to bit-width of lhs (that is 32u in this case), and in GLSL the result
+ // is undefined if rhs is greater than or equal to bit-width of lhs. The results of `x << y`
+ // in C++ and HLSL are different when `y >= 32u`, and the `s` and `e` defined above can be
+ // 32u, which are cases we must handle specially. Replace all `(x << y)` to
+ // `select(Tx(), x << y, y < 32u)`, in which `Tx` is the type of x, where y can be greater
+ // than or equal to 32u.
+ // WGSL polyfill function:
+ // fn tint_insert_bits(v : T, n : T, offset : u32, count : u32) -> T {
+ // let e = offset + count;
+ // let mask = (
+ // (select(0u, 1u << offset, offset < 32u) - 1u) ^
+ // (select(0u, 1u << e, e < 32u) - 1u)
+ // );
+ // return ((select(T(), n << offset, offset < 32u) & mask) | (v & ~(mask)));
+ // }
+
+ utils::Vector<const ast::Statement*, 8> body;
switch (polyfill.insert_bits) {
case Level::kFull:
- // let mask = ((1 << s) - 1) ^ ((1 << e) - 1)
+ // let e = offset + count;
+ body.Push(b.Decl(b.Let("e", b.Add("offset", "count"))));
+
+ // let mask = (
+ // (select(0u, 1u << offset, offset < 32u) - 1u) ^
+ // (select(0u, 1u << e, e < 32u) - 1u)
+ // );
body.Push(b.Decl(b.Let(
- "mask", b.Xor(b.Sub(b.Shl(1_u, "s"), 1_u), b.Sub(b.Shl(1_u, "e"), 1_u)))));
- // return ((n << s) & mask) | (v & ~mask)
- body.Push(b.Return(b.Or(b.And(b.Shl("n", U("s")), V("mask")),
- b.And("v", V(b.Complement("mask"))))));
+ "mask",
+ b.Xor( //
+ b.Sub(
+ b.Call("select", 0_u, b.Shl(1_u, "offset"), b.LessThan("offset", 32_u)),
+ 1_u),
+ b.Sub(b.Call("select", 0_u, b.Shl(1_u, "e"), b.LessThan("e", 32_u)),
+ 1_u) //
+ ))));
+
+ // return ((select(T(), n << offset, offset < 32u) & mask) | (v & ~(mask)));
+ body.Push(
+ b.Return(b.Or(b.And(b.Call("select", b.Construct(T(ty)),
+ b.Shl("n", U("offset")), b.LessThan("offset", 32_u)),
+ V("mask")),
+ b.And("v", V(b.Complement("mask"))))));
+
break;
case Level::kClampParameters:
+ body.Push(b.Decl(b.Let("s", b.Call("min", "offset", u32(W)))));
+ body.Push(b.Decl(b.Let("e", b.Call("min", u32(W), b.Add("s", "count")))));
body.Push(b.Return(b.Call("insertBits", "v", "n", "s", b.Sub("e", "s"))));
break;
default:
diff --git a/src/tint/transform/builtin_polyfill_test.cc b/src/tint/transform/builtin_polyfill_test.cc
index 3b7a42c..1e380d4 100644
--- a/src/tint/transform/builtin_polyfill_test.cc
+++ b/src/tint/transform/builtin_polyfill_test.cc
@@ -1722,10 +1722,9 @@
auto* expect = R"(
fn tint_insert_bits(v : i32, n : i32, offset : u32, count : u32) -> i32 {
- let s = min(offset, 32u);
- let e = min(32u, (s + count));
- let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
- return (((n << s) & i32(mask)) | (v & i32(~(mask))));
+ let e = (offset + count);
+ let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+ return ((select(i32(), (n << offset), (offset < 32u)) & i32(mask)) | (v & i32(~(mask))));
}
fn f() {
@@ -1749,10 +1748,9 @@
auto* expect = R"(
fn tint_insert_bits(v : u32, n : u32, offset : u32, count : u32) -> u32 {
- let s = min(offset, 32u);
- let e = min(32u, (s + count));
- let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
- return (((n << s) & mask) | (v & ~(mask)));
+ let e = (offset + count);
+ let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+ return ((select(u32(), (n << offset), (offset < 32u)) & mask) | (v & ~(mask)));
}
fn f() {
@@ -1776,10 +1774,9 @@
auto* expect = R"(
fn tint_insert_bits(v : vec3<i32>, n : vec3<i32>, offset : u32, count : u32) -> vec3<i32> {
- let s = min(offset, 32u);
- let e = min(32u, (s + count));
- let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
- return (((n << vec3<u32>(s)) & vec3<i32>(i32(mask))) | (v & vec3<i32>(i32(~(mask)))));
+ let e = (offset + count);
+ let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+ return ((select(vec3<i32>(), (n << vec3<u32>(offset)), (offset < 32u)) & vec3<i32>(i32(mask))) | (v & vec3<i32>(i32(~(mask)))));
}
fn f() {
@@ -1803,10 +1800,9 @@
auto* expect = R"(
fn tint_insert_bits(v : vec3<u32>, n : vec3<u32>, offset : u32, count : u32) -> vec3<u32> {
- let s = min(offset, 32u);
- let e = min(32u, (s + count));
- let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
- return (((n << vec3<u32>(s)) & vec3<u32>(mask)) | (v & vec3<u32>(~(mask))));
+ let e = (offset + count);
+ let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+ return ((select(vec3<u32>(), (n << vec3<u32>(offset)), (offset < 32u)) & vec3<u32>(mask)) | (v & vec3<u32>(~(mask))));
}
fn f() {
diff --git a/src/tint/transform/clamp_frag_depth_test.cc b/src/tint/transform/clamp_frag_depth_test.cc
index b94d5af..7c7019a 100644
--- a/src/tint/transform/clamp_frag_depth_test.cc
+++ b/src/tint/transform/clamp_frag_depth_test.cc
@@ -169,7 +169,7 @@
@fragment fn main() -> @builtin(frag_depth) f32 {
return 0.0;
}
- @fragment fn friend() -> @location(0) f32 {
+ @fragment fn other() -> @location(0) f32 {
return 0.0;
}
)";
@@ -194,7 +194,7 @@
}
@fragment
-fn friend() -> @location(0) f32 {
+fn other() -> @location(0) f32 {
return 0.0;
}
)";
diff --git a/src/tint/transform/multiplanar_external_texture_test.cc b/src/tint/transform/multiplanar_external_texture_test.cc
index 4416d35..8f2b014 100644
--- a/src/tint/transform/multiplanar_external_texture_test.cc
+++ b/src/tint/transform/multiplanar_external_texture_test.cc
@@ -580,9 +580,9 @@
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
- var signed = textureLoad(ext_tex, vec2<i32>(1));
- var unsigned = textureLoad(ext_tex, vec2<u32>(1));
- return signed + unsigned;
+ var val_signed = textureLoad(ext_tex, vec2<i32>(1));
+ var val_unsigned = textureLoad(ext_tex, vec2<u32>(1));
+ return val_signed + val_unsigned;
}
)";
@@ -652,9 +652,9 @@
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
- var signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
- var unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
- return (signed + unsigned);
+ var val_signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
+ var val_unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
+ return (val_signed + val_unsigned);
}
)";
@@ -670,9 +670,9 @@
auto* src = R"(
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
- var signed = textureLoad(ext_tex, vec2<i32>(1));
- var unsigned = textureLoad(ext_tex, vec2<u32>(1));
- return signed + unsigned;
+ var val_signed = textureLoad(ext_tex, vec2<i32>(1));
+ var val_unsigned = textureLoad(ext_tex, vec2<u32>(1));
+ return val_signed + val_unsigned;
}
@group(0) @binding(0) var ext_tex : texture_external;
@@ -742,9 +742,9 @@
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
- var signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
- var unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
- return (signed + unsigned);
+ var val_signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
+ var val_unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
+ return (val_signed + val_unsigned);
}
@group(0) @binding(0) var ext_tex : texture_2d<f32>;
diff --git a/src/tint/transform/robustness_test.cc b/src/tint/transform/robustness_test.cc
index 990bbde..8b9a8ad 100644
--- a/src/tint/transform/robustness_test.cc
+++ b/src/tint/transform/robustness_test.cc
@@ -787,7 +787,7 @@
@group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
@group(0) @binding(0) var tex_external : texture_external;
-fn signed() {
+fn idx_signed() {
var array_idx : i32;
var level_idx : i32;
var sample_idx : i32;
@@ -802,7 +802,7 @@
textureLoad(tex_external, vec2<i32>(1, 2));
}
-fn unsigned() {
+fn idx_unsigned() {
var array_idx : u32;
var level_idx : u32;
var sample_idx : u32;
@@ -836,7 +836,7 @@
@group(0) @binding(0) var tex_external : texture_external;
-fn signed() {
+fn idx_signed() {
var array_idx : i32;
var level_idx : i32;
var sample_idx : i32;
@@ -850,7 +850,7 @@
textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
}
-fn unsigned() {
+fn idx_unsigned() {
var array_idx : u32;
var level_idx : u32;
var sample_idx : u32;
@@ -873,7 +873,7 @@
// Clamp textureLoad() coord, array_index and level values
TEST_F(RobustnessTest, TextureLoad_Clamp_OutOfOrder) {
auto* src = R"(
-fn signed() {
+fn idx_signed() {
var array_idx : i32;
var level_idx : i32;
var sample_idx : i32;
@@ -888,7 +888,7 @@
textureLoad(tex_external, vec2<i32>(1, 2));
}
-fn unsigned() {
+fn idx_unsigned() {
var array_idx : u32;
var level_idx : u32;
var sample_idx : u32;
@@ -915,7 +915,7 @@
auto* expect =
R"(
-fn signed() {
+fn idx_signed() {
var array_idx : i32;
var level_idx : i32;
var sample_idx : i32;
@@ -929,7 +929,7 @@
textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
}
-fn unsigned() {
+fn idx_unsigned() {
var array_idx : u32;
var level_idx : u32;
var sample_idx : u32;
@@ -976,14 +976,14 @@
@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
-fn signed() {
+fn idx_signed() {
textureStore(tex1d, 10i, vec4<i32>());
textureStore(tex2d, vec2<i32>(10, 20), vec4<i32>());
textureStore(tex2d_arr, vec2<i32>(10, 20), 50i, vec4<i32>());
textureStore(tex3d, vec3<i32>(10, 20, 30), vec4<i32>());
}
-fn unsigned() {
+fn idx_unsigned() {
textureStore(tex1d, 10u, vec4<i32>());
textureStore(tex2d, vec2<u32>(10, 20), vec4<i32>());
textureStore(tex2d_arr, vec2<u32>(10, 20), 50u, vec4<i32>());
@@ -1000,14 +1000,14 @@
@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
-fn signed() {
+fn idx_signed() {
textureStore(tex1d, clamp(10i, 0, i32((u32(textureDimensions(tex1d)) - 1))), vec4<i32>());
textureStore(tex2d, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d)) - vec2(1)))), vec4<i32>());
textureStore(tex2d_arr, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1)))), clamp(50i, 0, i32((u32(textureNumLayers(tex2d_arr)) - 1))), vec4<i32>());
textureStore(tex3d, clamp(vec3<i32>(10, 20, 30), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex3d)) - vec3(1)))), vec4<i32>());
}
-fn unsigned() {
+fn idx_unsigned() {
textureStore(tex1d, min(10u, (u32(textureDimensions(tex1d)) - 1)), vec4<i32>());
textureStore(tex2d, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d)) - vec2(1))), vec4<i32>());
textureStore(tex2d_arr, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1))), min(50u, (u32(textureNumLayers(tex2d_arr)) - 1)), vec4<i32>());
@@ -1023,14 +1023,14 @@
// Clamp textureStore() coord, array_index and level values
TEST_F(RobustnessTest, TextureStore_Clamp_OutOfOrder) {
auto* src = R"(
-fn signed() {
+fn idx_signed() {
textureStore(tex1d, 10i, vec4<i32>());
textureStore(tex2d, vec2<i32>(10, 20), vec4<i32>());
textureStore(tex2d_arr, vec2<i32>(10, 20), 50i, vec4<i32>());
textureStore(tex3d, vec3<i32>(10, 20, 30), vec4<i32>());
}
-fn unsigned() {
+fn idx_unsigned() {
textureStore(tex1d, 10u, vec4<i32>());
textureStore(tex2d, vec2<u32>(10, 20), vec4<i32>());
textureStore(tex2d_arr, vec2<u32>(10, 20), 50u, vec4<i32>());
@@ -1048,14 +1048,14 @@
)";
auto* expect = R"(
-fn signed() {
+fn idx_signed() {
textureStore(tex1d, clamp(10i, 0, i32((u32(textureDimensions(tex1d)) - 1))), vec4<i32>());
textureStore(tex2d, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d)) - vec2(1)))), vec4<i32>());
textureStore(tex2d_arr, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1)))), clamp(50i, 0, i32((u32(textureNumLayers(tex2d_arr)) - 1))), vec4<i32>());
textureStore(tex3d, clamp(vec3<i32>(10, 20, 30), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex3d)) - vec3(1)))), vec4<i32>());
}
-fn unsigned() {
+fn idx_unsigned() {
textureStore(tex1d, min(10u, (u32(textureDimensions(tex1d)) - 1)), vec4<i32>());
textureStore(tex2d, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d)) - vec2(1))), vec4<i32>());
textureStore(tex2d_arr, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1))), min(50u, (u32(textureNumLayers(tex2d_arr)) - 1)), vec4<i32>());
diff --git a/src/tint/transform/single_entry_point_test.cc b/src/tint/transform/single_entry_point_test.cc
index 7451090..e9fa68f 100644
--- a/src/tint/transform/single_entry_point_test.cc
+++ b/src/tint/transform/single_entry_point_test.cc
@@ -240,32 +240,6 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(SingleEntryPointTest, WorkgroupSizeLetPreserved) {
- auto* src = R"(
-let size : i32 = 1;
-
-@compute @workgroup_size(size)
-fn main() {
-}
-)";
-
- auto* expect = R"(
-const size : i32 = 1;
-
-@compute @workgroup_size(size)
-fn main() {
-}
-)";
-
- SingleEntryPoint::Config cfg("main");
-
- DataMap data;
- data.Add<SingleEntryPoint::Config>(cfg);
- auto got = Run<SingleEntryPoint>(src, data);
-
- EXPECT_EQ(expect, str(got));
-}
-
TEST_F(SingleEntryPointTest, WorkgroupSizeConstPreserved) {
auto* src = R"(
const size : i32 = 1;
diff --git a/src/tint/transform/unshadow_test.cc b/src/tint/transform/unshadow_test.cc
index f5a8102..c731e76 100644
--- a/src/tint/transform/unshadow_test.cc
+++ b/src/tint/transform/unshadow_test.cc
@@ -384,82 +384,6 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(UnshadowTest, LocalShadowsGlobalLet) {
- auto* src = R"(
-let a : i32 = 1;
-
-fn X() {
- var a = (a == 123);
-}
-
-fn Y() {
- let a = (a == 321);
-}
-
-fn Z() {
- const a = 321;
-}
-)";
-
- auto* expect = R"(
-const a : i32 = 1;
-
-fn X() {
- var a_1 = (a == 123);
-}
-
-fn Y() {
- let a_2 = (a == 321);
-}
-
-fn Z() {
- const a_3 = 321;
-}
-)";
-
- auto got = Run<Unshadow>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(UnshadowTest, LocalShadowsGlobalLet_OutOfOrder) {
- auto* src = R"(
-fn X() {
- var a = (a == 123);
-}
-
-fn Y() {
- let a = (a == 321);
-}
-
-fn Z() {
- const a = 321;
-}
-
-let a : i32 = 1;
-)";
-
- auto* expect = R"(
-fn X() {
- var a_1 = (a == 123);
-}
-
-fn Y() {
- let a_2 = (a == 321);
-}
-
-fn Z() {
- const a_3 = 321;
-}
-
-const a : i32 = 1;
-)";
-
- auto got = Run<Unshadow>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
TEST_F(UnshadowTest, LocalShadowsGlobalConst) {
auto* src = R"(
const a : i32 = 1;
@@ -732,46 +656,6 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(UnshadowTest, ParamShadowsGlobalLet) {
- auto* src = R"(
-let a : i32 = 1;
-
-fn F(a : bool) {
-}
-)";
-
- auto* expect = R"(
-const a : i32 = 1;
-
-fn F(a_1 : bool) {
-}
-)";
-
- auto got = Run<Unshadow>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(UnshadowTest, ParamShadowsGlobalLet_OutOfOrder) {
- auto* src = R"(
-fn F(a : bool) {
-}
-
-let a : i32 = 1;
-)";
-
- auto* expect = R"(
-fn F(a_1 : bool) {
-}
-
-const a : i32 = 1;
-)";
-
- auto got = Run<Unshadow>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
TEST_F(UnshadowTest, ParamShadowsGlobalConst) {
auto* src = R"(
const a : i32 = 1;
