tint: Add support for workgroupUniformLoad

Accept any type in the intrinsics definition, and then manually
validate that there are no atomics in the type. Add manual E2E tests
for composite types.

Use the BuiltinPolyfill transform to implement it for all backends.

Update the uniformity analysis with special-case tags for the builtin.

Fixed: tint:1780
Change-Id: I95786dff4df70a0b16ed1c53b853b5d0ec6bc501
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/114862
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: James Price <jrprice@google.com>

src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc

@@ -657,6 +657,7 @@ bool WgslMutator::ReplaceFunctionCallWithBuiltin(std::string& wgsl_code) {
                                                "unpack2x16unorm",
                                                "unpack2x16float",
                                                "storageBarrier",
+                                               "workgroupUniformLoad",
                                                "workgroupBarrier"};
     wgsl_code.replace(left_bracket_pos + function_call_identifier.first,
                       function_call_identifier.second,

src/tint/intrinsics.def

@@ -275,6 +275,8 @@ match workgroup_or_storage
   | address_space.storage
 match storage
   : address_space.storage
+match workgroup
+  : address_space.workgroup
 
 ////////////////////////////////////////////////////////////////////////////////
 // Builtin Functions                                                          //
@@ -594,6 +596,7 @@ fn dot4U8Packed(u32, u32) -> u32
 @const fn unpack4x8snorm(u32) -> vec4<f32>
 @const fn unpack4x8unorm(u32) -> vec4<f32>
 @stage("compute") fn workgroupBarrier()
+@stage("compute") fn workgroupUniformLoad<T>(ptr<workgroup, T, read_write>) -> T
 
 fn textureDimensions<T: fiu32>(texture: texture_1d<T>) -> u32
 fn textureDimensions<T: fiu32, L: iu32>(texture: texture_1d<T>, level: L) -> u32

src/tint/resolver/builtin_validation_test.cc

@@ -647,5 +647,70 @@ TEST_F(ResolverDP4aExtensionValidationTest, Dot4U8PackedWithoutExtension) {
         R"(12:34 error: cannot call built-in function 'dot4U8Packed' without extension chromium_experimental_dp4a)");
 }
 
+TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_WrongAddressSpace) {
+    // @group(0) @binding(0) var<storage, read_write> v : i32;
+    // fn foo() {
+    //   workgroupUniformLoad(&v);
+    // }
+    GlobalVar("v", ty.i32(), ast::AddressSpace::kStorage, ast::Access::kReadWrite,
+              utils::Vector{Group(0_a), Binding(0_a)});
+    WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf(Source{{12, 34}}, "v"))));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(error: no matching call to workgroupUniformLoad(ptr<storage, i32, read_write>)
+
+1 candidate function:
+  workgroupUniformLoad(ptr<workgroup, T, read_write>) -> T
+)");
+}
+
+TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_Atomic) {
+    // var<workgroup> v : atomic<i32>;
+    // fn foo() {
+    //   workgroupUniformLoad(&v);
+    // }
+    GlobalVar("v", ty.atomic<i32>(), ast::AddressSpace::kWorkgroup);
+    WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf(Source{{12, 34}}, "v"))));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(12:34 error: workgroupUniformLoad must not be called with an argument that contains an atomic type)");
+}
+
+TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_AtomicInArray) {
+    // var<workgroup> v : array<atomic<i32>, 4>;
+    // fn foo() {
+    //   workgroupUniformLoad(&v);
+    // }
+    GlobalVar("v", ty.array(ty.atomic<i32>(), 4_a), ast::AddressSpace::kWorkgroup);
+    WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf(Source{{12, 34}}, "v"))));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(12:34 error: workgroupUniformLoad must not be called with an argument that contains an atomic type)");
+}
+
+TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_AtomicInStruct) {
+    // struct Inner { a : array<atomic<i32, 4> }
+    // struct S { i : Inner }
+    // var<workgroup> v : array<S, 4>;
+    // fn foo() {
+    //   workgroupUniformLoad(&v);
+    // }
+    Structure("Inner", utils::Vector{Member("a", ty.array(ty.atomic<i32>(), 4_a))});
+    Structure("S", utils::Vector{Member("i", ty.type_name("Inner"))});
+    GlobalVar(Source{{12, 34}}, "v", ty.array(ty.type_name("S"), 4_a),
+              ast::AddressSpace::kWorkgroup);
+    WrapInFunction(CallStmt(Call("workgroupUniformLoad", AddressOf("v"))));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(error: workgroupUniformLoad must not be called with an argument that contains an atomic type)");
+}
+
 }  // namespace
 }  // namespace tint::resolver

