Import Tint changes from Dawn
Changes:
- 2591c8dccc0c68bb1964a07728bb90068a945696 [tint][wgsl] Fix ICE in SingleEntryPoint transform by Ben Clayton <bclayton@google.com>
- de9f5234cc17387b62ae148a89aeaae7277c0593 [tint][core] Fix return type of HashCode() by Ben Clayton <bclayton@google.com>
- a54554a52ad83c193121115e7efe3fca21a76944 Tint/HLSL: Don't pack level in coords for textureLoad on ... by Jiawei Shao <jiawei.shao@intel.com>
- 4a85e5d207b31f69a537fa28f042fef645d26198 Tint/HLSL: Implement AST transform for pixel local variab... by Jiawei Shao <jiawei.shao@intel.com>
GitOrigin-RevId: 2591c8dccc0c68bb1964a07728bb90068a945696
Change-Id: I0fb56c94d17c10e70cb384e201e9d6963978f919
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/160681
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/core/ir/transform/direct_variable_access.cc b/src/tint/lang/core/ir/transform/direct_variable_access.cc
index a13d1b9..6a264b9 100644
--- a/src/tint/lang/core/ir/transform/direct_variable_access.cc
+++ b/src/tint/lang/core/ir/transform/direct_variable_access.cc
@@ -53,7 +53,7 @@
Var* var = nullptr;
/// @return a hash value for this object
- uint64_t HashCode() const { return Hash(var); }
+ size_t HashCode() const { return Hash(var); }
/// Inequality operator
bool operator!=(const RootModuleScopeVar& other) const { return var != other.var; }
@@ -66,7 +66,7 @@
const type::Pointer* type = nullptr;
/// @return a hash value for this object
- uint64_t HashCode() const { return Hash(type); }
+ size_t HashCode() const { return Hash(type); }
/// Inequality operator
bool operator!=(const RootPtrParameter& other) const { return type != other.type; }
@@ -81,7 +81,7 @@
const type::StructMember* member;
/// @return a hash member for this object
- uint64_t HashCode() const { return Hash(member); }
+ size_t HashCode() const { return Hash(member); }
/// Inequality operator
bool operator!=(const MemberAccess& other) const { return member != other.member; }
@@ -91,7 +91,7 @@
/// The ordered list of indices is passed by parameter.
struct IndexAccess {
/// @return a hash value for this object
- uint64_t HashCode() const { return 42; }
+ size_t HashCode() const { return 42; }
/// Inequality operator
bool operator!=(const IndexAccess&) const { return false; }
@@ -166,7 +166,7 @@
}
/// @return a hash value for this object
- uint64_t HashCode() const { return Hash(root, ops); }
+ size_t HashCode() const { return Hash(root, ops); }
/// Inequality operator
bool operator!=(const AccessShape& other) const {
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index 6ee9b5a..6f883e5 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -2814,10 +2814,16 @@
break;
case wgsl::BuiltinFn::kTextureLoad:
out << ".Load(";
- // Multisampled textures do not support mip-levels.
- if (!texture_type->Is<core::type::MultisampledTexture>()) {
- pack_level_in_coords = true;
+ // Multisampled textures and read-write storage textures do not support mip-levels.
+ if (texture_type->Is<core::type::MultisampledTexture>()) {
+ break;
}
+ if (auto* storage_texture_type = texture_type->As<core::type::StorageTexture>()) {
+ if (storage_texture_type->access() == core::Access::kReadWrite) {
+ break;
+ }
+ }
+ pack_level_in_coords = true;
break;
case wgsl::BuiltinFn::kTextureGather:
out << ".Gather";
diff --git a/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel b/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel
index 71d6a18..5ea675c 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel
+++ b/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel
@@ -43,6 +43,7 @@
"decompose_memory_access.cc",
"localize_struct_array_assignment.cc",
"num_workgroups_from_uniform.cc",
+ "pixel_local.cc",
"remove_continue_in_switch.cc",
"truncate_interstage_variables.cc",
],
@@ -51,6 +52,7 @@
"decompose_memory_access.h",
"localize_struct_array_assignment.h",
"num_workgroups_from_uniform.h",
+ "pixel_local.h",
"remove_continue_in_switch.h",
"truncate_interstage_variables.h",
],
@@ -91,6 +93,7 @@
"decompose_memory_access_test.cc",
"localize_struct_array_assignment_test.cc",
"num_workgroups_from_uniform_test.cc",
+ "pixel_local_test.cc",
"remove_continue_in_switch_test.cc",
"truncate_interstage_variables_test.cc",
],
diff --git a/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake b/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake
index 189d75a..a7a9437 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake
+++ b/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake
@@ -49,6 +49,8 @@
lang/hlsl/writer/ast_raise/localize_struct_array_assignment.h
lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.cc
lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.h
+ lang/hlsl/writer/ast_raise/pixel_local.cc
+ lang/hlsl/writer/ast_raise/pixel_local.h
lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc
lang/hlsl/writer/ast_raise/remove_continue_in_switch.h
lang/hlsl/writer/ast_raise/truncate_interstage_variables.cc
@@ -94,6 +96,7 @@
lang/hlsl/writer/ast_raise/decompose_memory_access_test.cc
lang/hlsl/writer/ast_raise/localize_struct_array_assignment_test.cc
lang/hlsl/writer/ast_raise/num_workgroups_from_uniform_test.cc
+ lang/hlsl/writer/ast_raise/pixel_local_test.cc
lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc
lang/hlsl/writer/ast_raise/truncate_interstage_variables_test.cc
)
diff --git a/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn b/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn
index 0e23db4..3200039 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn
+++ b/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn
@@ -52,6 +52,8 @@
"localize_struct_array_assignment.h",
"num_workgroups_from_uniform.cc",
"num_workgroups_from_uniform.h",
+ "pixel_local.cc",
+ "pixel_local.h",
"remove_continue_in_switch.cc",
"remove_continue_in_switch.h",
"truncate_interstage_variables.cc",
@@ -94,6 +96,7 @@
"decompose_memory_access_test.cc",
"localize_struct_array_assignment_test.cc",
"num_workgroups_from_uniform_test.cc",
+ "pixel_local_test.cc",
"remove_continue_in_switch_test.cc",
"truncate_interstage_variables_test.cc",
]
diff --git a/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
new file mode 100644
index 0000000..de4f17d
--- /dev/null
+++ b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
@@ -0,0 +1,517 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/hlsl/writer/ast_raise/pixel_local.h"
+
+#include <string>
+#include <utility>
+
+#include "src/tint/lang/core/fluent_types.h"
+#include "src/tint/lang/core/number.h"
+#include "src/tint/lang/wgsl/program/clone_context.h"
+#include "src/tint/lang/wgsl/resolver/resolve.h"
+#include "src/tint/lang/wgsl/sem/function.h"
+#include "src/tint/lang/wgsl/sem/module.h"
+#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/rtti/switch.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::PixelLocal);
+TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::PixelLocal::RasterizerOrderedView);
+TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::PixelLocal::Config);
+
+using namespace tint::core::fluent_types; // NOLINT
+
+namespace tint::hlsl::writer {
+
+/// PIMPL state for the transform
+struct PixelLocal::State {
+ /// The source program
+ const Program& src;
+ /// The semantic information of the source program
+ const sem::Info& sem;
+ /// The target program builder
+ ProgramBuilder b;
+ /// The clone context
+ program::CloneContext ctx = {&b, &src, /* auto_clone_symbols */ true};
+ /// The transform config
+ const Config& cfg;
+
+ /// Constructor
+ /// @param program the source program
+ /// @param config the transform config
+ State(const Program& program, const Config& config)
+ : src(program), sem(program.Sem()), cfg(config) {}
+
+ /// Runs the transform
+ /// @returns the new program or SkipTransform if the transform is not required
+ ApplyResult Run() {
+ // If the pixel local extension isn't enabled, then there must be no use of pixel_local
+ // variables, and so there's nothing for this transform to do.
+ if (!sem.Module()->Extensions().Contains(
+ wgsl::Extension::kChromiumExperimentalPixelLocal)) {
+ return SkipTransform;
+ }
+
+ bool made_changes = false;
+
+ // Change all module scope `var<pixel_local>` variables to `var<private>`.
+ // We need to do this even if the variable is not referenced by the entry point as later
+ // stages do not understand the pixel_local address space.
+ for (auto* global : src.AST().GlobalVariables()) {
+ if (auto* var = global->As<ast::Var>()) {
+ if (sem.Get(var)->AddressSpace() == core::AddressSpace::kPixelLocal) {
+ // Change the 'var<pixel_local>' to 'var<private>'
+ ctx.Replace(var->declared_address_space, b.Expr(core::AddressSpace::kPrivate));
+ made_changes = true;
+ }
+ }
+ }
+
+ // Find the single entry point
+ const sem::Function* entry_point = nullptr;
+ for (auto* fn : src.AST().Functions()) {
+ if (fn->IsEntryPoint()) {
+ if (entry_point != nullptr) {
+ TINT_ICE() << "PixelLocal transform requires that the SingleEntryPoint "
+ "transform has already been run";
+ return SkipTransform;
+ }
+ entry_point = sem.Get(fn);
+
+ // Look for a `var<pixel_local>` used by the entry point...
+ const tint::sem::GlobalVariable* pixel_local_variable = nullptr;
+ for (auto* global : entry_point->TransitivelyReferencedGlobals()) {
+ if (global->AddressSpace() == core::AddressSpace::kPixelLocal) {
+ pixel_local_variable = global;
+ made_changes = true;
+ break;
+ }
+ }
+ if (pixel_local_variable == nullptr) {
+ continue;
+ }
+
+ // 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);
+
+ break; // Only a single `var<pixel_local>` can be used by an entry point.
+ }
+ }
+
+ if (!made_changes) {
+ return SkipTransform;
+ }
+
+ ctx.Clone();
+ return resolver::Resolve(b);
+ }
+
+ /// Transforms the entry point @p entry_point to handle the direct or transitive usage of the
+ /// `var<pixel_local>` @p pixel_local_var.
+ /// @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) {
+ // 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();
+ auto entry_point_name = original_entry_point_fn->name->symbol.Name();
+
+ // Remove the @fragment attribute from the entry point
+ ctx.Remove(original_entry_point_fn->attributes,
+ ast::GetAttribute<ast::StageAttribute>(original_entry_point_fn->attributes));
+ // Rename the entry point.
+ auto inner_function_name = b.Symbols().New(entry_point_name + "_inner");
+ ctx.Replace(original_entry_point_fn->name, b.Ident(inner_function_name));
+
+ // Create a new function that wraps the entry point.
+ // This function has all the existing entry point parameters and an additional
+ // parameter for the input pixel local structure.
+ auto new_entry_point_params = ctx.Clone(original_entry_point_fn->params);
+
+ // Remove any entry-point attributes from the inner function.
+ // This must come after `ctx.Clone(fn->params)` as we want these attributes on the outer
+ // function.
+ for (auto* param : original_entry_point_fn->params) {
+ for (auto* attr : param->attributes) {
+ if (attr->IsAnyOf<ast::BuiltinAttribute, ast::LocationAttribute,
+ ast::InterpolateAttribute, ast::InvariantAttribute>()) {
+ ctx.Remove(param->attributes, attr);
+ }
+ }
+ }
+
+ // Declare the ROVs for the members of the pixel local variable and the functions to
+ // 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);
+
+ // Declare new entry point
+ Vector<const ast::Statement*, 5> new_entry_point_function_body;
+
+ // 1. let `hlsl_sv_position` be `@builtin(position)`
+ // Declare `@builtin(position)` in the input parameter of the new entry point if it is not
+ // declared in the original entry point.
+ auto sv_position_symbol = b.Symbols().New("hlsl_sv_position");
+ new_entry_point_function_body.Push(DeclareVariableWithBuiltinPosition(
+ new_entry_point_params, sv_position_symbol, entry_point));
+
+ // 2. Call `load_from_pixel_local_storage(hlsl_sv_position)`
+ new_entry_point_function_body.Push(
+ b.CallStmt(b.Call(load_rov_function_name, sv_position_symbol)));
+
+ // Declare the inner function
+ // Build the arguments to call the inner function
+ auto inner_function_call_args = tint::Transform(
+ original_entry_point_fn->params, [&](auto* p) { return b.Expr(ctx.Clone(p->name)); });
+
+ ast::Type new_entry_point_return_type;
+ if (original_entry_point_fn->return_type) {
+ // Create a structure to hold the combined flattened result of the entry point with
+ // `@location` attribute
+ auto new_entry_point_return_struct_name = b.Symbols().New(entry_point_name + "_res");
+ Vector<const ast::StructMember*, 8> members;
+ // arguments to the final `return` statement in the new entry point
+ Vector<const ast::Expression*, 8> new_entry_point_return_value_constructor_args;
+
+ auto add_member = [&](const core::type::Type* ty,
+ VectorRef<const ast::Attribute*> attrs) {
+ members.Push(b.Member("output_" + std::to_string(members.Length()),
+ CreateASTTypeFor(ctx, ty), std::move(attrs)));
+ };
+
+ Symbol inner_function_call_result = b.Symbols().New("result");
+ if (auto* str = entry_point->ReturnType()->As<sem::Struct>()) {
+ // The entry point returned a structure.
+ for (auto* member : str->Members()) {
+ auto& member_attrs = member->Declaration()->attributes;
+ add_member(member->Type(), ctx.Clone(member_attrs));
+ new_entry_point_return_value_constructor_args.Push(
+ b.MemberAccessor(inner_function_call_result, ctx.Clone(member->Name())));
+ if (auto* location = ast::GetAttribute<ast::LocationAttribute>(member_attrs)) {
+ // Remove the @location attribute from the member of the inner function's
+ // output structure.
+ // Note: This will break other entry points that share the same output
+ // structure, however this transform assumes that the SingleEntryPoint
+ // transform will have already been run.
+ ctx.Remove(member_attrs, location);
+ }
+ }
+ } else {
+ // The entry point returned a non-structure
+ add_member(entry_point->ReturnType(),
+ ctx.Clone(original_entry_point_fn->return_type_attributes));
+ new_entry_point_return_value_constructor_args.Push(
+ b.Expr(inner_function_call_result));
+
+ // Remove the @location from the inner function's return type attributes
+ ctx.Remove(original_entry_point_fn->return_type_attributes,
+ ast::GetAttribute<ast::LocationAttribute>(
+ original_entry_point_fn->return_type_attributes));
+ }
+
+ // 3. Call inner function and get the return value
+ new_entry_point_function_body.Push(
+ b.Decl(b.Let(inner_function_call_result,
+ b.Call(inner_function_name, std::move(inner_function_call_args)))));
+
+ // Declare the output structure
+ b.Structure(new_entry_point_return_struct_name, std::move(members));
+
+ // 4. Call `store_into_pixel_local_storage(hlsl_sv_position)`
+ new_entry_point_function_body.Push(
+ b.CallStmt(b.Call(store_rov_function_name, sv_position_symbol)));
+
+ // 5. Return the output structure
+ new_entry_point_function_body.Push(
+ b.Return(b.Call(new_entry_point_return_struct_name,
+ std::move(new_entry_point_return_value_constructor_args))));
+
+ new_entry_point_return_type = b.ty(new_entry_point_return_struct_name);
+ } else {
+ // 3. Call inner function without return value
+ new_entry_point_function_body.Push(
+ b.CallStmt(b.Call(inner_function_name, std::move(inner_function_call_args))));
+
+ // 4. Call `store_into_pixel_local_storage(hlsl_sv_position)`
+ new_entry_point_function_body.Push(
+ b.CallStmt(b.Call(store_rov_function_name, sv_position_symbol)));
+
+ new_entry_point_return_type = b.ty.void_();
+ }
+
+ // 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)});
+ }
+
+ /// Add the declarations of all the ROVs as a special type of read-write storage texture that
+ /// represent the pixel local variable, the functions to load data from them and store data into
+ /// them.
+ /// @param load_rov_function_name the name of the funtion that loads the data from the ROVs
+ /// @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) {
+ 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>())};
+
+ // 1 declaration of `rov_texcoord` and at most 4 texture[Load|Store] calls (now the maximum
+ // size of PLS is 16)
+ Vector<const ast::Statement*, 5> load_body;
+ Vector<const ast::Statement*, 5> store_body;
+
+ // let rov_texcoord = vec2u(my_input.xy);
+ auto rov_texcoord = b.Symbols().New("rov_texcoord");
+ load_body.Push(b.Decl(
+ b.Let(rov_texcoord, b.Call("vec2u", b.MemberAccessor(load_store_input_name, "xy")))));
+ store_body.Push(b.Decl(
+ b.Let(rov_texcoord, b.Call("vec2u", b.MemberAccessor(load_store_input_name, "xy")))));
+
+ for (auto* member : pixel_local_str->Members()) {
+ // Declare the read-write storage texture with RasterizerOrderedView attribute.
+ auto member_format = Switch(
+ member->Type(), //
+ [&](const core::type::U32*) { return core::TexelFormat::kR32Uint; },
+ [&](const core::type::I32*) { return core::TexelFormat::kR32Sint; },
+ [&](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);
+ 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()))),
+ b.Group(AInt(cfg.rov_group_index)), RasterizerOrderedView()});
+
+ // The function body of loading from PLS
+ // PLS_Private_Variable.member = textureLoad(pixel_local_member, rov_texcoord).x;
+ // Or
+ // PLS_Private_Variable.member =
+ // bitcast(textureLoad(pixel_local_member, rov_texcoord).x);
+ auto pixel_local_var_member_access_in_load_call =
+ b.MemberAccessor(pixel_local_variable_name, ctx.Clone(member->Name()));
+ auto load_call = b.Call(wgsl::BuiltinFn::kTextureLoad, rov_symbol_name, rov_texcoord);
+ auto to_scalar_call = b.MemberAccessor(load_call, "x");
+ if (rov_format == member_format) {
+ load_body.Push(
+ b.Assign(pixel_local_var_member_access_in_load_call, to_scalar_call));
+ } else {
+ auto member_ast_type = Switch(
+ member->Type(), //
+ [&](const core::type::U32*) { return b.ty.u32(); },
+ [&](const core::type::I32*) { return b.ty.i32(); },
+ [&](const core::type::F32*) { return b.ty.f32(); }, //
+ TINT_ICE_ON_NO_MATCH);
+
+ auto bitcast_to_member_type_call = b.Bitcast(member_ast_type, to_scalar_call);
+ load_body.Push(b.Assign(pixel_local_var_member_access_in_load_call,
+ bitcast_to_member_type_call));
+ }
+
+ // The function body of storing data into PLS
+ // textureStore(pixel_local_member, rov_texcoord, vec4u(PLS_Private_Variable.member));
+ // Or
+ // textureStore(
+ // pixel_local_member, rov_texcoord, vec4u(bitcast(PLS_Private_Variable.member)));
+ std::string rov_pixel_type;
+ switch (rov_format) {
+ case core::TexelFormat::kR32Uint:
+ rov_pixel_type = "vec4u";
+ break;
+ case core::TexelFormat::kR32Sint:
+ rov_pixel_type = "vec4i";
+ break;
+ case core::TexelFormat::kR32Float:
+ rov_pixel_type = "vec4f";
+ break;
+ default:
+ TINT_ICE() << "Invalid ROV format (now only R32Uint, R32Sint and R32Float are "
+ "supported)";
+ break;
+ }
+ auto pixel_local_var_member_access_in_store_call =
+ b.MemberAccessor(pixel_local_variable_name, ctx.Clone(member->Name()));
+ const ast::CallExpression* to_vec4_call = nullptr;
+ if (rov_format == member_format) {
+ 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) {
+ case core::TexelFormat::kR32Uint:
+ rov_pixel_ast_type = b.ty.u32();
+ break;
+ case core::TexelFormat::kR32Sint:
+ rov_pixel_ast_type = b.ty.i32();
+ break;
+ case core::TexelFormat::kR32Float:
+ rov_pixel_ast_type = b.ty.f32();
+ break;
+ default:
+ TINT_UNREACHABLE();
+ break;
+ }
+ auto bitcast_to_rov_format_call =
+ b.Bitcast(rov_pixel_ast_type, pixel_local_var_member_access_in_store_call);
+ to_vec4_call = b.Call(rov_pixel_type, bitcast_to_rov_format_call);
+ }
+ auto store_call =
+ b.Call(wgsl::BuiltinFn::kTextureStore, rov_symbol_name, rov_texcoord, to_vec4_call);
+ store_body.Push(b.CallStmt(store_call));
+ }
+
+ 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);
+ }
+
+ /// Find and get `@builtin(position)` which is needed for loading and storing data with ROVs
+ /// @returns the statement object that initializes `new_entry_point_params` with
+ /// `@builtin(position)`.
+ /// @param new_entry_point_params the input parameters of the new entry point.
+ /// `@builtin(position)` may be added if it is not in `new_entry_point_params`.
+ /// @param variable_with_position_symbol the name of the variable that will be initialized with
+ /// `@builtin(position)`.
+ /// @param entry_point the semantic information of the entry point
+ const ast::VariableDeclStatement* DeclareVariableWithBuiltinPosition(
+ tint::Vector<const ast::Parameter*, 8>& new_entry_point_params,
+ const tint::Symbol& variable_with_position_symbol,
+ const sem::Function* entry_point) {
+ for (size_t i = 0; i < entry_point->Parameters().Length(); ++i) {
+ auto* parameter = entry_point->Parameters()[i];
+ // 1. `@builtin(position)` is declared as a member of a structure
+ if (auto* struct_type = parameter->Type()->As<sem::Struct>()) {
+ for (auto* member : struct_type->Members()) {
+ if (member->Attributes().builtin == core::BuiltinValue::kPosition) {
+ return b.Decl(b.Let(
+ variable_with_position_symbol,
+ b.MemberAccessor(new_entry_point_params[i], member->Name().Name())));
+ }
+ }
+ }
+ // 2. `@builtin(position)` is declared as an individual input parameter
+ if (auto* attribute = ast::GetAttribute<ast::BuiltinAttribute>(
+ parameter->Declaration()->attributes)) {
+ auto builtin = sem.Get(attribute)->Value();
+ if (builtin == core::BuiltinValue::kPosition) {
+ return b.Decl(
+ b.Let(variable_with_position_symbol, b.Expr(new_entry_point_params[i])));
+ }
+ }
+ }
+
+ // 3. `@builtin(position)` is not declared in the input parameters and we should add one
+ auto* new_position = b.Param(b.Symbols().New("my_pos"), b.ty.vec4<f32>(),
+ Vector{b.Builtin(core::BuiltinValue::kPosition)});
+ new_entry_point_params.Push(new_position);
+ return b.Decl(b.Let(variable_with_position_symbol, b.Expr(new_position)));
+ }
+
+ /// @returns a new RasterizerOrderedView attribute
+ PixelLocal::RasterizerOrderedView* RasterizerOrderedView() {
+ return b.ASTNodes().Create<PixelLocal::RasterizerOrderedView>(b.ID(), b.AllocateNodeID());
+ }
+
+ /// @returns the register index for the pixel local field with the given index
+ /// @param field_index the pixel local field index
+ uint32_t ROVRegisterIndex(uint32_t field_index) {
+ auto idx = cfg.pls_member_to_rov_reg.Get(field_index);
+ if (TINT_UNLIKELY(!idx)) {
+ b.Diagnostics().add_error(diag::System::Transform,
+ "PixelLocal::Config::attachments missing entry for field " +
+ std::to_string(field_index));
+ return 0;
+ }
+ return *idx;
+ }
+
+ /// @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) {
+ auto format = cfg.pls_member_to_rov_format.Get(field_index);
+ if (TINT_UNLIKELY(!format)) {
+ b.Diagnostics().add_error(diag::System::Transform,
+ "PixelLocal::Config::attachments missing entry for field " +
+ std::to_string(field_index));
+ return core::TexelFormat::kUndefined;
+ }
+ return *format;
+ }
+};
+
+PixelLocal::PixelLocal() = default;
+
+PixelLocal::~PixelLocal() = default;
+
+ast::transform::Transform::ApplyResult PixelLocal::Apply(const Program& src,
+ const ast::transform::DataMap& inputs,
+ ast::transform::DataMap&) const {
+ auto* cfg = inputs.Get<Config>();
+ if (!cfg) {
+ ProgramBuilder b;
+ b.Diagnostics().add_error(diag::System::Transform,
+ "missing transform data for " + std::string(TypeInfo().name));
+ return resolver::Resolve(b);
+ }
+
+ return State(src, *cfg).Run();
+}
+
+PixelLocal::Config::Config() = default;
+
+PixelLocal::Config::Config(const Config&) = default;
+
+PixelLocal::Config::~Config() = default;
+
+PixelLocal::RasterizerOrderedView::RasterizerOrderedView(GenerationID pid, ast::NodeID nid)
+ : Base(pid, nid, Empty) {}
+
+PixelLocal::RasterizerOrderedView::~RasterizerOrderedView() = default;
+
+std::string PixelLocal::RasterizerOrderedView::InternalName() const {
+ return "rov";
+}
+
+const PixelLocal::RasterizerOrderedView* PixelLocal::RasterizerOrderedView::Clone(
+ ast::CloneContext& ctx) const {
+ return ctx.dst->ASTNodes().Create<RasterizerOrderedView>(ctx.dst->ID(),
+ ctx.dst->AllocateNodeID());
+}
+
+} // namespace tint::hlsl::writer
diff --git a/src/tint/lang/hlsl/writer/ast_raise/pixel_local.h b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.h
new file mode 100644
index 0000000..874b039
--- /dev/null
+++ b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.h
@@ -0,0 +1,116 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_HLSL_WRITER_AST_RAISE_PIXEL_LOCAL_H_
+#define SRC_TINT_LANG_HLSL_WRITER_AST_RAISE_PIXEL_LOCAL_H_
+
+#include <string>
+
+#include "src/tint/lang/core/texel_format.h"
+#include "src/tint/lang/wgsl/ast/internal_attribute.h"
+#include "src/tint/lang/wgsl/ast/transform/transform.h"
+
+namespace tint::hlsl::writer {
+
+/// PixelLocal transforms module-scope `var<pixel_local>`s and fragment entry point functions that
+/// use them:
+/// * `var<pixel_local>` will be transformed to `var<private>`, a bunch of read-write storage
+/// textures with a special `RasterizerOrderedView` attribute, a function to read the data from
+/// thes read-write storage textures into the `var<private>` (ROV Reader), and a function to write
+/// the data from the `var<private>` into the read-write storage textures (ROV Writer).
+/// * The entry point function will be wrapped with another function ('outer') that calls the
+/// 'inner' function.
+/// * The outer function will always have `@builtin(position)` in its input parameters.
+/// * The outer function will call the ROV reader before the original entry point function, and call
+/// the ROV writer after the original entry point function.
+/// @note PixelLocal requires that the SingleEntryPoint transform has already been run
+/// TODO(tint:2083): Optimize this in the future by inserting the load as late as possible and the
+/// store as early as possible.
+class PixelLocal final : public Castable<PixelLocal, ast::transform::Transform> {
+ public:
+ /// Transform configuration options
+ struct Config final : public Castable<Config, ast::transform::Data> {
+ /// Constructor
+ Config();
+
+ /// Copy Constructor
+ Config(const Config&);
+
+ /// Destructor
+ ~Config() override;
+
+ /// Index of pixel_local structure member index to ROV register
+ Hashmap<uint32_t, uint32_t, 8> pls_member_to_rov_reg;
+
+ /// Index of pixel_local structure member index to ROV format
+ Hashmap<uint32_t, core::TexelFormat, 8> pls_member_to_rov_format;
+
+ /// Group index of all ROVs
+ uint32_t rov_group_index = 0;
+ };
+
+ /// RasterizerOrderedView is an InternalAttribute that's used to decorate a read-write storage
+ /// texture object.
+ class RasterizerOrderedView final
+ : public Castable<RasterizerOrderedView, ast::InternalAttribute> {
+ public:
+ /// Constructor
+ /// @param pid the identifier of the program that owns this node
+ /// @param nid the unique node identifier
+ RasterizerOrderedView(GenerationID pid, ast::NodeID nid);
+
+ /// Destructor
+ ~RasterizerOrderedView() override;
+
+ /// @return a short description of the internal attribute which will be
+ /// displayed as `@internal(<name>)`
+ std::string InternalName() const override;
+
+ /// Performs a deep clone of this object using the program::CloneContext `ctx`.
+ /// @param ctx the clone context
+ /// @return the newly cloned object
+ const RasterizerOrderedView* Clone(ast::CloneContext& ctx) const override;
+ };
+
+ /// Constructor
+ PixelLocal();
+
+ /// Destructor
+ ~PixelLocal() override;
+
+ /// @copydoc ast::transform::Transform::Apply
+ ApplyResult Apply(const Program& program,
+ const ast::transform::DataMap& inputs,
+ ast::transform::DataMap& outputs) const override;
+
+ private:
+ struct State;
+};
+
+} // namespace tint::hlsl::writer
+
+#endif // SRC_TINT_LANG_HLSL_WRITER_AST_RAISE_PIXEL_LOCAL_H_
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
new file mode 100644
index 0000000..9f644a7
--- /dev/null
+++ b/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc
@@ -0,0 +1,1287 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include <utility>
+
+#include "src/tint/lang/hlsl/writer/ast_raise/pixel_local.h"
+
+#include "src/tint/lang/wgsl/ast/transform/helper_test.h"
+
+namespace tint::hlsl::writer {
+namespace {
+
+struct ROVBinding {
+ uint32_t field_index;
+ uint32_t register_index;
+ core::TexelFormat rov_format;
+};
+
+ast::transform::DataMap Bindings(std::initializer_list<ROVBinding> bindings,
+ uint32_t groupIndex = 0) {
+ PixelLocal::Config cfg;
+ cfg.rov_group_index = groupIndex;
+ for (auto& binding : bindings) {
+ cfg.pls_member_to_rov_reg.Add(binding.field_index, binding.register_index);
+ cfg.pls_member_to_rov_format.Add(binding.field_index, binding.rov_format);
+ }
+ ast::transform::DataMap data;
+ data.Add<PixelLocal::Config>(std::move(cfg));
+ return data;
+}
+
+using HLSLPixelLocalTest = ast::transform::TransformTest;
+
+TEST_F(HLSLPixelLocalTest, EmptyModule) {
+ auto* src = "";
+
+ EXPECT_FALSE(ShouldRun<PixelLocal>(src, Bindings({})));
+}
+
+TEST_F(HLSLPixelLocalTest, UseInEntryPoint_NoPosition) {
+ 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"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+struct F_res {
+ @location(0)
+ output_0 : vec4<f32>,
+}
+
+@fragment
+fn F(@builtin(position) my_pos : vec4<f32>) -> F_res {
+ let hlsl_sv_position = my_pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ let result = F_inner();
+ store_into_pixel_local_storage(hlsl_sv_position);
+ return F_res(result);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn F_inner() -> vec4f {
+ P.a += 42;
+ return vec4f(1, 0, 0, 1);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, UseInEntryPoint_SeparatePosition) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+@fragment
+fn F(@builtin(position) pos : vec4f) -> @location(0) vec4f {
+ P.a += 42;
+ return pos;
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+struct F_res {
+ @location(0)
+ output_0 : vec4<f32>,
+}
+
+@fragment
+fn F(@builtin(position) pos : vec4f) -> F_res {
+ let hlsl_sv_position = pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ let result = F_inner(pos);
+ store_into_pixel_local_storage(hlsl_sv_position);
+ return F_res(result);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn F_inner(pos : vec4f) -> vec4f {
+ P.a += 42;
+ return pos;
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, UseInEntryPoint_PositionInStruct) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0) input : vec4f,
+ @builtin(position) pos : vec4f,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> @location(0) vec4f {
+ P.a += 42;
+ return fragmentInput.input + fragmentInput.pos;
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+struct F_res {
+ @location(0)
+ output_0 : vec4<f32>,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> F_res {
+ let hlsl_sv_position = fragmentInput.pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ let result = F_inner(fragmentInput);
+ store_into_pixel_local_storage(hlsl_sv_position);
+ return F_res(result);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0)
+ input : vec4f,
+ @builtin(position)
+ pos : vec4f,
+}
+
+fn F_inner(fragmentInput : FragmentInput) -> vec4f {
+ P.a += 42;
+ return (fragmentInput.input + fragmentInput.pos);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, UseInEntryPoint_NonZeroROVGroupIndex) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0) input : vec4f,
+ @builtin(position) pos : vec4f,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> @location(0) vec4f {
+ P.a += 42;
+ return fragmentInput.input + fragmentInput.pos;
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(3) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+struct F_res {
+ @location(0)
+ output_0 : vec4<f32>,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> F_res {
+ let hlsl_sv_position = fragmentInput.pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ let result = F_inner(fragmentInput);
+ store_into_pixel_local_storage(hlsl_sv_position);
+ return F_res(result);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0)
+ input : vec4f,
+ @builtin(position)
+ pos : vec4f,
+}
+
+fn F_inner(fragmentInput : FragmentInput) -> vec4f {
+ P.a += 42;
+ return (fragmentInput.input + fragmentInput.pos);
+}
+)";
+
+ constexpr uint32_t kROVGroupIndex = 3u;
+ auto got =
+ Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}, kROVGroupIndex));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, UseInEntryPoint_Multiple_PixelLocal_Members) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+ b : i32,
+ c : f32,
+ d : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0) input : vec4f,
+ @builtin(position) pos : vec4f,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> @location(0) vec4f {
+ P.a += 42;
+ P.b -= 21;
+ P.c += 12.5f;
+ P.d -= 5;
+ return fragmentInput.input + fragmentInput.pos;
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+@binding(2) @group(0) @internal(rov) var pixel_local_b : texture_storage_2d<r32sint, read_write>;
+
+@binding(3) @group(0) @internal(rov) var pixel_local_c : texture_storage_2d<r32float, read_write>;
+
+@binding(4) @group(0) @internal(rov) var pixel_local_d : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+ P.b = textureLoad(pixel_local_b, rov_texcoord).x;
+ P.c = textureLoad(pixel_local_c, rov_texcoord).x;
+ P.d = textureLoad(pixel_local_d, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+ textureStore(pixel_local_b, rov_texcoord, vec4i(P.b));
+ textureStore(pixel_local_c, rov_texcoord, vec4f(P.c));
+ textureStore(pixel_local_d, rov_texcoord, vec4u(P.d));
+}
+
+struct F_res {
+ @location(0)
+ output_0 : vec4<f32>,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> F_res {
+ let hlsl_sv_position = fragmentInput.pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ let result = F_inner(fragmentInput);
+ store_into_pixel_local_storage(hlsl_sv_position);
+ return F_res(result);
+}
+
+struct PixelLocal {
+ a : u32,
+ b : i32,
+ c : f32,
+ d : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0)
+ input : vec4f,
+ @builtin(position)
+ pos : vec4f,
+}
+
+fn F_inner(fragmentInput : FragmentInput) -> vec4f {
+ P.a += 42;
+ P.b -= 21;
+ P.c += 12.5f;
+ P.d -= 5;
+ return (fragmentInput.input + fragmentInput.pos);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint},
+ {1, 2, core::TexelFormat::kR32Sint},
+ {2, 3, core::TexelFormat::kR32Float},
+ {3, 4, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, UseInEntryPoint_Multiple_PixelLocal_Members_And_Fragment_Output) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+ b : i32,
+ c : f32,
+ d : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0) input : vec4f,
+ @builtin(position) pos : vec4f,
+}
+
+struct FragmentOutput {
+ @location(0) color0 : vec4f,
+ @location(1) color1 : vec4f,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> FragmentOutput {
+ P.a += 42;
+ P.b -= 21;
+ P.c += 12.5f;
+ P.d -= 5;
+ return FragmentOutput(fragmentInput.input, fragmentInput.pos);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+@binding(2) @group(0) @internal(rov) var pixel_local_b : texture_storage_2d<r32sint, read_write>;
+
+@binding(3) @group(0) @internal(rov) var pixel_local_c : texture_storage_2d<r32float, read_write>;
+
+@binding(4) @group(0) @internal(rov) var pixel_local_d : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+ P.b = textureLoad(pixel_local_b, rov_texcoord).x;
+ P.c = textureLoad(pixel_local_c, rov_texcoord).x;
+ P.d = textureLoad(pixel_local_d, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+ textureStore(pixel_local_b, rov_texcoord, vec4i(P.b));
+ textureStore(pixel_local_c, rov_texcoord, vec4f(P.c));
+ textureStore(pixel_local_d, rov_texcoord, vec4u(P.d));
+}
+
+struct F_res {
+ @location(0)
+ output_0 : vec4<f32>,
+ @location(1)
+ output_1 : vec4<f32>,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> F_res {
+ let hlsl_sv_position = fragmentInput.pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ let result = F_inner(fragmentInput);
+ store_into_pixel_local_storage(hlsl_sv_position);
+ return F_res(result.color0, result.color1);
+}
+
+struct PixelLocal {
+ a : u32,
+ b : i32,
+ c : f32,
+ d : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0)
+ input : vec4f,
+ @builtin(position)
+ pos : vec4f,
+}
+
+struct FragmentOutput {
+ color0 : vec4f,
+ color1 : vec4f,
+}
+
+fn F_inner(fragmentInput : FragmentInput) -> FragmentOutput {
+ P.a += 42;
+ P.b -= 21;
+ P.c += 12.5f;
+ P.d -= 5;
+ return FragmentOutput(fragmentInput.input, fragmentInput.pos);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint},
+ {1, 2, core::TexelFormat::kR32Sint},
+ {2, 3, core::TexelFormat::kR32Float},
+ {3, 4, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, ROV_Format_Different_From_Pixel_Local_Member_Format) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+ b : i32,
+ c : f32,
+ d : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0) input : vec4f,
+ @builtin(position) pos : vec4f,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> @location(0) vec4f {
+ P.a += 42;
+ P.b -= 21;
+ P.c += 12.5f;
+ P.d -= 5;
+ return fragmentInput.input + fragmentInput.pos;
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32sint, read_write>;
+
+@binding(2) @group(0) @internal(rov) var pixel_local_b : texture_storage_2d<r32float, read_write>;
+
+@binding(3) @group(0) @internal(rov) var pixel_local_c : texture_storage_2d<r32uint, read_write>;
+
+@binding(4) @group(0) @internal(rov) var pixel_local_d : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = bitcast<u32>(textureLoad(pixel_local_a, rov_texcoord).x);
+ P.b = bitcast<i32>(textureLoad(pixel_local_b, rov_texcoord).x);
+ P.c = bitcast<f32>(textureLoad(pixel_local_c, rov_texcoord).x);
+ P.d = textureLoad(pixel_local_d, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4i(bitcast<i32>(P.a)));
+ textureStore(pixel_local_b, rov_texcoord, vec4f(bitcast<f32>(P.b)));
+ textureStore(pixel_local_c, rov_texcoord, vec4u(bitcast<u32>(P.c)));
+ textureStore(pixel_local_d, rov_texcoord, vec4u(P.d));
+}
+
+struct F_res {
+ @location(0)
+ output_0 : vec4<f32>,
+}
+
+@fragment
+fn F(fragmentInput : FragmentInput) -> F_res {
+ let hlsl_sv_position = fragmentInput.pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ let result = F_inner(fragmentInput);
+ store_into_pixel_local_storage(hlsl_sv_position);
+ return F_res(result);
+}
+
+struct PixelLocal {
+ a : u32,
+ b : i32,
+ c : f32,
+ d : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct FragmentInput {
+ @location(0)
+ input : vec4f,
+ @builtin(position)
+ pos : vec4f,
+}
+
+fn F_inner(fragmentInput : FragmentInput) -> vec4f {
+ P.a += 42;
+ P.b -= 21;
+ P.c += 12.5f;
+ P.d -= 5;
+ return (fragmentInput.input + fragmentInput.pos);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Sint},
+ {1, 2, core::TexelFormat::kR32Float},
+ {2, 3, core::TexelFormat::kR32Uint},
+ {3, 4, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, UseInCallee) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+fn X() {
+ P.a += 42;
+}
+
+@fragment
+fn F() {
+ X();
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(@builtin(position) my_pos : vec4<f32>) {
+ let hlsl_sv_position = my_pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner();
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn X() {
+ P.a += 42;
+}
+
+fn F_inner() {
+ X();
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, ROVLoadFunctionNameUsed) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ P.a += u32(my_input.x);
+}
+
+@fragment
+fn F() {
+ load_from_pixel_local_storage(vec4f(1.0, 0.0, 0.0, 1.0));
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage_1(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(@builtin(position) my_pos : vec4<f32>) {
+ let hlsl_sv_position = my_pos;
+ load_from_pixel_local_storage_1(hlsl_sv_position);
+ F_inner();
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ P.a += u32(my_input.x);
+}
+
+fn F_inner() {
+ load_from_pixel_local_storage(vec4f(1.0, 0.0, 0.0, 1.0));
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, ROVStoreFunctionNameUsed) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ P.a += u32(my_input.x);
+}
+
+@fragment
+fn F() {
+ store_into_pixel_local_storage(vec4f(1.0, 0.0, 0.0, 1.0));
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage_1(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(@builtin(position) my_pos : vec4<f32>) {
+ let hlsl_sv_position = my_pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner();
+ store_into_pixel_local_storage_1(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ P.a += u32(my_input.x);
+}
+
+fn F_inner() {
+ store_into_pixel_local_storage(vec4f(1.0, 0.0, 0.0, 1.0));
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, NewBuiltinPositionVariableNamesUsed) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+@fragment
+fn F(@location(0) my_pos : vec4f) {
+ let hlsl_sv_position = my_pos * 2;
+ P.a += u32(hlsl_sv_position.y);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(@location(0) my_pos : vec4f, @builtin(position) my_pos_1 : vec4<f32>) {
+ let hlsl_sv_position_1 = my_pos_1;
+ load_from_pixel_local_storage(hlsl_sv_position_1);
+ F_inner(my_pos);
+ store_into_pixel_local_storage(hlsl_sv_position_1);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn F_inner(my_pos : vec4f) {
+ let hlsl_sv_position = (my_pos * 2);
+ P.a += u32(hlsl_sv_position.y);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, WithInvariantBuiltinInputParameter) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+@fragment
+fn F(@invariant @builtin(position) pos : vec4f) {
+ P.a += u32(pos.x);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(@invariant @builtin(position) pos : vec4f) {
+ let hlsl_sv_position = pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner(pos);
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn F_inner(pos : vec4f) {
+ P.a += u32(pos.x);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, WithInvariantBuiltinInputStructParameter) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct In {
+ @invariant @builtin(position) pos : vec4f,
+}
+
+@fragment
+fn F(in : In) {
+ P.a += u32(in.pos.x);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(in : In) {
+ let hlsl_sv_position = in.pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner(in);
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct In {
+ @invariant @builtin(position)
+ pos : vec4f,
+}
+
+fn F_inner(in : In) {
+ P.a += u32(in.pos.x);
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, WithLocationInputParameter) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+@fragment
+fn F(@location(0) a : vec4f, @interpolate(flat) @location(1) b : vec4f) {
+ P.a += u32(a.x) + u32(b.y);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(@location(0) a : vec4f, @interpolate(flat) @location(1) b : vec4f, @builtin(position) my_pos : vec4<f32>) {
+ let hlsl_sv_position = my_pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner(a, b);
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn F_inner(a : vec4f, b : vec4f) {
+ P.a += (u32(a.x) + u32(b.y));
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, WithLocationInputStructParameter) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct In {
+ @location(0) a : vec4f,
+ @interpolate(flat) @location(1) b : vec4f,
+}
+
+@fragment
+fn F(in : In) {
+ P.a += u32(in.a.x) + u32(in.b.y);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(in : In, @builtin(position) my_pos : vec4<f32>) {
+ let hlsl_sv_position = my_pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner(in);
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct In {
+ @location(0)
+ a : vec4f,
+ @interpolate(flat) @location(1)
+ b : vec4f,
+}
+
+fn F_inner(in : In) {
+ P.a += (u32(in.a.x) + u32(in.b.y));
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, WithBuiltinAndLocationInputParameter) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+@fragment
+fn F(@builtin(position) pos : vec4f, @location(0) uv : vec4f) {
+ P.a += u32(pos.x) + u32(uv.x);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(@builtin(position) pos : vec4f, @location(0) uv : vec4f) {
+ let hlsl_sv_position = pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner(pos, uv);
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+fn F_inner(pos : vec4f, uv : vec4f) {
+ P.a += (u32(pos.x) + u32(uv.x));
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(HLSLPixelLocalTest, WithBuiltinAndLocationInputStructParameter) {
+ auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+struct In {
+ @builtin(position) pos : vec4f,
+ @location(0) uv : vec4f,
+}
+
+@fragment
+fn F(in : In) {
+ P.a += u32(in.pos.x) + u32(in.uv.x);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_pixel_local;
+
+@binding(1) @group(0) @internal(rov) var pixel_local_a : texture_storage_2d<r32uint, read_write>;
+
+fn load_from_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ P.a = textureLoad(pixel_local_a, rov_texcoord).x;
+}
+
+fn store_into_pixel_local_storage(my_input : vec4<f32>) {
+ let rov_texcoord = vec2u(my_input.xy);
+ textureStore(pixel_local_a, rov_texcoord, vec4u(P.a));
+}
+
+@fragment
+fn F(in : In) {
+ let hlsl_sv_position = in.pos;
+ load_from_pixel_local_storage(hlsl_sv_position);
+ F_inner(in);
+ store_into_pixel_local_storage(hlsl_sv_position);
+}
+
+struct PixelLocal {
+ a : u32,
+}
+
+var<private> P : PixelLocal;
+
+struct In {
+ @builtin(position)
+ pos : vec4f,
+ @location(0)
+ uv : vec4f,
+}
+
+fn F_inner(in : In) {
+ P.a += (u32(in.pos.x) + u32(in.uv.x));
+}
+)";
+
+ auto got = Run<PixelLocal>(src, Bindings({{0, 1, core::TexelFormat::kR32Uint}}));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+} // namespace
+} // namespace tint::hlsl::writer
diff --git a/src/tint/lang/wgsl/ast/transform/single_entry_point.cc b/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
index 3311ce4..53b00d0 100644
--- a/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
+++ b/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
@@ -132,6 +132,7 @@
// Drop requires directives as they are optional, and it's non-trivial to determine
// which features are needed for which entry points.
},
+ [&](const ConstAssert* ca) { b.AST().AddConstAssert(ctx.Clone(ca)); },
[&](const DiagnosticDirective* d) { b.AST().AddDiagnosticDirective(ctx.Clone(d)); }, //
TINT_ICE_ON_NO_MATCH);
}
diff --git a/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc b/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc
index f22ddb5..4d4ede3 100644
--- a/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc
@@ -669,5 +669,50 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(SingleEntryPointTest, ConstAssert_ModuleScope) {
+ // module scope const_assert is preserved
+ auto* src = R"(
+const C = 42;
+
+const_assert (C == 42);
+
+@compute @workgroup_size(1)
+fn main() {
+}
+)";
+
+ auto* expect = src;
+
+ SingleEntryPoint::Config cfg("main");
+
+ DataMap data;
+ data.Add<SingleEntryPoint::Config>(cfg);
+ auto got = Run<SingleEntryPoint>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(SingleEntryPointTest, ConstAssert_FnScope) {
+ // function scope const_assert is preserved
+ auto* src = R"(
+const C = 42;
+
+@compute @workgroup_size(1)
+fn main() {
+ const_assert (C == 42);
+}
+)";
+
+ auto* expect = src;
+
+ SingleEntryPoint::Config cfg("main");
+
+ DataMap data;
+ data.Add<SingleEntryPoint::Config>(cfg);
+ auto got = Run<SingleEntryPoint>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace tint::ast::transform
diff --git a/src/tint/utils/math/hash.h b/src/tint/utils/math/hash.h
index ac6bce9..e143d07 100644
--- a/src/tint/utils/math/hash.h
+++ b/src/tint/utils/math/hash.h
@@ -103,7 +103,10 @@
/// @returns a hash of the value
size_t operator()(const T& value) const {
if constexpr (detail::HasHashCodeMember<T>::value) {
- return value.HashCode();
+ auto hash = value.HashCode();
+ static_assert(std::is_same_v<decltype(hash), size_t>,
+ "T::HashCode() must return size_t");
+ return hash;
} else {
return std::hash<T>()(value);
}
