Import Tint changes from Dawn
Changes:
- 6488e26e6b3742a78e3918d8000a82e2d325c94c [tint][fuzz][wgsl] Add Raise fuzzer by Ben Clayton <bclayton@google.com>
- 50faf095a595281feb00d7cd05f225d46a9996fa [tint][fuzz][wgsl] Add PtrToRef fuzzer by Ben Clayton <bclayton@google.com>
- 2a8c10e51d79ec119dca7475137c28b6f7d5b4b7 [tint][fuzz][wgsl] Add RenameConflicts fuzzer by Ben Clayton <bclayton@google.com>
- 24aa34ee3fe4621a5be5cf0efeb3ced443f510f1 [tint][fuzz][wgsl] Add ValueToLet fuzzer by Ben Clayton <bclayton@google.com>
- d07a2002494ccc6ece8bbeb13814d3d98e4ab75c [spirv] Fix dynamic indexes into constant arrays by James Price <jrprice@google.com>
- 4f491bf8c370e4975672b3f5abd0f1de4a95f12e [ir] Validate block parameters by James Price <jrprice@google.com>
- 750f71d3234ac9d49e175d5e7558b052e6b7519b [tint][ir] Track owning block in BlockParam by James Price <jrprice@google.com>
- e5380b752f0f1f9d5950710af8cad0191fcd8122 [ir] Validate function parameters by James Price <jrprice@google.com>
- 21517e4b5ff7944caf09200658e7444f5232cccb [ir] Add source map for a function and its params by James Price <jrprice@google.com>
- e5fff2277474fbdc1b3d696c69346b088fc098aa [tint][ir] Track owning Function in FunctionParam by James Price <jrprice@google.com>
- b0e445c0a103a80d530c802a223a395e8782146a [tint][IRToProgram] Create phony assignment if a value is... by Ben Clayton <bclayton@google.com>
- 304e57a5b2863383ba037fb42bfd26162da8e37b Restore "[tint][fuzz] Enable AllowedFeatures::Everything()" by Ben Clayton <bclayton@google.com>
- 83f6d0d2d3f6f1af440347fffd4f184c32ca6e47 [tint][hlsl] Don't ICE if the PixelLocal transform is mis... by Ben Clayton <bclayton@google.com>
- a871f650869cecb5dadae4abb30fb28bc7ef5f76 [tint][ir] Correctly `enable chromium_internal_graphite;` by Ben Clayton <bclayton@google.com>
- f7c7b93ba92915bd75805e67681dc6310d453b0d [tint][ir] Ensure all enumerators are serialized by Ben Clayton <bclayton@google.com>
- ddd45561455cda0def30fcabf6748a89a44fd23c [tools][utils] Split BufferReader out to separate files by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: 6488e26e6b3742a78e3918d8000a82e2d325c94c
Change-Id: I4b4c0b7c6d46e996f6ee675f665976346d713e39
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/185801
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/cmd/fuzz/ir/fuzz.h b/src/tint/cmd/fuzz/ir/fuzz.h
index 533cac4..18ffdce 100644
--- a/src/tint/cmd/fuzz/ir/fuzz.h
+++ b/src/tint/cmd/fuzz/ir/fuzz.h
@@ -33,6 +33,7 @@
#include <tuple>
#include <utility>
+#include "src/tint/utils/bytes/buffer_reader.h"
#include "src/tint/utils/bytes/decoder.h"
#include "src/tint/utils/containers/slice.h"
#include "src/tint/utils/macros/static_init.h"
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.cmake b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
index d51bd4d..350a170 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.cmake
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
@@ -55,6 +55,7 @@
tint_lang_wgsl_program
tint_lang_wgsl_program_fuzz
tint_lang_wgsl_sem
+ tint_lang_wgsl_writer_raise_fuzz
tint_lang_wgsl_fuzz
tint_utils_bytes
tint_utils_cli
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.gn b/src/tint/cmd/fuzz/wgsl/BUILD.gn
index 389a98e..89d7ea9 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.gn
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.gn
@@ -98,6 +98,7 @@
"${tint_src_dir}/lang/wgsl/program",
"${tint_src_dir}/lang/wgsl/program:fuzz",
"${tint_src_dir}/lang/wgsl/sem",
+ "${tint_src_dir}/lang/wgsl/writer/raise:fuzz",
"${tint_src_dir}/utils/bytes",
"${tint_src_dir}/utils/cli",
"${tint_src_dir}/utils/containers",
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.cc b/src/tint/cmd/fuzz/wgsl/fuzz.cc
index 6e4d9b8..23bbe0f 100644
--- a/src/tint/cmd/fuzz/wgsl/fuzz.cc
+++ b/src/tint/cmd/fuzz/wgsl/fuzz.cc
@@ -30,6 +30,8 @@
#include <iostream>
#include <thread>
+#include "src/tint/lang/wgsl/common/allowed_features.h"
+#include "src/tint/lang/wgsl/reader/options.h"
#include "src/tint/lang/wgsl/reader/reader.h"
#include "src/tint/utils/containers/vector.h"
#include "src/tint/utils/macros/defer.h"
@@ -87,8 +89,14 @@
tint::Source::File file("test.wgsl", wgsl);
// Parse the WGSL program.
- auto program = tint::wgsl::reader::Parse(&file);
+ tint::wgsl::reader::Options parse_options;
+ parse_options.allowed_features = tint::wgsl::AllowedFeatures::Everything();
+ auto program = tint::wgsl::reader::Parse(&file, parse_options);
if (!program.IsValid()) {
+ if (options.verbose) {
+ std::cerr << "invalid WGSL program: " << std::endl
+ << program.Diagnostics() << std::endl;
+ }
return;
}
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.h b/src/tint/cmd/fuzz/wgsl/fuzz.h
index c01455b..e4bc7a0 100644
--- a/src/tint/cmd/fuzz/wgsl/fuzz.h
+++ b/src/tint/cmd/fuzz/wgsl/fuzz.h
@@ -33,6 +33,7 @@
#include <utility>
#include "src/tint/lang/wgsl/program/program.h"
+#include "src/tint/utils/bytes/buffer_reader.h"
#include "src/tint/utils/bytes/decoder.h"
#include "src/tint/utils/containers/slice.h"
#include "src/tint/utils/macros/static_init.h"
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index 825be15..5fbb84d 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -196,10 +196,13 @@
return Function::PipelineStage::kFragment;
case pb::PipelineStage::Vertex:
return Function::PipelineStage::kVertex;
- default:
- TINT_ICE() << "unhandled PipelineStage: " << stage;
- return Function::PipelineStage::kCompute;
+
+ case pb::PipelineStage::PipelineStage_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::PipelineStage::PipelineStage_INT_MAX_SENTINEL_DO_NOT_USE_:
+ break;
}
+ TINT_ICE() << "unhandled PipelineStage: " << stage;
+ return Function::PipelineStage::kCompute;
}
////////////////////////////////////////////////////////////////////////////
@@ -322,8 +325,7 @@
case pb::Instruction::KindCase::kUnreachable:
inst_out = CreateInstructionUnreachable(inst_in.unreachable());
break;
- default:
- TINT_UNIMPLEMENTED() << inst_in.kind_case();
+ case pb::Instruction::KindCase::KIND_NOT_SET:
break;
}
TINT_ASSERT_OR_RETURN_VALUE(inst_out, nullptr);
@@ -550,7 +552,7 @@
return CreateTypeExternalTexture(type_in.external_texture());
case pb::Type::KindCase::kSampler:
return CreateTypeSampler(type_in.sampler());
- default:
+ case pb::Type::KindCase::KIND_NOT_SET:
break;
}
TINT_ICE() << type_in.kind_case();
@@ -571,10 +573,13 @@
return mod_out_.Types().Get<f32>();
case pb::TypeBasic::f16:
return mod_out_.Types().Get<f16>();
- default:
- TINT_ICE() << "invalid TypeBasic: " << basic_in;
- return nullptr;
+
+ case pb::TypeBasic::TypeBasic_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::TypeBasic::TypeBasic_INT_MAX_SENTINEL_DO_NOT_USE_:
+ break;
}
+ TINT_ICE() << "invalid TypeBasic: " << basic_in;
+ return nullptr;
}
const type::Vector* CreateTypeVector(const pb::TypeVector& vector_in) {
@@ -712,10 +717,15 @@
case pb::Value::KindCase::kConstant:
value_out = b.Constant(ConstantValue(value_in.constant()));
break;
- default:
- TINT_ICE() << "invalid TypeDecl.kind: " << value_in.kind_case();
- return nullptr;
+ case pb::Value::KindCase::KIND_NOT_SET:
+ break;
}
+
+ if (!value_out) {
+ TINT_ICE() << "invalid TypeDecl.kind: " << value_in.kind_case();
+ return nullptr;
+ }
+
return value_out;
}
@@ -788,10 +798,11 @@
return CreateConstantComposite(value_in.composite());
case pb::ConstantValue::KindCase::kSplat:
return CreateConstantSplat(value_in.splat());
- default:
- TINT_ICE() << "invalid ConstantValue.kind: " << value_in.kind_case();
- return nullptr;
+ case pb::ConstantValue::KindCase::KIND_NOT_SET:
+ break;
}
+ TINT_ICE() << "invalid ConstantValue.kind: " << value_in.kind_case();
+ return nullptr;
}
const core::constant::Value* CreateConstantScalar(const pb::ConstantValueScalar& value_in) {
@@ -806,10 +817,11 @@
return b.ConstantValue(f32(value_in.f32()));
case pb::ConstantValueScalar::KindCase::kF16:
return b.ConstantValue(f16(value_in.f16()));
- default:
- TINT_ICE() << "invalid ConstantValueScalar.kind: " << value_in.kind_case();
- return nullptr;
+ case pb::ConstantValueScalar::KindCase::KIND_NOT_SET:
+ break;
}
+ TINT_ICE() << "invalid ConstantValueScalar.kind: " << value_in.kind_case();
+ return nullptr;
}
const core::constant::Value* CreateConstantComposite(
@@ -874,10 +886,13 @@
return core::AddressSpace::kUniform;
case pb::AddressSpace::workgroup:
return core::AddressSpace::kWorkgroup;
- default:
- TINT_ICE() << "invalid AddressSpace: " << in;
- return core::AddressSpace::kUndefined;
+
+ case pb::AddressSpace::AddressSpace_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::AddressSpace::AddressSpace_INT_MAX_SENTINEL_DO_NOT_USE_:
+ break;
}
+ TINT_ICE() << "invalid AddressSpace: " << in;
+ return core::AddressSpace::kUndefined;
}
core::Access AccessControl(pb::AccessControl in) {
@@ -888,10 +903,13 @@
return core::Access::kWrite;
case pb::AccessControl::read_write:
return core::Access::kReadWrite;
- default:
- TINT_ICE() << "invalid Access: " << in;
- return core::Access::kUndefined;
+
+ case pb::AccessControl::AccessControl_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::AccessControl::AccessControl_INT_MAX_SENTINEL_DO_NOT_USE_:
+ break;
}
+ TINT_ICE() << "invalid Access: " << in;
+ return core::Access::kUndefined;
}
core::UnaryOp UnaryOp(pb::UnaryOp in) {
@@ -907,10 +925,12 @@
case pb::UnaryOp::not_:
return core::UnaryOp::kNot;
- default:
- TINT_ICE() << "invalid UnaryOp: " << in;
- return core::UnaryOp::kComplement;
+ case pb::UnaryOp::UnaryOp_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::UnaryOp::UnaryOp_INT_MAX_SENTINEL_DO_NOT_USE_:
+ break;
}
+ TINT_ICE() << "invalid UnaryOp: " << in;
+ return core::UnaryOp::kComplement;
}
core::BinaryOp BinaryOp(pb::BinaryOp in) {
@@ -947,11 +967,17 @@
return core::BinaryOp::kShiftLeft;
case pb::BinaryOp::shift_right:
return core::BinaryOp::kShiftRight;
+ case pb::BinaryOp::logical_and:
+ return core::BinaryOp::kLogicalAnd;
+ case pb::BinaryOp::logical_or:
+ return core::BinaryOp::kLogicalOr;
- default:
- TINT_ICE() << "invalid BinaryOp: " << in;
- return core::BinaryOp::kAdd;
+ case pb::BinaryOp::BinaryOp_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::BinaryOp::BinaryOp_INT_MAX_SENTINEL_DO_NOT_USE_:
+ break;
}
+ TINT_ICE() << "invalid BinaryOp: " << in;
+ return core::BinaryOp::kAdd;
}
core::type::TextureDimension TextureDimension(pb::TextureDimension in) {
@@ -968,7 +994,9 @@
return core::type::TextureDimension::kCube;
case pb::TextureDimension::cube_array:
return core::type::TextureDimension::kCubeArray;
- default:
+
+ case pb::TextureDimension::TextureDimension_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::TextureDimension::TextureDimension_INT_MAX_SENTINEL_DO_NOT_USE_:
break;
}
@@ -980,6 +1008,8 @@
switch (in) {
case pb::TexelFormat::bgra8_unorm:
return core::TexelFormat::kBgra8Unorm;
+ case pb::TexelFormat::r8_unorm:
+ return core::TexelFormat::kR8Unorm;
case pb::TexelFormat::r32_float:
return core::TexelFormat::kR32Float;
case pb::TexelFormat::r32_sint:
@@ -1012,7 +1042,9 @@
return core::TexelFormat::kRgba8Uint;
case pb::TexelFormat::rgba8_unorm:
return core::TexelFormat::kRgba8Unorm;
- default:
+
+ case pb::TexelFormat::TexelFormat_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::TexelFormat::TexelFormat_INT_MAX_SENTINEL_DO_NOT_USE_:
break;
}
@@ -1026,7 +1058,9 @@
return core::type::SamplerKind::kSampler;
case pb::SamplerKind::comparison:
return core::type::SamplerKind::kComparisonSampler;
- default:
+
+ case pb::SamplerKind::SamplerKind_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::SamplerKind::SamplerKind_INT_MAX_SENTINEL_DO_NOT_USE_:
break;
}
@@ -1042,7 +1076,9 @@
return core::InterpolationType::kLinear;
case pb::InterpolationType::perspective:
return core::InterpolationType::kPerspective;
- default:
+
+ case pb::InterpolationType::InterpolationType_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::InterpolationType::InterpolationType_INT_MAX_SENTINEL_DO_NOT_USE_:
break;
}
TINT_ICE() << "invalid InterpolationType: " << in;
@@ -1057,7 +1093,9 @@
return core::InterpolationSampling::kCentroid;
case pb::InterpolationSampling::sample:
return core::InterpolationSampling::kSample;
- default:
+
+ case pb::InterpolationSampling::InterpolationSampling_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::InterpolationSampling::InterpolationSampling_INT_MAX_SENTINEL_DO_NOT_USE_:
break;
}
TINT_ICE() << "invalid InterpolationSampling: " << in;
@@ -1096,7 +1134,9 @@
return core::BuiltinValue::kVertexIndex;
case pb::BuiltinValue::workgroup_id:
return core::BuiltinValue::kWorkgroupId;
- default:
+
+ case pb::BuiltinValue::BuiltinValue_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::BuiltinValue::BuiltinValue_INT_MAX_SENTINEL_DO_NOT_USE_:
break;
}
TINT_ICE() << "invalid BuiltinValue: " << in;
@@ -1225,6 +1265,14 @@
return core::BuiltinFn::kPack4X8Snorm;
case pb::BuiltinFn::pack4x8_unorm:
return core::BuiltinFn::kPack4X8Unorm;
+ case pb::BuiltinFn::pack4xi8:
+ return core::BuiltinFn::kPack4XI8;
+ case pb::BuiltinFn::pack4xu8:
+ return core::BuiltinFn::kPack4XU8;
+ case pb::BuiltinFn::pack4xi8_clamp:
+ return core::BuiltinFn::kPack4XI8Clamp;
+ case pb::BuiltinFn::pack4xu8_clamp:
+ return core::BuiltinFn::kPack4XU8Clamp;
case pb::BuiltinFn::pow:
return core::BuiltinFn::kPow;
case pb::BuiltinFn::quantize_to_f16:
@@ -1275,6 +1323,10 @@
return core::BuiltinFn::kUnpack4X8Snorm;
case pb::BuiltinFn::unpack4x8_unorm:
return core::BuiltinFn::kUnpack4X8Unorm;
+ case pb::BuiltinFn::unpack4xi8:
+ return core::BuiltinFn::kUnpack4XI8;
+ case pb::BuiltinFn::unpack4xu8:
+ return core::BuiltinFn::kUnpack4XU8;
case pb::BuiltinFn::workgroup_barrier:
return core::BuiltinFn::kWorkgroupBarrier;
case pb::BuiltinFn::texture_barrier:
@@ -1335,7 +1387,9 @@
return core::BuiltinFn::kSubgroupBallot;
case pb::BuiltinFn::subgroup_broadcast:
return core::BuiltinFn::kSubgroupBroadcast;
- default:
+
+ case pb::BuiltinFn::BuiltinFn_INT_MIN_SENTINEL_DO_NOT_USE_:
+ case pb::BuiltinFn::BuiltinFn_INT_MAX_SENTINEL_DO_NOT_USE_:
break;
}
TINT_ICE() << "invalid BuiltinFn: " << in;
diff --git a/src/tint/lang/core/ir/binary/encode.cc b/src/tint/lang/core/ir/binary/encode.cc
index ef364ce..562a568 100644
--- a/src/tint/lang/core/ir/binary/encode.cc
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -162,10 +162,11 @@
return pb::PipelineStage::Fragment;
case Function::PipelineStage::kVertex:
return pb::PipelineStage::Vertex;
- default:
- TINT_ICE() << "unhandled PipelineStage: " << stage;
- return pb::PipelineStage::Compute;
+ case Function::PipelineStage::kUndefined:
+ break;
}
+ TINT_ICE() << "unhandled PipelineStage: " << stage;
+ return pb::PipelineStage::Compute;
}
////////////////////////////////////////////////////////////////////////////
@@ -655,10 +656,14 @@
return pb::AddressSpace::uniform;
case core::AddressSpace::kWorkgroup:
return pb::AddressSpace::workgroup;
- default:
- TINT_ICE() << "invalid AddressSpace: " << in;
- return pb::AddressSpace::function;
+
+ case core::AddressSpace::kUndefined:
+ case core::AddressSpace::kIn:
+ case core::AddressSpace::kOut:
+ break;
}
+ TINT_ICE() << "invalid AddressSpace: " << in;
+ return pb::AddressSpace::function;
}
pb::AccessControl AccessControl(core::Access in) {
@@ -669,10 +674,11 @@
return pb::AccessControl::write;
case core::Access::kReadWrite:
return pb::AccessControl::read_write;
- default:
- TINT_ICE() << "invalid Access: " << in;
- return pb::AccessControl::read;
+ case core::Access::kUndefined:
+ break;
}
+ TINT_ICE() << "invalid Access: " << in;
+ return pb::AccessControl::read;
}
pb::UnaryOp UnaryOp(core::UnaryOp in) {
@@ -750,7 +756,7 @@
return pb::TextureDimension::cube;
case core::type::TextureDimension::kCubeArray:
return pb::TextureDimension::cube_array;
- default:
+ case core::type::TextureDimension::kNone:
break;
}
@@ -768,6 +774,8 @@
return pb::TexelFormat::r32_sint;
case core::TexelFormat::kR32Uint:
return pb::TexelFormat::r32_uint;
+ case core::TexelFormat::kR8Unorm:
+ return pb::TexelFormat::r8_unorm;
case core::TexelFormat::kRg32Float:
return pb::TexelFormat::rg32_float;
case core::TexelFormat::kRg32Sint:
@@ -794,7 +802,7 @@
return pb::TexelFormat::rgba8_uint;
case core::TexelFormat::kRgba8Unorm:
return pb::TexelFormat::rgba8_unorm;
- default:
+ case core::TexelFormat::kUndefined:
break;
}
@@ -822,7 +830,7 @@
return pb::InterpolationType::linear;
case core::InterpolationType::kPerspective:
return pb::InterpolationType::perspective;
- default:
+ case core::InterpolationType::kUndefined:
break;
}
TINT_ICE() << "invalid InterpolationType: " << in;
@@ -837,7 +845,7 @@
return pb::InterpolationSampling::centroid;
case core::InterpolationSampling::kSample:
return pb::InterpolationSampling::sample;
- default:
+ case core::InterpolationSampling::kUndefined:
break;
}
TINT_ICE() << "invalid InterpolationSampling: " << in;
@@ -876,7 +884,7 @@
return pb::BuiltinValue::vertex_index;
case core::BuiltinValue::kWorkgroupId:
return pb::BuiltinValue::workgroup_id;
- default:
+ case core::BuiltinValue::kUndefined:
break;
}
TINT_ICE() << "invalid BuiltinValue: " << in;
@@ -1005,6 +1013,14 @@
return pb::BuiltinFn::pack4x8_snorm;
case core::BuiltinFn::kPack4X8Unorm:
return pb::BuiltinFn::pack4x8_unorm;
+ case core::BuiltinFn::kPack4XI8:
+ return pb::BuiltinFn::pack4xi8;
+ case core::BuiltinFn::kPack4XU8:
+ return pb::BuiltinFn::pack4xu8;
+ case core::BuiltinFn::kPack4XI8Clamp:
+ return pb::BuiltinFn::pack4xi8_clamp;
+ case core::BuiltinFn::kPack4XU8Clamp:
+ return pb::BuiltinFn::pack4xu8_clamp;
case core::BuiltinFn::kPow:
return pb::BuiltinFn::pow;
case core::BuiltinFn::kQuantizeToF16:
@@ -1055,6 +1071,10 @@
return pb::BuiltinFn::unpack4x8_snorm;
case core::BuiltinFn::kUnpack4X8Unorm:
return pb::BuiltinFn::unpack4x8_unorm;
+ case core::BuiltinFn::kUnpack4XI8:
+ return pb::BuiltinFn::unpack4xi8;
+ case core::BuiltinFn::kUnpack4XU8:
+ return pb::BuiltinFn::unpack4xu8;
case core::BuiltinFn::kWorkgroupBarrier:
return pb::BuiltinFn::workgroup_barrier;
case core::BuiltinFn::kTextureBarrier:
@@ -1115,7 +1135,7 @@
return pb::BuiltinFn::subgroup_ballot;
case core::BuiltinFn::kSubgroupBroadcast:
return pb::BuiltinFn::subgroup_broadcast;
- default:
+ case core::BuiltinFn::kNone:
break;
}
TINT_ICE() << "invalid BuiltinFn: " << in;
diff --git a/src/tint/lang/core/ir/binary/ir.proto b/src/tint/lang/core/ir/binary/ir.proto
index e59a9cb..19f3b0a 100644
--- a/src/tint/lang/core/ir/binary/ir.proto
+++ b/src/tint/lang/core/ir/binary/ir.proto
@@ -247,31 +247,30 @@
InstructionReturn return = 3;
InstructionUnary unary = 4;
InstructionBinary binary = 5;
- InstructionBuiltin builtin = 6;
- InstructionDiscard discard = 7;
- InstructionLet let = 8;
- InstructionVar var = 9;
- InstructionBitcast bitcast = 10;
- InstructionConstruct construct = 11;
- InstructionConvert convert = 12;
- InstructionAccess access = 13;
- InstructionUserCall user_call = 14;
- InstructionBuiltinCall builtin_call = 15;
- InstructionLoad load = 16;
- InstructionStore store = 17;
- InstructionLoadVectorElement load_vector_element = 18;
- InstructionStoreVectorElement store_vector_element = 19;
- InstructionSwizzle swizzle = 20;
- InstructionIf if = 21;
- InstructionSwitch switch = 22;
- InstructionLoop loop = 23;
- InstructionExitIf exit_if = 24;
- InstructionExitSwitch exit_switch = 25;
- InstructionExitLoop exit_loop = 26;
- InstructionNextIteration next_iteration = 27;
- InstructionContinue continue = 28;
- InstructionBreakIf break_if = 29;
- InstructionUnreachable unreachable = 30;
+ InstructionDiscard discard = 6;
+ InstructionLet let = 7;
+ InstructionVar var = 8;
+ InstructionBitcast bitcast = 9;
+ InstructionConstruct construct = 10;
+ InstructionConvert convert = 11;
+ InstructionAccess access = 12;
+ InstructionUserCall user_call = 13;
+ InstructionBuiltinCall builtin_call = 14;
+ InstructionLoad load = 15;
+ InstructionStore store = 16;
+ InstructionLoadVectorElement load_vector_element = 17;
+ InstructionStoreVectorElement store_vector_element = 18;
+ InstructionSwizzle swizzle = 19;
+ InstructionIf if = 20;
+ InstructionSwitch switch = 21;
+ InstructionLoop loop = 22;
+ InstructionExitIf exit_if = 23;
+ InstructionExitSwitch exit_switch = 24;
+ InstructionExitLoop exit_loop = 25;
+ InstructionNextIteration next_iteration = 26;
+ InstructionContinue continue = 27;
+ InstructionBreakIf break_if = 28;
+ InstructionUnreachable unreachable = 29;
}
}
@@ -285,8 +284,6 @@
BinaryOp op = 1;
}
-message InstructionBuiltin {}
-
message InstructionBitcast {}
message InstructionConstruct {}
@@ -450,22 +447,23 @@
enum TexelFormat {
bgra8_unorm = 0;
- r32_float = 1;
- r32_sint = 2;
- r32_uint = 3;
- rg32_float = 4;
- rg32_sint = 5;
- rg32_uint = 6;
- rgba16_float = 7;
- rgba16_sint = 8;
- rgba16_uint = 9;
- rgba32_float = 10;
- rgba32_sint = 11;
- rgba32_uint = 12;
- rgba8_sint = 13;
- rgba8_snorm = 14;
- rgba8_uint = 15;
- rgba8_unorm = 16;
+ r8_unorm = 1;
+ r32_float = 2;
+ r32_sint = 3;
+ r32_uint = 4;
+ rg32_float = 5;
+ rg32_sint = 6;
+ rg32_uint = 7;
+ rgba16_float = 8;
+ rgba16_sint = 9;
+ rgba16_uint = 10;
+ rgba32_float = 11;
+ rgba32_sint = 12;
+ rgba32_uint = 13;
+ rgba8_sint = 14;
+ rgba8_snorm = 15;
+ rgba8_uint = 16;
+ rgba8_unorm = 17;
}
enum SamplerKind {
@@ -564,59 +562,65 @@
pack2x16_unorm = 57;
pack4x8_snorm = 58;
pack4x8_unorm = 59;
- pow = 60;
- quantize_to_f16 = 61;
- radians = 62;
- reflect = 63;
- refract = 64;
- reverse_bits = 65;
- round = 66;
- saturate = 67;
- select = 68;
- sign = 69;
- sin = 70;
- sinh = 71;
- smoothstep = 72;
- sqrt = 73;
- step = 74;
- storage_barrier = 75;
- tan = 76;
- tanh = 77;
- transpose = 78;
- trunc = 79;
- unpack2x16_float = 80;
- unpack2x16_snorm = 81;
- unpack2x16_unorm = 82;
- unpack4x8_snorm = 83;
- unpack4x8_unorm = 84;
- workgroup_barrier = 85;
- texture_barrier = 86;
- texture_dimensions = 87;
- texture_gather = 88;
- texture_gather_compare = 89;
- texture_num_layers = 90;
- texture_num_levels = 91;
- texture_num_samples = 92;
- texture_sample = 93;
- texture_sample_bias = 94;
- texture_sample_compare = 95;
- texture_sample_compare_level = 96;
- texture_sample_grad = 97;
- texture_sample_level = 98;
- texture_sample_base_clamp_to_edge = 99;
- texture_store = 100;
- texture_load = 101;
- atomic_load = 102;
- atomic_store = 103;
- atomic_add = 104;
- atomic_sub = 105;
- atomic_max = 106;
- atomic_min = 107;
- atomic_and = 108;
- atomic_or = 109;
- atomic_xor = 110;
- atomic_exchange = 111;
- atomic_compare_exchange_weak = 112;
- subgroup_ballot = 113;
- subgroup_broadcast = 114;
+ pack4xi8 = 60;
+ pack4xu8 = 61;
+ pack4xi8_clamp = 62;
+ pack4xu8_clamp = 63;
+ pow = 64;
+ quantize_to_f16 = 65;
+ radians = 66;
+ reflect = 67;
+ refract = 68;
+ reverse_bits = 69;
+ round = 70;
+ saturate = 71;
+ select = 72;
+ sign = 73;
+ sin = 74;
+ sinh = 75;
+ smoothstep = 76;
+ sqrt = 77;
+ step = 78;
+ storage_barrier = 79;
+ tan = 80;
+ tanh = 81;
+ transpose = 82;
+ trunc = 83;
+ unpack2x16_float = 84;
+ unpack2x16_snorm = 85;
+ unpack2x16_unorm = 86;
+ unpack4x8_snorm = 87;
+ unpack4x8_unorm = 88;
+ unpack4xi8 = 89;
+ unpack4xu8 = 90;
+ workgroup_barrier = 91;
+ texture_barrier = 92;
+ texture_dimensions = 93;
+ texture_gather = 94;
+ texture_gather_compare = 95;
+ texture_num_layers = 96;
+ texture_num_levels = 97;
+ texture_num_samples = 98;
+ texture_sample = 99;
+ texture_sample_bias = 100;
+ texture_sample_compare = 101;
+ texture_sample_compare_level = 102;
+ texture_sample_grad = 103;
+ texture_sample_level = 104;
+ texture_sample_base_clamp_to_edge = 105;
+ texture_store = 106;
+ texture_load = 107;
+ atomic_load = 108;
+ atomic_store = 109;
+ atomic_add = 110;
+ atomic_sub = 111;
+ atomic_max = 112;
+ atomic_min = 113;
+ atomic_and = 114;
+ atomic_or = 115;
+ atomic_xor = 116;
+ atomic_exchange = 117;
+ atomic_compare_exchange_weak = 118;
+ subgroup_ballot = 119;
+ subgroup_broadcast = 120;
}
diff --git a/src/tint/lang/core/ir/block_param.h b/src/tint/lang/core/ir/block_param.h
index 009e022..0b697a2 100644
--- a/src/tint/lang/core/ir/block_param.h
+++ b/src/tint/lang/core/ir/block_param.h
@@ -31,25 +31,42 @@
#include "src/tint/lang/core/ir/value.h"
#include "src/tint/utils/rtti/castable.h"
+// Forward declarations
+namespace tint::core::ir {
+class MultiInBlock;
+} // namespace tint::core::ir
+
namespace tint::core::ir {
-/// An instruction in the IR.
+/// A block parameter in the IR.
class BlockParam : public Castable<BlockParam, Value> {
public:
/// Constructor
- /// @param type the type of the var
+ /// @param type the type of the parameter
explicit BlockParam(const core::type::Type* type);
~BlockParam() override;
- /// @returns the type of the var
+ /// @returns the type of the parameter
const core::type::Type* Type() const override { return type_; }
+ /// Sets the block that this parameter belongs to.
+ /// @param block the block
+ void SetBlock(MultiInBlock* block) { block_ = block; }
+
+ /// @returns the block that this parameter belongs to, or nullptr
+ MultiInBlock* Block() { return block_; }
+
+ /// @returns the block that this parameter belongs to, or nullptr
+ const MultiInBlock* Block() const { return block_; }
+
/// @copydoc Instruction::Clone()
BlockParam* Clone(CloneContext& ctx) override;
private:
- /// the result type of the instruction
+ /// the type of the parameter
const core::type::Type* type_ = nullptr;
+ /// the block that the parameter belongs to
+ MultiInBlock* block_ = nullptr;
};
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index c705d0c..99e94d1 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -219,6 +219,37 @@
cb();
}
+ /// Calls @p cb with the builder inserting after @p val
+ /// @param val the insertion point for new instructions
+ /// @param cb the function to call with the builder inserting new instructions after @p val
+ template <typename FUNCTION>
+ void InsertAfter(ir::Value* val, FUNCTION&& cb) {
+ tint::Switch(
+ val,
+ [&](core::ir::InstructionResult* result) {
+ const TINT_SCOPED_ASSIGNMENT(insertion_point_,
+ InsertionPoints::InsertAfter{result->Instruction()});
+ cb();
+ },
+ [&](core::ir::FunctionParam* param) {
+ auto* body = param->Function()->Block();
+ if (body->IsEmpty()) {
+ Append(body, cb);
+ } else {
+ InsertBefore(body->Front(), cb);
+ }
+ },
+ [&](core::ir::BlockParam* param) {
+ auto* block = param->Block();
+ if (block->IsEmpty()) {
+ Append(block, cb);
+ } else {
+ InsertBefore(block->Front(), cb);
+ }
+ },
+ TINT_ICE_ON_NO_MATCH);
+ }
+
/// Adds and returns the instruction @p instruction to the current insertion point. If there
/// is no current insertion point set, then @p instruction is just returned.
/// @param instruction the instruction to append
diff --git a/src/tint/lang/core/ir/disassembler.cc b/src/tint/lang/core/ir/disassembler.cc
index 2e95c2c..281cb61 100644
--- a/src/tint/lang/core/ir/disassembler.cc
+++ b/src/tint/lang/core/ir/disassembler.cc
@@ -188,7 +188,14 @@
if (auto* merge = blk->As<MultiInBlock>()) {
if (!merge->Params().IsEmpty()) {
out_ << " (";
- EmitValueList(merge->Params().Slice());
+ for (auto* p : merge->Params()) {
+ if (p != merge->Params().Front()) {
+ out_ << ", ";
+ }
+ SourceMarker psm(this);
+ EmitValue(p);
+ psm.Store(p);
+ }
out_ << ")";
}
}
@@ -302,7 +309,12 @@
in_function_ = true;
std::string fn_id = IdOf(func);
- Indent() << "%" << fn_id << " =";
+ {
+ SourceMarker sm(this);
+ Indent() << "%" << fn_id;
+ sm.Store(func);
+ }
+ out_ << " =";
if (func->Stage() != Function::PipelineStage::kUndefined) {
out_ << " @" << func->Stage();
@@ -318,7 +330,9 @@
if (p != func->Params().Front()) {
out_ << ", ";
}
+ SourceMarker sm(this);
out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+ sm.Store(p);
EmitParamAttributes(p);
}
@@ -790,15 +804,6 @@
);
}
-void Disassembler::EmitValueList(tint::Slice<const Value* const> values) {
- for (size_t i = 0, n = values.Length(); i < n; i++) {
- if (i > 0) {
- out_ << ", ";
- }
- EmitValue(values[i]);
- }
-}
-
void Disassembler::EmitBinary(const Binary* b) {
SourceMarker sm(this);
EmitValueWithType(b);
diff --git a/src/tint/lang/core/ir/disassembler.h b/src/tint/lang/core/ir/disassembler.h
index d22b8a1..6b6f7c6 100644
--- a/src/tint/lang/core/ir/disassembler.h
+++ b/src/tint/lang/core/ir/disassembler.h
@@ -32,6 +32,7 @@
#include "src/tint/lang/core/ir/binary.h"
#include "src/tint/lang/core/ir/block.h"
+#include "src/tint/lang/core/ir/block_param.h"
#include "src/tint/lang/core/ir/call.h"
#include "src/tint/lang/core/ir/if.h"
#include "src/tint/lang/core/ir/loop.h"
@@ -100,10 +101,26 @@
/// @returns the source for the result
Source ResultSource(IndexedValue result) { return result_to_src_.GetOr(result, Source{}); }
- /// @param blk teh block to retrieve
+ /// @param blk the block to retrieve
/// @returns the source for the block
Source BlockSource(const Block* blk) { return block_to_src_.GetOr(blk, Source{}); }
+ /// @param param the block parameter to retrieve
+ /// @returns the source for the parameter
+ Source BlockParamSource(const BlockParam* param) {
+ return block_param_to_src_.GetOr(param, Source{});
+ }
+
+ /// @param func the function to retrieve
+ /// @returns the source for the function
+ Source FunctionSource(const Function* func) { return function_to_src_.GetOr(func, Source{}); }
+
+ /// @param param the function parameter to retrieve
+ /// @returns the source for the parameter
+ Source FunctionParamSource(const FunctionParam* param) {
+ return function_param_to_src_.GetOr(param, Source{});
+ }
+
/// Stores the given @p src location for @p inst instruction
/// @param inst the instruction to store
/// @param src the source location
@@ -114,6 +131,23 @@
/// @param src the source location
void SetSource(const Block* blk, Source src) { block_to_src_.Add(blk, src); }
+ /// Stores the given @p src location for @p param block parameter
+ /// @param param the block parameter to store
+ /// @param src the source location
+ void SetSource(const BlockParam* param, Source src) { block_param_to_src_.Add(param, src); }
+
+ /// Stores the given @p src location for @p func function
+ /// @param func the function to store
+ /// @param src the source location
+ void SetSource(const Function* func, Source src) { function_to_src_.Add(func, src); }
+
+ /// Stores the given @p src location for @p param function parameter
+ /// @param param the function parameter to store
+ /// @param src the source location
+ void SetSource(const FunctionParam* param, Source src) {
+ function_param_to_src_.Add(param, src);
+ }
+
/// Stores the given @p src location for @p op operand
/// @param op the operand to store
/// @param src the source location
@@ -137,6 +171,12 @@
void Store(const Block* blk) { dis_->SetSource(blk, MakeSource()); }
+ void Store(const BlockParam* param) { dis_->SetSource(param, MakeSource()); }
+
+ void Store(const Function* func) { dis_->SetSource(func, MakeSource()); }
+
+ void Store(const FunctionParam* param) { dis_->SetSource(param, MakeSource()); }
+
void Store(IndexedValue operand) { dis_->SetSource(operand, MakeSource()); }
void StoreResult(IndexedValue result) { dis_->SetResultSource(result, MakeSource()); }
@@ -168,7 +208,6 @@
void EmitValueWithType(const Instruction* val);
void EmitValueWithType(const Value* val);
void EmitValue(const Value* val);
- void EmitValueList(tint::Slice<const ir::Value* const> values);
void EmitBinary(const Binary* b);
void EmitUnary(const Unary* b);
void EmitTerminator(const Terminator* b);
@@ -193,9 +232,12 @@
uint32_t current_output_start_pos_ = 0;
Hashmap<const Block*, Source, 8> block_to_src_;
+ Hashmap<const BlockParam*, Source, 8> block_param_to_src_;
Hashmap<const Instruction*, Source, 8> instruction_to_src_;
Hashmap<IndexedValue, Source, 8> operand_to_src_;
Hashmap<IndexedValue, Source, 8> result_to_src_;
+ Hashmap<const Function*, Source, 8> function_to_src_;
+ Hashmap<const FunctionParam*, Source, 8> function_param_to_src_;
Hashmap<const If*, std::string, 8> if_names_;
Hashmap<const Loop*, std::string, 8> loop_names_;
Hashmap<const Switch*, std::string, 8> switch_names_;
diff --git a/src/tint/lang/core/ir/function.cc b/src/tint/lang/core/ir/function.cc
index 51f7928..5b669c2 100644
--- a/src/tint/lang/core/ir/function.cc
+++ b/src/tint/lang/core/ir/function.cc
@@ -53,7 +53,7 @@
auto* new_func =
ctx.ir.allocators.values.Create<Function>(return_.type, pipeline_stage_, workgroup_size_);
new_func->block_ = ctx.ir.blocks.Create<ir::Block>();
- new_func->params_ = ctx.Clone<1>(params_.Slice());
+ new_func->SetParams(ctx.Clone<1>(params_.Slice()));
new_func->return_.builtin = return_.builtin;
new_func->return_.location = return_.location;
new_func->return_.invariant = return_.invariant;
@@ -66,13 +66,25 @@
}
void Function::SetParams(VectorRef<FunctionParam*> params) {
+ for (auto* param : params_) {
+ param->SetFunction(nullptr);
+ }
params_ = std::move(params);
- TINT_ASSERT(!params_.Any(IsNull));
+ TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+ for (auto* param : params_) {
+ param->SetFunction(this);
+ }
}
void Function::SetParams(std::initializer_list<FunctionParam*> params) {
+ for (auto* param : params_) {
+ param->SetFunction(nullptr);
+ }
params_ = params;
- TINT_ASSERT(!params_.Any(IsNull));
+ TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+ for (auto* param : params_) {
+ param->SetFunction(this);
+ }
}
void Function::Destroy() {
diff --git a/src/tint/lang/core/ir/function_param.h b/src/tint/lang/core/ir/function_param.h
index 6cefb10..f323914 100644
--- a/src/tint/lang/core/ir/function_param.h
+++ b/src/tint/lang/core/ir/function_param.h
@@ -38,6 +38,11 @@
#include "src/tint/utils/ice/ice.h"
#include "src/tint/utils/rtti/castable.h"
+// Forward declarations
+namespace tint::core::ir {
+class Function;
+} // namespace tint::core::ir
+
namespace tint::core::ir {
/// A function parameter in the IR.
@@ -48,6 +53,16 @@
explicit FunctionParam(const core::type::Type* type);
~FunctionParam() override;
+ /// Sets the function that this parameter belongs to.
+ /// @param func the function
+ void SetFunction(ir::Function* func) { func_ = func; }
+
+ /// @returns the function that this parameter belongs to, or nullptr
+ ir::Function* Function() { return func_; }
+
+ /// @returns the function that this parameter belongs to, or nullptr
+ const ir::Function* Function() const { return func_; }
+
/// @returns the type of the var
const core::type::Type* Type() const override { return type_; }
@@ -99,6 +114,7 @@
std::optional<struct BindingPoint> BindingPoint() const { return binding_point_; }
private:
+ ir::Function* func_ = nullptr;
const core::type::Type* type_ = nullptr;
std::optional<core::BuiltinValue> builtin_;
std::optional<struct Location> location_;
diff --git a/src/tint/lang/core/ir/function_test.cc b/src/tint/lang/core/ir/function_test.cc
index 8873648..27f332e 100644
--- a/src/tint/lang/core/ir/function_test.cc
+++ b/src/tint/lang/core/ir/function_test.cc
@@ -126,6 +126,12 @@
// Cloned functions are not automatically added to the module.
EXPECT_EQ(mod.functions.Length(), 1u);
+
+ // Check parameter ownership is correct.
+ EXPECT_EQ(param1->Function(), f);
+ EXPECT_EQ(param2->Function(), f);
+ EXPECT_EQ(new_param1->Function(), new_f);
+ EXPECT_EQ(new_param2->Function(), new_f);
}
TEST_F(IR_FunctionTest, CloneWithExits) {
@@ -138,5 +144,23 @@
EXPECT_EQ(new_f, new_f->Block()->Front()->As<Return>()->Func());
}
+TEST_F(IR_FunctionTest, Parameters) {
+ auto* f = b.Function("my_func", mod.Types().void_());
+
+ auto* param1 = b.FunctionParam("a", mod.Types().i32());
+ auto* param2 = b.FunctionParam("b", mod.Types().f32());
+ auto* param3 = b.FunctionParam("b", mod.Types().f32());
+
+ f->SetParams({param1, param2});
+ EXPECT_EQ(param1->Function(), f);
+ EXPECT_EQ(param2->Function(), f);
+ EXPECT_EQ(param3->Function(), nullptr);
+
+ f->SetParams({param1, param3});
+ EXPECT_EQ(param1->Function(), f);
+ EXPECT_EQ(param2->Function(), nullptr);
+ EXPECT_EQ(param3->Function(), f);
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/multi_in_block.cc b/src/tint/lang/core/ir/multi_in_block.cc
index c57c622..411d06c 100644
--- a/src/tint/lang/core/ir/multi_in_block.cc
+++ b/src/tint/lang/core/ir/multi_in_block.cc
@@ -55,11 +55,25 @@
}
void MultiInBlock::SetParams(VectorRef<BlockParam*> params) {
+ for (auto* param : params_) {
+ param->SetBlock(nullptr);
+ }
params_ = std::move(params);
+ TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+ for (auto* param : params_) {
+ param->SetBlock(this);
+ }
}
void MultiInBlock::SetParams(std::initializer_list<BlockParam*> params) {
+ for (auto* param : params_) {
+ param->SetBlock(nullptr);
+ }
params_ = std::move(params);
+ TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+ for (auto* param : params_) {
+ param->SetBlock(this);
+ }
}
void MultiInBlock::AddInboundSiblingBranch(ir::Terminator* node) {
diff --git a/src/tint/lang/core/ir/multi_in_block_test.cc b/src/tint/lang/core/ir/multi_in_block_test.cc
index 790a90e..a377c44 100644
--- a/src/tint/lang/core/ir/multi_in_block_test.cc
+++ b/src/tint/lang/core/ir/multi_in_block_test.cc
@@ -54,7 +54,9 @@
auto* blk = b.MultiInBlock();
auto* add = b.Add(mod.Types().i32(), 1_i, 2_i);
blk->Append(add);
- blk->SetParams({b.BlockParam(mod.Types().i32()), b.BlockParam(mod.Types().f32())});
+ auto* param1 = b.BlockParam(mod.Types().i32());
+ auto* param2 = b.BlockParam(mod.Types().f32());
+ blk->SetParams({param1, param2});
blk->SetParent(loop);
auto* terminate = b.TerminateInvocation();
@@ -75,6 +77,12 @@
EXPECT_NE(add, new_blk->Front());
EXPECT_TRUE(new_blk->Front()->Is<Binary>());
EXPECT_EQ(BinaryOp::kAdd, new_blk->Front()->As<Binary>()->Op());
+
+ // Check parameter ownership is correct.
+ EXPECT_EQ(param1->Block(), blk);
+ EXPECT_EQ(param2->Block(), blk);
+ EXPECT_EQ(new_blk->Params()[0]->Block(), new_blk);
+ EXPECT_EQ(new_blk->Params()[1]->Block(), new_blk);
}
TEST_F(IR_MultiInBlockTest, CloneEmpty) {
@@ -86,5 +94,23 @@
EXPECT_EQ(0u, new_blk->Params().Length());
}
+TEST_F(IR_MultiInBlockTest, Parameters) {
+ auto* blk = b.MultiInBlock();
+
+ auto* param1 = b.BlockParam("a", mod.Types().i32());
+ auto* param2 = b.BlockParam("b", mod.Types().f32());
+ auto* param3 = b.BlockParam("b", mod.Types().f32());
+
+ blk->SetParams({param1, param2});
+ EXPECT_EQ(param1->Block(), blk);
+ EXPECT_EQ(param2->Block(), blk);
+ EXPECT_EQ(param3->Block(), nullptr);
+
+ blk->SetParams({param1, param3});
+ EXPECT_EQ(param1->Block(), blk);
+ EXPECT_EQ(param2->Block(), nullptr);
+ EXPECT_EQ(param3->Block(), blk);
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 3f1a681..24d594d 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -158,11 +158,26 @@
/// @returns the diagnostic
diag::Diagnostic& AddResultError(const Instruction* inst, size_t idx);
- /// Adds an error the @p block and highlights the block header in the disassembly
+ /// Adds an error for the @p block and highlights the block header in the disassembly
/// @param blk the block
/// @returns the diagnostic
diag::Diagnostic& AddError(const Block* blk);
+ /// Adds an error for the @p param and highlights the parameter in the disassembly
+ /// @param param the parameter
+ /// @returns the diagnostic
+ diag::Diagnostic& AddError(const BlockParam* param);
+
+ /// Adds an error for the @p func and highlights the function in the disassembly
+ /// @param func the function
+ /// @returns the diagnostic
+ diag::Diagnostic& AddError(const Function* func);
+
+ /// Adds an error for the @p param and highlights the parameter in the disassembly
+ /// @param param the parameter
+ /// @returns the diagnostic
+ diag::Diagnostic& AddError(const FunctionParam* param);
+
/// Adds an error the @p block and highlights the block header in the disassembly
/// @param src the source lines to highlight
/// @returns the diagnostic
@@ -172,6 +187,10 @@
/// @param inst the instruction
diag::Diagnostic& AddNote(const Instruction* inst);
+ /// Adds a note to @p func and highlights the function in the disassembly
+ /// @param func the function
+ diag::Diagnostic& AddNote(const Function* func);
+
/// Adds a note to @p inst for operand @p idx and highlights the operand in the
/// disassembly
/// @param inst the instruction
@@ -346,8 +365,8 @@
for (auto& func : mod_.functions) {
if (!all_functions_.Add(func.Get())) {
- AddError(Source{}) << "function " << style::Function(Name(func.Get()))
- << " added to module multiple times";
+ AddError(func) << "function " << style::Function(Name(func.Get()))
+ << " added to module multiple times";
}
}
@@ -413,12 +432,36 @@
return AddError(src);
}
+diag::Diagnostic& Validator::AddError(const BlockParam* param) {
+ DisassembleIfNeeded();
+ auto src = dis_.BlockParamSource(param);
+ return AddError(src);
+}
+
+diag::Diagnostic& Validator::AddError(const Function* func) {
+ DisassembleIfNeeded();
+ auto src = dis_.FunctionSource(func);
+ return AddError(src);
+}
+
+diag::Diagnostic& Validator::AddError(const FunctionParam* param) {
+ DisassembleIfNeeded();
+ auto src = dis_.FunctionParamSource(param);
+ return AddError(src);
+}
+
diag::Diagnostic& Validator::AddNote(const Instruction* inst) {
DisassembleIfNeeded();
auto src = dis_.InstructionSource(inst);
return AddNote(src);
}
+diag::Diagnostic& Validator::AddNote(const Function* func) {
+ DisassembleIfNeeded();
+ auto src = dis_.FunctionSource(func);
+ return AddNote(src);
+}
+
diag::Diagnostic& Validator::AddNote(const Instruction* inst, size_t idx) {
DisassembleIfNeeded();
auto src = dis_.OperandSource(Disassembler::IndexedValue{inst, static_cast<uint32_t>(idx)});
@@ -488,22 +531,50 @@
void Validator::CheckFunction(const Function* func) {
CheckBlock(func->Block());
- // References not allowed on function signatures even with Capability::kAllowRefTypes
for (auto* param : func->Params()) {
+ if (!param->Alive()) {
+ AddError(param) << "destroyed parameter found in function parameter list";
+ return;
+ }
+ if (!param->Function()) {
+ AddError(param) << "function parameter has nullptr parent function";
+ return;
+ } else if (param->Function() != func) {
+ AddError(param) << "function parameter has incorrect parent function";
+ AddNote(param->Function()) << "parent function declared here";
+ return;
+ }
+
+ // References not allowed on function signatures even with Capability::kAllowRefTypes
if (HoldsType<type::Reference>(param->Type())) {
- // TODO(dsinclair): Parameters need a source mapping.
- AddError(Source{}) << "references are not permitted as parameter types";
+ AddError(param) << "references are not permitted as parameter types";
}
}
if (HoldsType<type::Reference>(func->ReturnType())) {
- // TODO(dsinclair): Function need a source mapping.
- AddError(Source{}) << "references are not permitted as return types";
+ AddError(func) << "references are not permitted as return types";
}
}
void Validator::CheckBlock(const Block* blk) {
TINT_SCOPED_ASSIGNMENT(current_block_, blk);
+ if (auto* mb = blk->As<MultiInBlock>()) {
+ for (auto* param : mb->Params()) {
+ if (!param->Alive()) {
+ AddError(param) << "destroyed parameter found in block parameter list";
+ return;
+ }
+ if (!param->Block()) {
+ AddError(param) << "block parameter has nullptr parent block";
+ return;
+ } else if (param->Block() != mb) {
+ AddError(param) << "block parameter has incorrect parent block";
+ AddNote(param->Block()) << "parent block declared here";
+ return;
+ }
+ }
+ }
+
if (!blk->Terminator()) {
AddError(blk) << "block: does not end in a terminator instruction";
}
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 3e1a708..8742f8a 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -139,7 +139,10 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(error: function 'my_func' added to module multiple times
+ R"(:1:1 error: function 'my_func' added to module multiple times
+%my_func = func(%2:i32, %3:f32):void -> %b1 {
+^^^^^^^^
+
note: # Disassembly
%my_func = func(%2:i32, %3:f32):void -> %b1 {
%b1 = block {
@@ -154,6 +157,88 @@
)");
}
+TEST_F(IR_ValidatorTest, Function_DeadParameter) {
+ auto* f = b.Function("my_func", ty.void_());
+ auto* p = b.FunctionParam("my_param", ty.f32());
+ f->SetParams({p});
+ f->Block()->Append(b.Return(f));
+
+ p->Destroy();
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:17 error: destroyed parameter found in function parameter list
+%my_func = func(%my_param:f32):void -> %b1 {
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func(%my_param:f32):void -> %b1 {
+ %b1 = block {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_ParameterWithNullFunction) {
+ auto* f = b.Function("my_func", ty.void_());
+ auto* p = b.FunctionParam("my_param", ty.f32());
+ f->SetParams({p});
+ f->Block()->Append(b.Return(f));
+
+ p->SetFunction(nullptr);
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:17 error: function parameter has nullptr parent function
+%my_func = func(%my_param:f32):void -> %b1 {
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func(%my_param:f32):void -> %b1 {
+ %b1 = block {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_ParameterUsedInMultipleFunctions) {
+ auto* p = b.FunctionParam("my_param", ty.f32());
+ auto* f1 = b.Function("my_func1", ty.void_());
+ auto* f2 = b.Function("my_func2", ty.void_());
+ f1->SetParams({p});
+ f2->SetParams({p});
+ f1->Block()->Append(b.Return(f1));
+ f2->Block()->Append(b.Return(f2));
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:18 error: function parameter has incorrect parent function
+%my_func1 = func(%my_param:f32):void -> %b1 {
+ ^^^^^^^^^^^^^
+
+:6:1 note: parent function declared here
+%my_func2 = func(%my_param:f32):void -> %b2 {
+^^^^^^^^^
+
+note: # Disassembly
+%my_func1 = func(%my_param:f32):void -> %b1 {
+ %b1 = block {
+ ret
+ }
+}
+%my_func2 = func(%my_param:f32):void -> %b2 {
+ %b2 = block {
+ ret
+ }
+}
+)");
+}
+
TEST_F(IR_ValidatorTest, CallToFunctionOutsideModule) {
auto* f = b.Function("f", ty.void_());
auto* g = b.Function("g", ty.void_());
@@ -393,6 +478,115 @@
)");
}
+TEST_F(IR_ValidatorTest, Block_DeadParameter) {
+ auto* f = b.Function("my_func", ty.void_());
+
+ auto* p = b.BlockParam("my_param", ty.f32());
+ b.Append(f->Block(), [&] {
+ auto* l = b.Loop();
+ l->Body()->SetParams({p});
+ b.Append(l->Body(), [&] { b.ExitLoop(l); });
+ b.Return(f);
+ });
+
+ p->Destroy();
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:4:20 error: destroyed parameter found in block parameter list
+ %b2 = block (%my_param:f32) { # body
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func():void -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%my_param:f32) { # body
+ exit_loop # loop_1
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Block_ParameterWithNullBlock) {
+ auto* f = b.Function("my_func", ty.void_());
+
+ auto* p = b.BlockParam("my_param", ty.f32());
+ b.Append(f->Block(), [&] {
+ auto* l = b.Loop();
+ l->Body()->SetParams({p});
+ b.Append(l->Body(), [&] { b.ExitLoop(l); });
+ b.Return(f);
+ });
+
+ p->SetBlock(nullptr);
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:4:20 error: block parameter has nullptr parent block
+ %b2 = block (%my_param:f32) { # body
+ ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func():void -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%my_param:f32) { # body
+ exit_loop # loop_1
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Block_ParameterUsedInMultipleBlocks) {
+ auto* f = b.Function("my_func", ty.void_());
+
+ auto* p = b.BlockParam("my_param", ty.f32());
+ b.Append(f->Block(), [&] {
+ auto* l = b.Loop();
+ l->Body()->SetParams({p});
+ b.Append(l->Body(), [&] { b.Continue(l, p); });
+ l->Continuing()->SetParams({p});
+ b.Append(l->Continuing(), [&] { b.NextIteration(l, p); });
+ b.Return(f);
+ });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:4:20 error: block parameter has incorrect parent block
+ %b2 = block (%my_param:f32) { # body
+ ^^^^^^^^^^^^^
+
+:7:7 note: parent block declared here
+ %b3 = block (%my_param:f32) { # continuing
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func():void -> %b1 {
+ %b1 = block {
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block (%my_param:f32) { # body
+ continue %b3 %my_param:f32
+ }
+ %b3 = block (%my_param:f32) { # continuing
+ next_iteration %b2 %my_param:f32
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
TEST_F(IR_ValidatorTest, Access_NegativeIndex) {
auto* f = b.Function("my_func", ty.void_());
auto* obj = b.FunctionParam(ty.vec3<f32>());
diff --git a/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
index 04e7b14..fd59677 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
@@ -39,7 +39,10 @@
#include "src/tint/lang/wgsl/sem/statement.h"
#include "src/tint/lang/wgsl/sem/struct.h"
#include "src/tint/utils/containers/transform.h"
+#include "src/tint/utils/diagnostic/diagnostic.h"
+#include "src/tint/utils/result/result.h"
#include "src/tint/utils/rtti/switch.h"
+#include "src/tint/utils/text/text_style.h"
TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::PixelLocal);
TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::PixelLocal::RasterizerOrderedView);
@@ -120,7 +123,12 @@
// Obtain struct of the pixel local.
auto* pixel_local_str =
pixel_local_variable->Type()->UnwrapRef()->As<sem::Struct>();
- TransformEntryPoint(entry_point, pixel_local_variable, pixel_local_str);
+ if (auto res =
+ TransformEntryPoint(entry_point, pixel_local_variable, pixel_local_str);
+ res != Success) {
+ b.Diagnostics().Add(res.Failure().reason);
+ made_changes = true;
+ }
break; // Only a single `var<pixel_local>` can be used by an entry point.
}
@@ -139,9 +147,9 @@
/// @param entry_point the entry point
/// @param pixel_local_var the `var<pixel_local>`
/// @param pixel_local_str the struct type of the var
- void TransformEntryPoint(const sem::Function* entry_point,
- const sem::GlobalVariable* pixel_local_var,
- const sem::Struct* pixel_local_str) {
+ Result<SuccessType> TransformEntryPoint(const sem::Function* entry_point,
+ const sem::GlobalVariable* pixel_local_var,
+ const sem::Struct* pixel_local_str) {
// Wrap the old entry point "fn" into a new entry point where functions to load and store
// ROV data are called.
auto* original_entry_point_fn = entry_point->Declaration();
@@ -175,9 +183,12 @@
// load data from and store data into the ROVs.
auto load_rov_function_name = b.Symbols().New("load_from_pixel_local_storage");
auto store_rov_function_name = b.Symbols().New("store_into_pixel_local_storage");
- DeclareROVsAndLoadStoreFunctions(load_rov_function_name, store_rov_function_name,
- pixel_local_var->Declaration()->name->symbol.Name(),
- pixel_local_str);
+ if (auto res = DeclareROVsAndLoadStoreFunctions(
+ load_rov_function_name, store_rov_function_name,
+ pixel_local_var->Declaration()->name->symbol.Name(), pixel_local_str);
+ res != Success) {
+ return res.Failure();
+ }
// Declare new entry point
Vector<const ast::Statement*, 5> new_entry_point_function_body;
@@ -276,6 +287,7 @@
// Declare the new entry point that calls the inner function
b.Func(entry_point_name, std::move(new_entry_point_params), new_entry_point_return_type,
new_entry_point_function_body, Vector{b.Stage(ast::PipelineStage::kFragment)});
+ return Success;
}
/// Add the declarations of all the ROVs as a special type of read-write storage texture that
@@ -285,10 +297,11 @@
/// @param store_rov_function_name the name of the function that stores the data into the ROVs
/// @param pixel_local_variable_name the name of the pixel local variable
/// @param pixel_local_str the struct type of the pixel local variable
- void DeclareROVsAndLoadStoreFunctions(const Symbol& load_rov_function_name,
- const Symbol& store_rov_function_name,
- const std::string& pixel_local_variable_name,
- const sem::Struct* pixel_local_str) {
+ Result<SuccessType> DeclareROVsAndLoadStoreFunctions(
+ const Symbol& load_rov_function_name,
+ const Symbol& store_rov_function_name,
+ const std::string& pixel_local_variable_name,
+ const sem::Struct* pixel_local_str) {
std::string_view load_store_input_name = "my_input";
Vector load_parameters{b.Param(load_store_input_name, b.ty.vec4<f32>())};
Vector store_parameters{b.Param(load_store_input_name, b.ty.vec4<f32>())};
@@ -314,8 +327,11 @@
[&](const core::type::F32*) { return core::TexelFormat::kR32Float; },
TINT_ICE_ON_NO_MATCH);
auto rov_format = ROVTexelFormat(member->Index());
- auto rov_type = b.ty.storage_texture(core::type::TextureDimension::k2d, rov_format,
- core::Access::kReadWrite);
+ if (TINT_UNLIKELY(rov_format != Success)) {
+ return rov_format.Failure();
+ }
+ auto rov_type = b.ty.storage_texture(core::type::TextureDimension::k2d,
+ rov_format.Get(), core::Access::kReadWrite);
auto rov_symbol_name = b.Symbols().New("pixel_local_" + member->Name().Name());
b.GlobalVar(rov_symbol_name, rov_type,
tint::Vector{b.Binding(AInt(ROVRegisterIndex(member->Index()))),
@@ -352,7 +368,7 @@
// textureStore(
// pixel_local_member, rov_texcoord, vec4u(bitcast(PLS_Private_Variable.member)));
std::string rov_pixel_type;
- switch (rov_format) {
+ switch (rov_format.Get()) {
case core::TexelFormat::kR32Uint:
rov_pixel_type = "vec4u";
break;
@@ -363,8 +379,7 @@
rov_pixel_type = "vec4f";
break;
default:
- TINT_ICE() << "Invalid ROV format (now only R32Uint, R32Sint and R32Float are "
- "supported)";
+ TINT_UNREACHABLE();
break;
}
auto pixel_local_var_member_access_in_store_call =
@@ -374,7 +389,7 @@
to_vec4_call = b.Call(rov_pixel_type, pixel_local_var_member_access_in_store_call);
} else {
ast::Type rov_pixel_ast_type;
- switch (rov_format) {
+ switch (rov_format.Get()) {
case core::TexelFormat::kR32Uint:
rov_pixel_ast_type = b.ty.u32();
break;
@@ -399,6 +414,7 @@
b.Func(load_rov_function_name, std::move(load_parameters), b.ty.void_(), load_body);
b.Func(store_rov_function_name, std::move(store_parameters), b.ty.void_(), store_body);
+ return Success;
}
/// Find and get `@builtin(position)` which is needed for loading and storing data with ROVs
@@ -462,12 +478,14 @@
/// @returns the texel format for the pixel local field with the given index
/// @param field_index the pixel local field index
- core::TexelFormat ROVTexelFormat(uint32_t field_index) {
+ Result<core::TexelFormat> ROVTexelFormat(uint32_t field_index) {
auto format = cfg.pls_member_to_rov_format.Get(field_index);
if (TINT_UNLIKELY(!format)) {
- b.Diagnostics().AddError(diag::System::Transform, Source{})
- << "PixelLocal::Config::attachments missing entry for field " << field_index;
- return core::TexelFormat::kUndefined;
+ diag::Diagnostic err;
+ err.severity = diag::Severity::Error;
+ err.message << "PixelLocal::Config::attachments missing entry for field "
+ << field_index;
+ return Failure{std::move(err)};
}
return *format;
}
diff --git a/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc b/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc
index 9f644a7..f34138a 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc
@@ -61,6 +61,31 @@
EXPECT_FALSE(ShouldRun<PixelLocal>(src, Bindings({})));
}
+TEST_F(HLSLPixelLocalTest, MissingBindings) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+@fragment
+fn F() -> @location(0) vec4f {
+ P.a += 42;
+ return vec4f(1, 0, 0, 1);
+}
+)";
+
+ auto* expect = R"(error: PixelLocal::Config::attachments missing entry for field 0)";
+ ast::transform::DataMap data;
+ data.Add<PixelLocal::Config>();
+ auto got = Run<PixelLocal>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(HLSLPixelLocalTest, UseInEntryPoint_NoPosition) {
auto* src = R"(
enable chromium_experimental_pixel_local;
diff --git a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
index 901b4bb..98b3067 100644
--- a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
+++ b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
@@ -134,7 +134,7 @@
}
// Replace each access instruction that we recorded.
- Hashmap<core::ir::Value*, core::ir::Value*, 4> object_to_local;
+ Hashmap<core::ir::Value*, core::ir::Value*, 4> object_to_var;
Hashmap<PartialAccess, core::ir::Value*, 4> source_object_to_value;
for (const auto& to_replace : worklist) {
auto* access = to_replace.access;
@@ -145,24 +145,58 @@
if (to_replace.first_dynamic_index > 0) {
PartialAccess partial_access = {
access->Object(), access->Indices().Truncate(to_replace.first_dynamic_index)};
- source_object = source_object_to_value.GetOrAdd(partial_access, [&] {
- auto* intermediate_source = builder.Access(to_replace.dynamic_index_source_type,
- source_object, partial_access.indices);
- intermediate_source->InsertBefore(access);
- return intermediate_source->Result(0);
- });
+ source_object =
+ source_object_to_value.GetOrAdd(partial_access, [&]() -> core::ir::Value* {
+ // If the source is a constant, then the partial access will also produce a
+ // constant. Extract the constant::Value and use that as the new source object.
+ if (source_object->Is<core::ir::Constant>()) {
+ for (const auto& i : partial_access.indices) {
+ auto idx = i->As<core::ir::Constant>()->Value()->ValueAs<uint32_t>();
+ source_object = builder.Constant(
+ source_object->As<core::ir::Constant>()->Value()->Index(idx));
+ }
+ return source_object;
+ }
+
+ // Extract a non-constant intermediate source using an access instruction that
+ // we insert immediately after the definition of the root source object.
+ auto* intermediate_source =
+ builder.Access(to_replace.dynamic_index_source_type, source_object,
+ partial_access.indices);
+ builder.InsertAfter(source_object,
+ [&] { builder.Append(intermediate_source); });
+ return intermediate_source->Result(0);
+ });
}
- // Declare a local variable and copy the source object to it.
- auto* local = object_to_local.GetOrAdd(source_object, [&] {
- auto* decl = builder.Var(ir.Types().ptr(
- core::AddressSpace::kFunction, source_object->Type(), core::Access::kReadWrite));
+ // Declare a variable and copy the source object to it.
+ auto* var = object_to_var.GetOrAdd(source_object, [&] {
+ // If the source object is a constant we use a module-scope variable, as it could be
+ // indexed by multiple functions. Otherwise, we declare a function-scope variable
+ // immediately after the definition of the source object.
+ core::ir::Var* decl = nullptr;
+ if (source_object->Is<core::ir::Constant>()) {
+ decl = builder.Var(ir.Types().ptr(core::AddressSpace::kPrivate,
+ source_object->Type(), core::Access::kReadWrite));
+ ir.root_block->Append(decl);
+ } else {
+ builder.InsertAfter(source_object, [&] {
+ decl = builder.Var(ir.Types().ptr(core::AddressSpace::kFunction,
+ source_object->Type(),
+ core::Access::kReadWrite));
+
+ // If we ever support value declarations at module-scope, we will need to modify
+ // the partial access logic above since `access` instructions cannot be used in
+ // the root block.
+ TINT_ASSERT(decl->Block() != ir.root_block);
+ });
+ }
+
decl->SetInitializer(source_object);
- decl->InsertBefore(access);
return decl->Result(0);
});
- // Create a new access instruction using the local variable as the source.
+ // Create a new access instruction using the new variable as the source.
Vector<core::ir::Value*, 4> indices{
access->Indices().Offset(to_replace.first_dynamic_index)};
const core::type::Type* access_type = access->Result(0)->Type();
@@ -178,9 +212,9 @@
vector_index = indices.Pop();
}
+ auto addrspace = var->Type()->As<core::type::Pointer>()->AddressSpace();
core::ir::Instruction* new_access = builder.Access(
- ir.Types().ptr(core::AddressSpace::kFunction, access_type, core::Access::kReadWrite),
- local, indices);
+ ir.Types().ptr(addrspace, access_type, core::Access::kReadWrite), var, indices);
new_access->InsertBefore(access);
core::ir::Instruction* load = nullptr;
diff --git a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
index e4adae9..657488b 100644
--- a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
+++ b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
@@ -430,5 +430,449 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToFuncParam_FromDifferentBlocks) {
+ auto* arr = b.FunctionParam(ty.array<i32, 4>());
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({arr, cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_b));
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:array<i32, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %6:i32 = access %2, %4
+ ret %6
+ }
+ %b3 = block { # false
+ %7:i32 = access %2, %5
+ ret %7
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:array<i32, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ %6:ptr<function, array<i32, 4>, read_write> = var, %2
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %7:ptr<function, i32, read_write> = access %6, %4
+ %8:i32 = load %7
+ ret %8
+ }
+ %b3 = block { # false
+ %9:ptr<function, i32, read_write> = access %6, %5
+ %10:i32 = load %9
+ ret %10
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+ MultipleAccessesToFuncParam_FromDifferentBlocks_WithLeadingConstantIndex) {
+ auto* arr = b.FunctionParam(ty.array(ty.array<i32, 4>(), 4));
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({arr, cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_b));
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:array<array<i32, 4>, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %6:i32 = access %2, 0u, %4
+ ret %6
+ }
+ %b3 = block { # false
+ %7:i32 = access %2, 0u, %5
+ ret %7
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:array<array<i32, 4>, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ %6:array<i32, 4> = access %2, 0u
+ %7:ptr<function, array<i32, 4>, read_write> = var, %6
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %8:ptr<function, i32, read_write> = access %7, %4
+ %9:i32 = load %8
+ ret %9
+ }
+ %b3 = block { # false
+ %10:ptr<function, i32, read_write> = access %7, %5
+ %11:i32 = load %10
+ ret %11
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToBlockParam_FromDifferentBlocks) {
+ auto* arr = b.BlockParam(ty.array<i32, 4>());
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ loop->Body()->SetParams({arr});
+ b.Append(loop->Body(), [&] {
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_b));
+ });
+ b.Unreachable();
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<i32, 4>) { # body
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %6:i32 = access %5:array<i32, 4>, %3
+ ret %6
+ }
+ %b4 = block { # false
+ %7:i32 = access %5:array<i32, 4>, %4
+ ret %7
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<i32, 4>) { # body
+ %6:ptr<function, array<i32, 4>, read_write> = var, %5:array<i32, 4>
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %7:ptr<function, i32, read_write> = access %6, %3
+ %8:i32 = load %7
+ ret %8
+ }
+ %b4 = block { # false
+ %9:ptr<function, i32, read_write> = access %6, %4
+ %10:i32 = load %9
+ ret %10
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+ MultipleAccessesToBlockParam_FromDifferentBlocks_WithLeadingConstantIndex) {
+ auto* arr = b.BlockParam(ty.array(ty.array<i32, 4>(), 4));
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ loop->Body()->SetParams({arr});
+ b.Append(loop->Body(), [&] {
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_b));
+ });
+ b.Unreachable();
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<array<i32, 4>, 4>) { # body
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %6:i32 = access %5:array<array<i32, 4>, 4>, 0u, %3
+ ret %6
+ }
+ %b4 = block { # false
+ %7:i32 = access %5:array<array<i32, 4>, 4>, 0u, %4
+ ret %7
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<array<i32, 4>, 4>) { # body
+ %6:array<i32, 4> = access %5:array<array<i32, 4>, 4>, 0u
+ %7:ptr<function, array<i32, 4>, read_write> = var, %6
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %8:ptr<function, i32, read_write> = access %7, %3
+ %9:i32 = load %8
+ ret %9
+ }
+ %b4 = block { # false
+ %10:ptr<function, i32, read_write> = access %7, %4
+ %11:i32 = load %10
+ ret %11
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToConstant_FromDifferentFunctions) {
+ auto* arr = b.Constant(mod.constant_values.Zero(ty.array<i32, 4>()));
+
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* func_a = b.Function("func_a", ty.i32());
+ func_a->SetParams({idx_a});
+ b.Append(func_a->Block(), [&] { //
+ b.Return(func_a, b.Access(ty.i32(), arr, idx_a));
+ });
+
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func_b = b.Function("func_b", ty.i32());
+ func_b->SetParams({idx_b});
+ b.Append(func_b->Block(), [&] { //
+ b.Return(func_b, b.Access(ty.i32(), arr, idx_b));
+ });
+
+ auto* idx_c = b.FunctionParam(ty.i32());
+ auto* func_c = b.Function("func_c", ty.i32());
+ func_c->SetParams({idx_c});
+ b.Append(func_c->Block(), [&] { //
+ b.Return(func_c, b.Access(ty.i32(), arr, idx_c));
+ });
+
+ auto* src = R"(
+%func_a = func(%2:i32):i32 -> %b1 {
+ %b1 = block {
+ %3:i32 = access array<i32, 4>(0i), %2
+ ret %3
+ }
+}
+%func_b = func(%5:i32):i32 -> %b2 {
+ %b2 = block {
+ %6:i32 = access array<i32, 4>(0i), %5
+ ret %6
+ }
+}
+%func_c = func(%8:i32):i32 -> %b3 {
+ %b3 = block {
+ %9:i32 = access array<i32, 4>(0i), %8
+ ret %9
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %1:ptr<private, array<i32, 4>, read_write> = var, array<i32, 4>(0i)
+}
+
+%func_a = func(%3:i32):i32 -> %b2 {
+ %b2 = block {
+ %4:ptr<private, i32, read_write> = access %1, %3
+ %5:i32 = load %4
+ ret %5
+ }
+}
+%func_b = func(%7:i32):i32 -> %b3 {
+ %b3 = block {
+ %8:ptr<private, i32, read_write> = access %1, %7
+ %9:i32 = load %8
+ ret %9
+ }
+}
+%func_c = func(%11:i32):i32 -> %b4 {
+ %b4 = block {
+ %12:ptr<private, i32, read_write> = access %1, %11
+ %13:i32 = load %12
+ ret %13
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+ MultipleAccessesToConstant_FromDifferentFunctions_WithLeadingConstantIndex) {
+ auto* arr = b.Constant(mod.constant_values.Zero(ty.array(ty.array<i32, 4>(), 4)));
+
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* func_a = b.Function("func_a", ty.i32());
+ func_a->SetParams({idx_a});
+ b.Append(func_a->Block(), [&] { //
+ b.Return(func_a, b.Access(ty.i32(), arr, 0_u, idx_a));
+ });
+
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func_b = b.Function("func_b", ty.i32());
+ func_b->SetParams({idx_b});
+ b.Append(func_b->Block(), [&] { //
+ b.Return(func_b, b.Access(ty.i32(), arr, 0_u, idx_b));
+ });
+
+ auto* idx_c = b.FunctionParam(ty.i32());
+ auto* func_c = b.Function("func_c", ty.i32());
+ func_c->SetParams({idx_c});
+ b.Append(func_c->Block(), [&] { //
+ b.Return(func_c, b.Access(ty.i32(), arr, 0_u, idx_c));
+ });
+
+ auto* src = R"(
+%func_a = func(%2:i32):i32 -> %b1 {
+ %b1 = block {
+ %3:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %2
+ ret %3
+ }
+}
+%func_b = func(%5:i32):i32 -> %b2 {
+ %b2 = block {
+ %6:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %5
+ ret %6
+ }
+}
+%func_c = func(%8:i32):i32 -> %b3 {
+ %b3 = block {
+ %9:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %8
+ ret %9
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %1:ptr<private, array<i32, 4>, read_write> = var, array<i32, 4>(0i)
+}
+
+%func_a = func(%3:i32):i32 -> %b2 {
+ %b2 = block {
+ %4:ptr<private, i32, read_write> = access %1, %3
+ %5:i32 = load %4
+ ret %5
+ }
+}
+%func_b = func(%7:i32):i32 -> %b3 {
+ %b3 = block {
+ %8:ptr<private, i32, read_write> = access %1, %7
+ %9:i32 = load %8
+ ret %9
+ }
+}
+%func_c = func(%11:i32):i32 -> %b4 {
+ %b4 = block {
+ %12:ptr<private, i32, read_write> = access %1, %11
+ %13:i32 = load %12
+ ret %13
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::writer::raise
diff --git a/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc b/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
index b745f56..3b41233 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
@@ -46,7 +46,9 @@
return;
}
- auto dst = tint::wgsl::writer::IRToProgram(ir);
+ writer::ProgramOptions program_options;
+ program_options.allowed_features = AllowedFeatures::Everything();
+ auto dst = tint::wgsl::writer::IRToProgram(ir, program_options);
if (!dst.IsValid()) {
std::cerr << "IR:\n" << core::ir::Disassemble(ir) << std::endl;
if (auto result = tint::wgsl::writer::Generate(dst, {}); result == Success) {
diff --git a/src/tint/lang/wgsl/ir_roundtrip_test.cc b/src/tint/lang/wgsl/ir_roundtrip_test.cc
index 5cb3f99..6573c32 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_test.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_test.cc
@@ -196,6 +196,14 @@
)");
}
+TEST_F(IRToProgramRoundtripTest, SingleFunction_UnrestrictedPointerParameters) {
+ RUN_TEST(R"(
+fn f(p : ptr<uniform, i32>) -> i32 {
+ return *(p);
+}
+)");
+}
+
////////////////////////////////////////////////////////////////////////////////
// Struct declaration
////////////////////////////////////////////////////////////////////////////////
@@ -1921,6 +1929,21 @@
)");
}
+TEST_F(IRToProgramRoundtripTest, PhonyAssign_Conversion) {
+ RUN_TEST(R"(
+fn f() {
+ let i = 42i;
+ _ = u32(i);
+}
+)",
+ R"(
+fn f() {
+ let i = 42i;
+ _ = u32(i);
+}
+)");
+}
+
////////////////////////////////////////////////////////////////////////////////
// let
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/wgsl/reader/lower/lower.cc b/src/tint/lang/wgsl/reader/lower/lower.cc
index 21782ee..94f2bac 100644
--- a/src/tint/lang/wgsl/reader/lower/lower.cc
+++ b/src/tint/lang/wgsl/reader/lower/lower.cc
@@ -35,6 +35,7 @@
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/lang/wgsl/builtin_fn.h"
#include "src/tint/lang/wgsl/ir/builtin_call.h"
+#include "src/tint/utils/ice/ice.h"
namespace tint::wgsl::reader {
namespace {
@@ -166,10 +167,15 @@
CASE(kAtomicCompareExchangeWeak)
CASE(kSubgroupBallot)
CASE(kSubgroupBroadcast)
- default:
- TINT_ICE() << "unhandled builtin function: " << fn;
- return core::BuiltinFn::kNone;
+
+ case tint::wgsl::BuiltinFn::kBitcast: // should lower to ir::Bitcast
+ case tint::wgsl::BuiltinFn::kWorkgroupUniformLoad: // should be handled in Lower()
+ case tint::wgsl::BuiltinFn::kTintMaterialize:
+ case tint::wgsl::BuiltinFn::kNone:
+ break;
}
+ TINT_ICE() << "unhandled builtin function: " << fn;
+ return core::BuiltinFn::kNone;
}
} // namespace
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index 8bf72e7..8840d93 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -68,6 +68,7 @@
#include "src/tint/lang/core/ir/user_call.h"
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/lang/core/ir/var.h"
+#include "src/tint/lang/core/texel_format.h"
#include "src/tint/lang/core/type/atomic.h"
#include "src/tint/lang/core/type/depth_multisampled_texture.h"
#include "src/tint/lang/core/type/depth_texture.h"
@@ -75,7 +76,9 @@
#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/lang/core/type/reference.h"
#include "src/tint/lang/core/type/sampler.h"
+#include "src/tint/lang/core/type/storage_texture.h"
#include "src/tint/lang/core/type/texture.h"
+#include "src/tint/lang/core/type/type.h"
#include "src/tint/lang/wgsl/ir/builtin_call.h"
#include "src/tint/lang/wgsl/ir/unary.h"
#include "src/tint/lang/wgsl/program/program_builder.h"
@@ -968,6 +971,10 @@
return b.ty.sampled_texture(t->dim(), el);
},
[&](const core::type::StorageTexture* t) {
+ if (RequiresChromiumInternalGraphite(t)) {
+ Enable(wgsl::Extension::kChromiumInternalGraphite);
+ }
+
return b.ty.storage_texture(t->dim(), t->texel_format(), t->access());
},
[&](const core::type::Sampler* s) { return b.ty.sampler(s->kind()); },
@@ -1051,12 +1058,17 @@
});
}
- /// Associates the IR value @p value with the AST expression @p expr.
+ /// Associates the IR value @p value with the AST expression @p expr if it is used, otherwise
+ /// creates a phony assignment with @p expr.
void Bind(const core::ir::Value* value, const ast::Expression* expr) {
TINT_ASSERT(value);
- // Value will be inlined at its place of usage.
- if (TINT_UNLIKELY(!bindings_.Add(value, InlinedValue{expr}))) {
- TINT_ICE() << "Bind(" << value->TypeInfo().name << ") called twice for same value";
+ if (value->IsUsed()) {
+ // Value will be inlined at its place of usage.
+ if (TINT_UNLIKELY(!bindings_.Add(value, InlinedValue{expr}))) {
+ TINT_ICE() << "Bind(" << value->TypeInfo().name << ") called twice for same value";
+ }
+ } else {
+ Append(b.Assign(b.Phony(), expr));
}
}
@@ -1192,6 +1204,12 @@
return false;
}
}
+
+ /// @returns true if the storage texture type requires the kChromiumInternalGraphite extension
+ /// to be enabled.
+ bool RequiresChromiumInternalGraphite(const core::type::StorageTexture* tex) {
+ return tex->texel_format() == core::TexelFormat::kR8Unorm;
+ }
};
} // namespace
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
index 10e88ca..30e6795 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
@@ -30,8 +30,12 @@
#include <sstream>
#include <string>
+#include "src/tint/lang/core/access.h"
+#include "src/tint/lang/core/address_space.h"
#include "src/tint/lang/core/ir/disassembler.h"
+#include "src/tint/lang/core/texel_format.h"
#include "src/tint/lang/core/type/storage_texture.h"
+#include "src/tint/lang/core/type/texture_dimension.h"
#include "src/tint/lang/wgsl/ir/builtin_call.h"
#include "src/tint/lang/wgsl/ir/unary.h"
#include "src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.h"
@@ -2629,5 +2633,23 @@
)");
}
+////////////////////////////////////////////////////////////////////////////////
+// chromium_internal_graphite
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRToProgramTest, Enable_ChromiumInternalGraphite_SubgroupBallot) {
+ b.Append(b.ir.root_block, [&] {
+ auto t = b.Var("T", ty.ref<core::AddressSpace::kHandle>(ty.Get<core::type::StorageTexture>(
+ core::type::TextureDimension::k2d, core::TexelFormat::kR8Unorm,
+ core::Access::kRead, ty.f32())));
+ t->SetBindingPoint(0, 0);
+ });
+
+ EXPECT_WGSL(R"(
+enable chromium_internal_graphite;
+
+@group(0) @binding(0) var T : texture_storage_2d<r8unorm, read>;
+)");
+}
+
} // namespace
} // namespace tint::wgsl::writer
diff --git a/src/tint/lang/wgsl/writer/raise/BUILD.cmake b/src/tint/lang/wgsl/writer/raise/BUILD.cmake
index d2ae1d1..ff5b855 100644
--- a/src/tint/lang/wgsl/writer/raise/BUILD.cmake
+++ b/src/tint/lang/wgsl/writer/raise/BUILD.cmake
@@ -112,3 +112,38 @@
tint_target_add_external_dependencies(tint_lang_wgsl_writer_raise_test test
"gtest"
)
+
+################################################################################
+# Target: tint_lang_wgsl_writer_raise_fuzz
+# Kind: fuzz
+################################################################################
+tint_add_target(tint_lang_wgsl_writer_raise_fuzz fuzz
+ lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc
+ lang/wgsl/writer/raise/raise_fuzz.cc
+ lang/wgsl/writer/raise/rename_conflicts_fuzz.cc
+ lang/wgsl/writer/raise/value_to_let_fuzz.cc
+)
+
+tint_target_add_dependencies(tint_lang_wgsl_writer_raise_fuzz fuzz
+ tint_api_common
+ tint_cmd_fuzz_ir_fuzz
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_ir
+ tint_lang_core_type
+ tint_lang_wgsl_writer_raise
+ tint_utils_bytes
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_id
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_reflection
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_symbol
+ tint_utils_text
+ tint_utils_traits
+)
diff --git a/src/tint/lang/wgsl/writer/raise/BUILD.gn b/src/tint/lang/wgsl/writer/raise/BUILD.gn
index 095b41e..1f5cfa8 100644
--- a/src/tint/lang/wgsl/writer/raise/BUILD.gn
+++ b/src/tint/lang/wgsl/writer/raise/BUILD.gn
@@ -112,3 +112,35 @@
]
}
}
+
+tint_fuzz_source_set("fuzz") {
+ sources = [
+ "ptr_to_ref_fuzz.cc",
+ "raise_fuzz.cc",
+ "rename_conflicts_fuzz.cc",
+ "value_to_let_fuzz.cc",
+ ]
+ deps = [
+ "${tint_src_dir}/api/common",
+ "${tint_src_dir}/cmd/fuzz/ir:fuzz",
+ "${tint_src_dir}/lang/core",
+ "${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/wgsl/writer/raise",
+ "${tint_src_dir}/utils/bytes",
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/diagnostic",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/id",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/reflection",
+ "${tint_src_dir}/utils/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/symbol",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+}
diff --git a/src/tint/lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc b/src/tint/lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc
new file mode 100644
index 0000000..b67fb22
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc
@@ -0,0 +1,50 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT\ OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/wgsl/writer/raise/ptr_to_ref.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void PtrToRefFuzzer(core::ir::Module& module) {
+ if (auto res = PtrToRef(module); res != Success) {
+ return;
+ }
+
+ core::ir::Capabilities capabilities{core::ir::Capability::kAllowRefTypes};
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of PtrToRef failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::PtrToRefFuzzer);
diff --git a/src/tint/lang/wgsl/writer/raise/raise.cc b/src/tint/lang/wgsl/writer/raise/raise.cc
index 46cfe7e..dcb3e4e 100644
--- a/src/tint/lang/wgsl/writer/raise/raise.cc
+++ b/src/tint/lang/wgsl/writer/raise/raise.cc
@@ -108,6 +108,10 @@
CASE(kPack2X16Unorm)
CASE(kPack4X8Snorm)
CASE(kPack4X8Unorm)
+ CASE(kPack4XI8)
+ CASE(kPack4XU8)
+ CASE(kPack4XI8Clamp)
+ CASE(kPack4XU8Clamp)
CASE(kPow)
CASE(kQuantizeToF16)
CASE(kRadians)
@@ -133,6 +137,8 @@
CASE(kUnpack2X16Unorm)
CASE(kUnpack4X8Snorm)
CASE(kUnpack4X8Unorm)
+ CASE(kUnpack4XI8)
+ CASE(kUnpack4XU8)
CASE(kWorkgroupBarrier)
CASE(kTextureBarrier)
CASE(kTextureDimensions)
@@ -163,10 +169,11 @@
CASE(kAtomicCompareExchangeWeak)
CASE(kSubgroupBallot)
CASE(kSubgroupBroadcast)
- default:
- TINT_ICE() << "unhandled builtin function: " << fn;
- return wgsl::BuiltinFn::kNone;
+ case core::BuiltinFn::kNone:
+ break;
}
+ TINT_ICE() << "unhandled builtin function: " << fn;
+ return wgsl::BuiltinFn::kNone;
}
void ReplaceBuiltinFnCall(core::ir::Module& mod, core::ir::CoreBuiltinCall* call) {
diff --git a/src/tint/lang/wgsl/writer/raise/raise_fuzz.cc b/src/tint/lang/wgsl/writer/raise/raise_fuzz.cc
new file mode 100644
index 0000000..e63783c
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/raise_fuzz.cc
@@ -0,0 +1,50 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT\ OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/wgsl/writer/raise/raise.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void RaiseFuzzer(core::ir::Module& module) {
+ if (auto res = Raise(module); res != Success) {
+ return;
+ }
+
+ core::ir::Capabilities capabilities{core::ir::Capability::kAllowRefTypes};
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of Raise failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::RaiseFuzzer);
diff --git a/src/tint/lang/wgsl/writer/raise/rename_conflicts_fuzz.cc b/src/tint/lang/wgsl/writer/raise/rename_conflicts_fuzz.cc
new file mode 100644
index 0000000..8ac74e0
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/rename_conflicts_fuzz.cc
@@ -0,0 +1,50 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT\ OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/wgsl/writer/raise/rename_conflicts.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void RenameConflictsFuzzer(core::ir::Module& module) {
+ if (auto res = RenameConflicts(module); res != Success) {
+ return;
+ }
+
+ core::ir::Capabilities capabilities;
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of RenameConflicts failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::RenameConflictsFuzzer);
diff --git a/src/tint/lang/wgsl/writer/raise/value_to_let_fuzz.cc b/src/tint/lang/wgsl/writer/raise/value_to_let_fuzz.cc
new file mode 100644
index 0000000..84755d0
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/value_to_let_fuzz.cc
@@ -0,0 +1,50 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT\ OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/wgsl/writer/raise/value_to_let.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void ValueToLetFuzzer(core::ir::Module& module) {
+ if (auto res = ValueToLet(module); res != Success) {
+ return;
+ }
+
+ core::ir::Capabilities capabilities;
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of ValueToLet failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::ValueToLetFuzzer);
diff --git a/src/tint/utils/bytes/BUILD.bazel b/src/tint/utils/bytes/BUILD.bazel
index 3b0bb76..3daeace 100644
--- a/src/tint/utils/bytes/BUILD.bazel
+++ b/src/tint/utils/bytes/BUILD.bazel
@@ -39,11 +39,13 @@
cc_library(
name = "bytes",
srcs = [
+ "buffer_reader.cc",
"bytes.cc",
"reader.cc",
"writer.cc",
],
hdrs = [
+ "buffer_reader.h",
"buffer_writer.h",
"decoder.h",
"endianness.h",
@@ -71,9 +73,9 @@
name = "test",
alwayslink = True,
srcs = [
+ "buffer_reader_test.cc",
"buffer_writer_test.cc",
"decoder_test.cc",
- "reader_test.cc",
"swap_test.cc",
],
deps = [
diff --git a/src/tint/utils/bytes/BUILD.cmake b/src/tint/utils/bytes/BUILD.cmake
index 6d919d3..3438bbb 100644
--- a/src/tint/utils/bytes/BUILD.cmake
+++ b/src/tint/utils/bytes/BUILD.cmake
@@ -39,6 +39,8 @@
# Kind: lib
################################################################################
tint_add_target(tint_utils_bytes lib
+ utils/bytes/buffer_reader.cc
+ utils/bytes/buffer_reader.h
utils/bytes/buffer_writer.h
utils/bytes/bytes.cc
utils/bytes/decoder.h
@@ -69,9 +71,9 @@
# Kind: test
################################################################################
tint_add_target(tint_utils_bytes_test test
+ utils/bytes/buffer_reader_test.cc
utils/bytes/buffer_writer_test.cc
utils/bytes/decoder_test.cc
- utils/bytes/reader_test.cc
utils/bytes/swap_test.cc
)
diff --git a/src/tint/utils/bytes/BUILD.gn b/src/tint/utils/bytes/BUILD.gn
index d3babe2..6618381 100644
--- a/src/tint/utils/bytes/BUILD.gn
+++ b/src/tint/utils/bytes/BUILD.gn
@@ -44,6 +44,8 @@
libtint_source_set("bytes") {
sources = [
+ "buffer_reader.cc",
+ "buffer_reader.h",
"buffer_writer.h",
"bytes.cc",
"decoder.h",
@@ -71,9 +73,9 @@
if (tint_build_unittests) {
tint_unittests_source_set("unittests") {
sources = [
+ "buffer_reader_test.cc",
"buffer_writer_test.cc",
"decoder_test.cc",
- "reader_test.cc",
"swap_test.cc",
]
deps = [
diff --git a/src/tint/utils/bytes/buffer_reader.cc b/src/tint/utils/bytes/buffer_reader.cc
new file mode 100644
index 0000000..4294fef
--- /dev/null
+++ b/src/tint/utils/bytes/buffer_reader.cc
@@ -0,0 +1,46 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/utils/bytes/buffer_reader.h"
+
+namespace tint::bytes {
+
+BufferReader::~BufferReader() = default;
+
+size_t BufferReader::Read(std::byte* out, size_t count) {
+ size_t n = std::min(count, bytes_remaining_);
+ memcpy(out, data_, n);
+ data_ += n;
+ bytes_remaining_ -= n;
+ return n;
+}
+
+bool BufferReader::IsEOF() const {
+ return bytes_remaining_ == 0;
+}
+
+} // namespace tint::bytes
diff --git a/src/tint/utils/bytes/buffer_reader.h b/src/tint/utils/bytes/buffer_reader.h
new file mode 100644
index 0000000..0367bb0
--- /dev/null
+++ b/src/tint/utils/bytes/buffer_reader.h
@@ -0,0 +1,81 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_UTILS_BYTES_BUFFER_READER_H_
+#define SRC_TINT_UTILS_BYTES_BUFFER_READER_H_
+
+#include <algorithm>
+#include <string>
+
+#include "src/tint/utils/bytes/reader.h"
+#include "src/tint/utils/ice/ice.h"
+
+namespace tint::bytes {
+
+/// BufferReader is an implementation of the Reader interface backed by a buffer.
+class BufferReader final : public Reader {
+ public:
+ // Destructor
+ ~BufferReader() override;
+
+ /// Constructor
+ /// @param data the data to read from
+ /// @param size the number of bytes in the buffer
+ BufferReader(const std::byte* data, size_t size) : data_(data), bytes_remaining_(size) {
+ TINT_ASSERT(data);
+ }
+
+ /// Constructor
+ /// @param string the string to read from
+ explicit BufferReader(std::string_view string)
+ : data_(reinterpret_cast<const std::byte*>(string.data())),
+ bytes_remaining_(string.length()) {}
+
+ /// Constructor
+ /// @param slice the byte slice to read from
+ explicit BufferReader(Slice<const std::byte> slice)
+ : data_(slice.data), bytes_remaining_(slice.len) {
+ TINT_ASSERT(slice.data);
+ }
+
+ /// @copydoc Reader::Read
+ size_t Read(std::byte* out, size_t count) override;
+
+ /// @copydoc Reader::IsEOF
+ bool IsEOF() const override;
+
+ private:
+ /// The data to read from
+ const std::byte* data_ = nullptr;
+
+ /// The number of bytes remaining
+ size_t bytes_remaining_ = 0;
+};
+
+} // namespace tint::bytes
+
+#endif // SRC_TINT_UTILS_BYTES_BUFFER_READER_H_
diff --git a/src/tint/utils/bytes/reader_test.cc b/src/tint/utils/bytes/buffer_reader_test.cc
similarity index 97%
rename from src/tint/utils/bytes/reader_test.cc
rename to src/tint/utils/bytes/buffer_reader_test.cc
index 2e90e88..b0b72d8 100644
--- a/src/tint/utils/bytes/reader_test.cc
+++ b/src/tint/utils/bytes/buffer_reader_test.cc
@@ -1,4 +1,4 @@
-// Copyright 2023 The Dawn & Tint Authors
+// Copyright 2024 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
@@ -25,7 +25,7 @@
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-#include "src/tint/utils/bytes/reader.h"
+#include "src/tint/utils/bytes/buffer_reader.h"
#include "gtest/gtest.h"
@@ -33,7 +33,7 @@
namespace {
template <typename... ARGS>
-auto Data(ARGS&&... args) {
+auto Data(ARGS... args) {
return std::array{std::byte{static_cast<uint8_t>(args)}...};
}
diff --git a/src/tint/utils/bytes/decoder_test.cc b/src/tint/utils/bytes/decoder_test.cc
index 125d7e6..11f663a 100644
--- a/src/tint/utils/bytes/decoder_test.cc
+++ b/src/tint/utils/bytes/decoder_test.cc
@@ -34,6 +34,7 @@
#include <utility>
#include "gmock/gmock.h"
+#include "src/tint/utils/bytes/buffer_reader.h"
#include "src/tint/utils/result/result.h"
namespace tint {
diff --git a/src/tint/utils/bytes/reader.cc b/src/tint/utils/bytes/reader.cc
index fc2dfc2..e12a020 100644
--- a/src/tint/utils/bytes/reader.cc
+++ b/src/tint/utils/bytes/reader.cc
@@ -31,18 +31,4 @@
Reader::~Reader() = default;
-BufferReader::~BufferReader() = default;
-
-size_t BufferReader::Read(std::byte* out, size_t count) {
- size_t n = std::min(count, bytes_remaining_);
- memcpy(out, data_, n);
- data_ += n;
- bytes_remaining_ -= n;
- return n;
-}
-
-bool BufferReader::IsEOF() const {
- return bytes_remaining_ == 0;
-}
-
} // namespace tint::bytes
diff --git a/src/tint/utils/bytes/reader.h b/src/tint/utils/bytes/reader.h
index 3b5f51a..f0d074a 100644
--- a/src/tint/utils/bytes/reader.h
+++ b/src/tint/utils/bytes/reader.h
@@ -28,15 +28,11 @@
#ifndef SRC_TINT_UTILS_BYTES_READER_H_
#define SRC_TINT_UTILS_BYTES_READER_H_
-#include <algorithm>
-#include <cstdint>
#include <string>
#include "src/tint/utils/bytes/endianness.h"
#include "src/tint/utils/bytes/swap.h"
-#include "src/tint/utils/containers/slice.h"
#include "src/tint/utils/result/result.h"
-
namespace tint::bytes {
/// A binary stream reader interface
@@ -114,46 +110,6 @@
}
};
-/// BufferReader is an implementation of the Reader interface backed by a buffer.
-class BufferReader final : public Reader {
- public:
- // Destructor
- ~BufferReader() override;
-
- /// Constructor
- /// @param data the data to read from
- /// @param size the number of bytes in the buffer
- BufferReader(const std::byte* data, size_t size) : data_(data), bytes_remaining_(size) {
- TINT_ASSERT(data);
- }
-
- /// Constructor
- /// @param string the string to read from
- explicit BufferReader(std::string_view string)
- : data_(reinterpret_cast<const std::byte*>(string.data())),
- bytes_remaining_(string.length()) {}
-
- /// Constructor
- /// @param slice the byte slice to read from
- explicit BufferReader(Slice<const std::byte> slice)
- : data_(slice.data), bytes_remaining_(slice.len) {
- TINT_ASSERT(slice.data);
- }
-
- /// @copydoc Reader::Read
- size_t Read(std::byte* out, size_t count) override;
-
- /// @copydoc Reader::IsEOF
- bool IsEOF() const override;
-
- private:
- /// The data to read from
- const std::byte* data_ = nullptr;
-
- /// The number of bytes remaining
- size_t bytes_remaining_ = 0;
-};
-
} // namespace tint::bytes
#endif // SRC_TINT_UTILS_BYTES_READER_H_