src/tint/resolver/resolver.cc

@@ -2386,6 +2386,12 @@ sem::Call* Resolver::BuiltinCall(const ast::CallExpression* expr,
         CollectTextureSamplerPairs(builtin.sem, call->Arguments());
     }
 
+    if (builtin_type == sem::BuiltinType::kWorkgroupUniformLoad) {
+        if (!validator_.WorkgroupUniformLoad(call)) {
+            return nullptr;
+        }
+    }
+
     if (!validator_.BuiltinCall(call)) {
         return nullptr;
     }

src/tint/resolver/side_effects_test.cc

@@ -176,6 +176,7 @@ TEST_P(SideEffectsBuiltinTest, Test) {
     GlobalVar("arr", ty.array<f32, 10>(), ast::AddressSpace::kPrivate);
     GlobalVar("storage_arr", ty.array<f32>(), ast::AddressSpace::kStorage, Group(0_a),
               Binding(AInt(next_binding++)));
+    GlobalVar("workgroup_arr", ty.array<f32, 4>(), ast::AddressSpace::kWorkgroup);
     GlobalVar("a", ty.atomic(ty.i32()), ast::AddressSpace::kStorage, ast::Access::kReadWrite,
               Group(0_a), Binding(AInt(next_binding++)));
     if (c.pipeline_stage != ast::PipelineStage::kCompute) {
@@ -199,6 +200,9 @@ TEST_P(SideEffectsBuiltinTest, Test) {
 
     utils::Vector<const ast::Statement*, 4> stmts;
     stmts.Push(Decl(Let("pstorage_arr", AddressOf("storage_arr"))));
+    if (c.pipeline_stage == ast::PipelineStage::kCompute) {
+        stmts.Push(Decl(Let("pworkgroup_arr", AddressOf("workgroup_arr"))));
+    }
     stmts.Push(Decl(Let("pa", AddressOf("a"))));
 
     utils::Vector<const ast::Expression*, 5> args;
@@ -339,6 +343,10 @@ INSTANTIATE_TEST_SUITE_P(
         C("atomicSub", utils::Vector{"pa", "i"}, true),                       //
         C("atomicXor", utils::Vector{"pa", "i"}, true),                       //
         C("textureStore", utils::Vector{"tstorage2d", "vi2", "vf4"}, true),   //
+        C("workgroupUniformLoad",
+          utils::Vector{"pworkgroup_arr"},
+          true,
+          ast::PipelineStage::kCompute),  //
 
         // Unimplemented builtins
         // C("quantizeToF16", utils::Vector{"f"}, false), //

src/tint/resolver/uniformity.cc

@@ -23,6 +23,7 @@
 #include "src/tint/resolver/dependency_graph.h"
 #include "src/tint/scope_stack.h"
 #include "src/tint/sem/block_statement.h"
+#include "src/tint/sem/builtin.h"
 #include "src/tint/sem/for_loop_statement.h"
 #include "src/tint/sem/function.h"
 #include "src/tint/sem/if_statement.h"
@@ -1454,6 +1455,8 @@ class UniformityGraph {
                 // some texture sampling builtins, and atomics.
                 if (builtin->IsBarrier()) {
                     callsite_tag = CallSiteRequiredToBeUniform;
+                } else if (builtin->Type() == sem::BuiltinType::kWorkgroupUniformLoad) {
+                    callsite_tag = CallSiteRequiredToBeUniform;
                 } else if (builtin->IsDerivative() ||
                            builtin->Type() == sem::BuiltinType::kTextureSample ||
                            builtin->Type() == sem::BuiltinType::kTextureSampleBias ||
@@ -1489,9 +1492,6 @@ class UniformityGraph {
                 TINT_ICE(Resolver, diagnostics_) << "unhandled function call target: " << name;
             });
 
-        if (callsite_tag == CallSiteRequiredToBeUniform) {
-            current_function_->required_to_be_uniform->AddEdge(call_node);
-        }
         cf_after->AddEdge(call_node);
 
         if (function_tag == ReturnValueMayBeNonUniform) {
@@ -1562,12 +1562,25 @@ class UniformityGraph {
                     current_function_->variables.Set(root_ident, ptr_result);
                 }
             } else {
-                // All builtin function parameters are RequiredToBeUniformForReturnValue, as are
-                // parameters for type initializers and type conversions.
-                result->AddEdge(args[i]);
+                auto* builtin = sem->Target()->As<sem::Builtin>();
+                if (builtin && builtin->Type() == sem::BuiltinType::kWorkgroupUniformLoad) {
+                    // The workgroupUniformLoad builtin requires its parameter to be uniform.
+                    current_function_->required_to_be_uniform->AddEdge(args[i]);
+                } else {
+                    // All other builtin function parameters are RequiredToBeUniformForReturnValue,
+                    // as are parameters for type initializers and type conversions.
+                    result->AddEdge(args[i]);
+                }
             }
         }
 
+        // Add the callsite requirement last.
+        // We traverse edges in reverse order, so this makes the callsite requirement take highest
+        // priority when reporting violations.
+        if (callsite_tag == CallSiteRequiredToBeUniform) {
+            current_function_->required_to_be_uniform->AddEdge(call_node);
+        }
+
         return {cf_after, result};
     }
 
@@ -1799,9 +1812,13 @@ class UniformityGraph {
 
         if (cause->type == Node::kFunctionCallArgumentValue) {
             // The requirement was on a function parameter.
-            auto param_name = NameFor(target->Parameters()[cause->arg_index]->Declaration());
+            auto* ast_param = target->Parameters()[cause->arg_index]->Declaration();
+            std::string param_name;
+            if (ast_param) {
+                param_name = " '" + NameFor(ast_param) + "'";
+            }
             report(call->args[cause->arg_index]->source,
-                   "parameter '" + param_name + "' of '" + func_name + "' must be uniform");
+                   "parameter" + param_name + " of '" + func_name + "' must be uniform");
 
             // If this is a call to a user-defined function, add a note to show the reason that the
             // parameter is required to be uniform.

src/tint/resolver/uniformity_test.cc

@@ -116,6 +116,7 @@ class BasicTest : public UniformityAnalysisTestBase,
         kUserRequiredToBeUniform,
         kWorkgroupBarrier,
         kStorageBarrier,
+        kWorkgroupUniformLoad,
         kTextureSample,
         kTextureSampleBias,
         kTextureSampleCompare,
@@ -184,6 +185,8 @@ class BasicTest : public UniformityAnalysisTestBase,
                 return "workgroupBarrier()";
             case kStorageBarrier:
                 return "storageBarrier()";
+            case kWorkgroupUniformLoad:
+                return "workgroupUniformLoad(&w)";
             case kTextureSample:
                 return "textureSample(t, s, vec2(0.5, 0.5))";
             case kTextureSampleBias:
@@ -257,6 +260,7 @@ class BasicTest : public UniformityAnalysisTestBase,
             CASE(kUserRequiredToBeUniform);
             CASE(kWorkgroupBarrier);
             CASE(kStorageBarrier);
+            CASE(kWorkgroupUniformLoad);
             CASE(kTextureSample);
             CASE(kTextureSampleBias);
             CASE(kTextureSampleCompare);
@@ -7591,6 +7595,106 @@ test:4:48 note: reading from read_write storage buffer 'arr' may result in a non
 )");
 }
 
+TEST_F(UniformityAnalysisTest, WorkgroupUniformLoad) {
+    std::string src = R"(
+const wgsize = 4;
+var<workgroup> data : array<u32, wgsize>;
+
+@compute @workgroup_size(wgsize)
+fn main(@builtin(local_invocation_index) idx : u32) {
+  data[idx] = idx + 1;
+  if (workgroupUniformLoad(&data[0]) > 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, WorkgroupUniformLoad_ViaPtrArg) {
+    std::string src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+const wgsize = 4;
+var<workgroup> data : array<u32, wgsize>;
+
+fn foo(p : ptr<workgroup, u32>) -> u32 {
+  return workgroupUniformLoad(p);
+}
+
+@compute @workgroup_size(wgsize)
+fn main(@builtin(local_invocation_index) idx : u32) {
+  data[idx] = idx + 1;
+  if (foo(&data[0]) > 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, WorkgroupUniformLoad_NonUniformPtr) {
+    std::string src = R"(
+const wgsize = 4;
+var<workgroup> data : array<u32, wgsize>;
+
+@compute @workgroup_size(wgsize)
+fn main(@builtin(local_invocation_index) idx : u32) {
+  data[idx] = idx + 1;
+  if (workgroupUniformLoad(&data[idx]) > 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, false);
+    EXPECT_EQ(error_, R"(test:8:28 warning: parameter of 'workgroupUniformLoad' must be uniform
+  if (workgroupUniformLoad(&data[idx]) > 0) {
+                           ^
+
+test:8:34 note: reading from builtin 'idx' may result in a non-uniform value
+  if (workgroupUniformLoad(&data[idx]) > 0) {
+                                 ^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, WorkgroupUniformLoad_NonUniformPtr_ViaPtrArg) {
+    std::string src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+const wgsize = 4;
+var<workgroup> data : array<u32, wgsize>;
+
+fn foo(p : ptr<workgroup, u32>) -> u32 {
+  return workgroupUniformLoad(p);
+}
+
+@compute @workgroup_size(wgsize)
+fn main(@builtin(local_invocation_index) idx : u32) {
+  data[idx] = idx + 1;
+  if (foo(&data[idx]) > 0) {
+    workgroupBarrier();
+  }
+}
+)";
+
+    RunTest(src, false);
+    EXPECT_EQ(error_, R"(test:14:11 warning: parameter 'p' of 'foo' must be uniform
+  if (foo(&data[idx]) > 0) {
+          ^
+
+test:8:31 note: parameter of 'workgroupUniformLoad' must be uniform
+  return workgroupUniformLoad(p);
+                              ^
+
+test:14:17 note: reading from builtin 'idx' may result in a non-uniform value
+  if (foo(&data[idx]) > 0) {
+                ^^^
+)");
+}
+
 TEST_F(UniformityAnalysisTest, WorkgroupAtomics) {
     std::string src = R"(
 var<workgroup> a : atomic<i32>;
@@ -7863,5 +7967,44 @@ test:17:7 note: return value of 'foo' may be non-uniform
 )");
 }
 
+TEST_F(UniformityAnalysisTest, Error_CallsiteAndParameterRequireUniformity) {
+    // Test that we report a violation for the callsite of a function when it has multiple
+    // uniformity requirements.
+    std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo(v : i32) {
+  if (v == 0) {
+    workgroupBarrier();
+  }
+}
+
+fn main() {
+  if (non_uniform == 42) {
+    foo(0);
+  }
+}
+)";
+
+    RunTest(src, false);
+    EXPECT_EQ(error_,
+              R"(test:12:5 warning: 'foo' must only be called from uniform control flow
+    foo(0);
+    ^^^
+
+test:6:5 note: 'foo' requires uniformity because it calls workgroupBarrier
+    workgroupBarrier();
+    ^^^^^^^^^^^^^^^^
+
+test:11:3 note: control flow depends on non-uniform value
+  if (non_uniform == 42) {
+  ^^
+
+test:11:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
+  if (non_uniform == 42) {
+      ^^^^^^^^^^^
+)");
+}
+
 }  // namespace
 }  // namespace tint::resolver

src/tint/resolver/validator.cc

@@ -1604,6 +1604,28 @@ bool Validator::TextureBuiltinFunction(const sem::Call* call) const {
            check_arg_is_constexpr(sem::ParameterUsage::kComponent, 0, 3);
 }
 
+bool Validator::WorkgroupUniformLoad(const sem::Call* call) const {
+    auto* builtin = call->Target()->As<sem::Builtin>();
+    if (!builtin) {
+        return false;
+    }
+
+    TINT_ASSERT(Resolver, call->Arguments().Length() > 0);
+    auto* arg = call->Arguments()[0];
+    auto* ptr = arg->Type()->As<type::Pointer>();
+    TINT_ASSERT(Resolver, ptr != nullptr);
+    auto* ty = ptr->StoreType();
+
+    if (ty->Is<type::Atomic>() || atomic_composite_info_.Contains(ty)) {
+        AddError(
+            "workgroupUniformLoad must not be called with an argument that contains an atomic type",
+            arg->Declaration()->source);
+        return false;
+    }
+
+    return true;
+}
+
 bool Validator::RequiredExtensionForBuiltinFunction(const sem::Call* call) const {
     const auto* builtin = call->Target()->As<sem::Builtin>();
     if (!builtin) {

src/tint/resolver/validator.h

@@ -439,6 +439,11 @@ class Validator {
     /// @returns true on success, false otherwise
     bool TextureBuiltinFunction(const sem::Call* call) const;
 
+    /// Validates a workgroupUniformLoad builtin function
+    /// @param call the builtin call to validate
+    /// @returns true on success, false otherwise
+    bool WorkgroupUniformLoad(const sem::Call* call) const;
+
     /// Validates an optional builtin function and its required extension.
     /// @param call the builtin call to validate
     /// @returns true on success, false otherwise

src/tint/sem/builtin.cc

@@ -171,6 +171,7 @@ bool Builtin::HasSideEffects() const {
         case sem::BuiltinType::kAtomicSub:
         case sem::BuiltinType::kAtomicXor:
         case sem::BuiltinType::kTextureStore:
+        case sem::BuiltinType::kWorkgroupUniformLoad:
             return true;
         default:
             break;

src/tint/sem/builtin_test.cc

@@ -116,7 +116,8 @@ INSTANTIATE_TEST_SUITE_P(
                     BuiltinData{"unpack2x16unorm", BuiltinType::kUnpack2X16Unorm},
                     BuiltinData{"unpack4x8snorm", BuiltinType::kUnpack4X8Snorm},
                     BuiltinData{"unpack4x8unorm", BuiltinType::kUnpack4X8Unorm},
-                    BuiltinData{"workgroupBarrier", BuiltinType::kWorkgroupBarrier}));
+                    BuiltinData{"workgroupBarrier", BuiltinType::kWorkgroupBarrier},
+                    BuiltinData{"workgroupUniformLoad", BuiltinType::kWorkgroupUniformLoad}));
 
 TEST_F(BuiltinTypeTest, ParseNoMatch) {
     EXPECT_EQ(ParseBuiltinType("not_builtin"), BuiltinType::kNone);

src/tint/sem/builtin_type.cc

@@ -285,6 +285,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
     if (name == "workgroupBarrier") {
         return BuiltinType::kWorkgroupBarrier;
     }
+    if (name == "workgroupUniformLoad") {
+        return BuiltinType::kWorkgroupUniformLoad;
+    }
     if (name == "textureDimensions") {
         return BuiltinType::kTextureDimensions;
     }
@@ -545,6 +548,8 @@ const char* str(BuiltinType i) {
             return "unpack4x8unorm";
         case BuiltinType::kWorkgroupBarrier:
             return "workgroupBarrier";
+        case BuiltinType::kWorkgroupUniformLoad:
+            return "workgroupUniformLoad";
         case BuiltinType::kTextureDimensions:
             return "textureDimensions";
         case BuiltinType::kTextureGather:

src/tint/sem/builtin_type.h

@@ -117,6 +117,7 @@ enum class BuiltinType {
     kUnpack4X8Snorm,
     kUnpack4X8Unorm,
     kWorkgroupBarrier,
+    kWorkgroupUniformLoad,
     kTextureDimensions,
     kTextureGather,
     kTextureGatherCompare,

src/tint/transform/builtin_polyfill.cc

@@ -39,7 +39,14 @@ struct BuiltinPolyfill::State {
     /// Constructor
     /// @param c the CloneContext
     /// @param p the builtins to polyfill
-    State(CloneContext& c, Builtins p) : ctx(c), polyfill(p) {}
+    State(CloneContext& c, Builtins p) : ctx(c), polyfill(p) {
+        has_full_ptr_params = false;
+        for (auto* enable : c.src->AST().Enables()) {
+            if (enable->extension == ast::Extension::kChromiumExperimentalFullPtrParameters) {
+                has_full_ptr_params = true;
+            }
+        }
+    }
 
     ////////////////////////////////////////////////////////////////////////////
     // Function polyfills
@@ -660,6 +667,29 @@ struct BuiltinPolyfill::State {
         return name;
     }
 
+    /// Builds the polyfill function for the `workgroupUniformLoad` builtin.
+    /// @param type the type being loaded
+    /// @return the polyfill function name
+    Symbol workgroupUniformLoad(const type::Type* type) {
+        if (!has_full_ptr_params) {
+            b.Enable(ast::Extension::kChromiumExperimentalFullPtrParameters);
+            has_full_ptr_params = true;
+        }
+        auto name = b.Symbols().New("tint_workgroupUniformLoad");
+        b.Func(name,
+               utils::Vector{
+                   b.Param("p", b.ty.pointer(T(type), ast::AddressSpace::kWorkgroup)),
+               },
+               T(type),
+               utils::Vector{
+                   b.CallStmt(b.Call("workgroupBarrier")),
+                   b.Decl(b.Let("result", b.Deref("p"))),
+                   b.CallStmt(b.Call("workgroupBarrier")),
+                   b.Return("result"),
+               });
+        return name;
+    }
+
     ////////////////////////////////////////////////////////////////////////////
     // Inline polyfills
     ////////////////////////////////////////////////////////////////////////////
@@ -756,6 +786,9 @@ struct BuiltinPolyfill::State {
     // Polyfill functions for binary operators.
     utils::Hashmap<BinaryOpSignature, Symbol, 8> binary_op_polyfills;
 
+    // Tracks whether the chromium_experimental_full_ptr_parameters extension has been enabled.
+    bool has_full_ptr_params;
+
     /// @returns the AST type for the given sem type
     const ast::Type* T(const type::Type* ty) const { return CreateASTTypeFor(ctx, ty); }
 
@@ -913,6 +946,13 @@ Transform::ApplyResult BuiltinPolyfill::Apply(const Program* src,
                     }
                     break;
 
+                case sem::BuiltinType::kWorkgroupUniformLoad:
+                    if (polyfill.workgroup_uniform_load) {
+                        fn = builtin_polyfills.GetOrCreate(
+                            builtin, [&] { return s.workgroupUniformLoad(builtin->ReturnType()); });
+                    }
+                    break;
+
                 default:
                     break;
             }

src/tint/transform/builtin_polyfill.h

@@ -75,6 +75,8 @@ class BuiltinPolyfill final : public Castable<BuiltinPolyfill, Transform> {
         /// Should the vector form of `quantizeToF16()` be polyfilled with a scalar implementation?
         /// See crbug.com/tint/1741
         bool quantize_to_vec_f16 = false;
+        /// Should `workgroupUniformLoad()` be polyfilled?
+        bool workgroup_uniform_load = false;
     };
 
     /// Config is consumed by the BuiltinPolyfill transform.

src/tint/transform/builtin_polyfill_test.cc

@@ -2942,6 +2942,169 @@ fn f() {
     EXPECT_EQ(expect, str(got));
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// workgroupUniformLoad
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillWorkgroupUniformLoad() {
+    BuiltinPolyfill::Builtins builtins;
+    builtins.workgroup_uniform_load = true;
+    DataMap data;
+    data.Add<BuiltinPolyfill::Config>(builtins);
+    return data;
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunWorkgroupUniformLoad) {
+    auto* src = R"(
+var<workgroup> v : i32;
+
+fn f() {
+  workgroupUniformLoad(&v);
+}
+)";
+
+    EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+    EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(src, polyfillWorkgroupUniformLoad()));
+}
+
+TEST_F(BuiltinPolyfillTest, WorkgroupUniformLoad_i32) {
+    auto* src = R"(
+var<workgroup> v : i32;
+
+fn f() {
+  let r = workgroupUniformLoad(&v);
+}
+)";
+
+    auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn tint_workgroupUniformLoad(p : ptr<workgroup, i32>) -> i32 {
+  workgroupBarrier();
+  let result = *(p);
+  workgroupBarrier();
+  return result;
+}
+
+var<workgroup> v : i32;
+
+fn f() {
+  let r = tint_workgroupUniformLoad(&(v));
+}
+)";
+
+    auto got = Run<BuiltinPolyfill>(src, polyfillWorkgroupUniformLoad());
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, WorkgroupUniformLoad_ComplexType) {
+    auto* src = R"(
+struct Inner {
+  b : bool,
+  v : vec4<i32>,
+  m : mat3x3<f32>,
+}
+
+struct Outer {
+  a : array<Inner, 4>,
+}
+
+var<workgroup> v : Outer;
+
+fn f() {
+  let r = workgroupUniformLoad(&v);
+}
+)";
+
+    auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn tint_workgroupUniformLoad(p : ptr<workgroup, Outer>) -> Outer {
+  workgroupBarrier();
+  let result = *(p);
+  workgroupBarrier();
+  return result;
+}
+
+struct Inner {
+  b : bool,
+  v : vec4<i32>,
+  m : mat3x3<f32>,
+}
+
+struct Outer {
+  a : array<Inner, 4>,
+}
+
+var<workgroup> v : Outer;
+
+fn f() {
+  let r = tint_workgroupUniformLoad(&(v));
+}
+)";
+
+    auto got = Run<BuiltinPolyfill>(src, polyfillWorkgroupUniformLoad());
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, WorkgroupUniformLoad_AvoidDuplicateEnables) {
+    auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : i32;
+var<workgroup> b : u32;
+var<workgroup> c : f32;
+
+fn f() {
+  let ra = workgroupUniformLoad(&a);
+  let rb = workgroupUniformLoad(&b);
+  let rc = workgroupUniformLoad(&c);
+}
+)";
+
+    auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn tint_workgroupUniformLoad(p : ptr<workgroup, i32>) -> i32 {
+  workgroupBarrier();
+  let result = *(p);
+  workgroupBarrier();
+  return result;
+}
+
+fn tint_workgroupUniformLoad_1(p : ptr<workgroup, u32>) -> u32 {
+  workgroupBarrier();
+  let result = *(p);
+  workgroupBarrier();
+  return result;
+}
+
+fn tint_workgroupUniformLoad_2(p : ptr<workgroup, f32>) -> f32 {
+  workgroupBarrier();
+  let result = *(p);
+  workgroupBarrier();
+  return result;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : u32;
+
+var<workgroup> c : f32;
+
+fn f() {
+  let ra = tint_workgroupUniformLoad(&(a));
+  let rb = tint_workgroupUniformLoad_1(&(b));
+  let rc = tint_workgroupUniformLoad_2(&(c));
+}
+)";
+
+    auto got = Run<BuiltinPolyfill>(src, polyfillWorkgroupUniformLoad());
+
+    EXPECT_EQ(expect, str(got));
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // quantizeToF16
 ////////////////////////////////////////////////////////////////////////////////

src/tint/writer/glsl/generator_impl.cc

@@ -200,6 +200,7 @@ SanitizedResult Sanitize(const Program* in,
         polyfills.int_div_mod = true;
         polyfills.saturate = true;
         polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
+        polyfills.workgroup_uniform_load = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

src/tint/writer/hlsl/generator_impl.cc

@@ -178,6 +178,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.insert_bits = transform::BuiltinPolyfill::Level::kFull;
         polyfills.int_div_mod = true;
         polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
+        polyfills.workgroup_uniform_load = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

src/tint/writer/msl/generator_impl.cc

@@ -183,6 +183,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.int_div_mod = true;
         polyfills.sign_int = true;
         polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
+        polyfills.workgroup_uniform_load = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

src/tint/writer/spirv/generator_impl.cc

@@ -69,6 +69,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.saturate = true;
         polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
         polyfills.quantize_to_vec_f16 = true;  // crbug.com/tint/1741
+        polyfills.workgroup_uniform_load = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }