Implement barrier intrinsics

Fixed: tint:658
Change-Id: I28d5225f42dacb2b6b0cb51ce9f15951b31f0fc9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/45284
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: Alan Baker <alanbaker@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

src/CMakeLists.txt

@@ -548,6 +548,7 @@ if(${TINT_BUILD_TESTS})
       reader/spirv/function_misc_test.cc
       reader/spirv/function_var_test.cc
       reader/spirv/namer_test.cc
+      reader/spirv/parser_impl_barrier_test.cc
       reader/spirv/parser_impl_convert_member_decoration_test.cc
       reader/spirv/parser_impl_convert_type_test.cc
       reader/spirv/parser_impl_function_decl_test.cc

src/intrinsic_table.cc

@@ -1055,6 +1055,7 @@ Impl::Impl() {
   Register(I::kSqrt,            vecN_f32,    {vecN_f32}                                               ); // NOLINT
   Register(I::kStep,            f32,         {f32, f32}                                               ); // NOLINT
   Register(I::kStep,            vecN_f32,    {vecN_f32, vecN_f32}                                     ); // NOLINT
+  Register(I::kStorageBarrier,  void_,       {}                                                       ); // NOLINT
   Register(I::kTan,             f32,         {f32}                                                    ); // NOLINT
   Register(I::kTan,             vecN_f32,    {vecN_f32}                                               ); // NOLINT
   Register(I::kTanh,            f32,         {f32}                                                    ); // NOLINT
@@ -1066,6 +1067,7 @@ Impl::Impl() {
   Register(I::kUnpack2x16Unorm, vec2_f32,    {u32}                                                    ); // NOLINT
   Register(I::kUnpack4x8Snorm,  vec4_f32,    {u32}                                                    ); // NOLINT
   Register(I::kUnpack4x8Unorm,  vec4_f32,    {u32}                                                    ); // NOLINT
+  Register(I::kWorkgroupBarrier,void_,       {}                                                       ); // NOLINT
   // clang-format on
 
   auto* tex_1d_f32 = sampled_texture(Dim::k1d, f32);

src/reader/spirv/function.cc

@@ -3104,6 +3104,9 @@ bool FunctionEmitter::EmitStatement(const spvtools::opt::Instruction& inst) {
     case SpvOpFunctionCall:
       return EmitFunctionCall(inst);
 
+    case SpvOpControlBarrier:
+      return EmitControlBarrier(inst);
+
     case SpvOpExtInst:
       if (parser_impl_.IsIgnoredExtendedInstruction(inst)) {
         return true;
@@ -4133,6 +4136,53 @@ bool FunctionEmitter::EmitFunctionCall(const spvtools::opt::Instruction& inst) {
   return EmitConstDefOrWriteToHoistedVar(inst, {result_type, call_expr});
 }
 
+bool FunctionEmitter::EmitControlBarrier(
+    const spvtools::opt::Instruction& inst) {
+  uint32_t operands[3];
+  for (int i = 0; i < 3; i++) {
+    if (auto* op = MakeOperand(inst, i).expr) {
+      auto* lit = As<ast::ScalarConstructorExpression>(op)->literal();
+      if (auto* int_lit = lit->As<ast::IntLiteral>()) {
+        operands[i] = int_lit->value_as_u32();
+        continue;
+      }
+    }
+    return Fail() << "invalid or missing operands for control barrier";
+  }
+
+  uint32_t execution = operands[0];
+  uint32_t memory = operands[1];
+  uint32_t semantics = operands[2];
+
+  if (execution != SpvScopeWorkgroup) {
+    return Fail() << "unsupported control barrier execution scope: "
+                  << "expected Workgroup (2), got: " << execution;
+  }
+  if (semantics & SpvMemorySemanticsAcquireReleaseMask) {
+    semantics &= ~SpvMemorySemanticsAcquireReleaseMask;
+  } else {
+    return Fail() << "control barrier semantics requires acquire and release";
+  }
+  if (semantics & SpvMemorySemanticsWorkgroupMemoryMask) {
+    if (memory != SpvScopeWorkgroup) {
+      return Fail() << "workgroupBarrier requires workgroup memory scope";
+    }
+    AddStatement(create<ast::CallStatement>(builder_.Call("workgroupBarrier")));
+    semantics &= ~SpvMemorySemanticsWorkgroupMemoryMask;
+  }
+  if (semantics & SpvMemorySemanticsUniformMemoryMask) {
+    if (memory != SpvScopeDevice) {
+      return Fail() << "storageBarrier requires device memory scope";
+    }
+    AddStatement(create<ast::CallStatement>(builder_.Call("storageBarrier")));
+    semantics &= ~SpvMemorySemanticsUniformMemoryMask;
+  }
+  if (semantics) {
+    return Fail() << "unsupported control barrier semantics: " << semantics;
+  }
+  return true;
+}
+
 TypedExpression FunctionEmitter::MakeIntrinsicCall(
     const spvtools::opt::Instruction& inst) {
   const auto intrinsic = GetIntrinsic(inst.opcode());

src/reader/spirv/function.h

@@ -883,6 +883,12 @@ class FunctionEmitter {
   /// @returns false if emission failed
   bool EmitFunctionCall(const spvtools::opt::Instruction& inst);
 
+  /// Emits a control barrier intrinsic.  On failure, emits a diagnostic and
+  /// returns false.
+  /// @param inst the SPIR-V control barrier instruction
+  /// @returns false if emission failed
+  bool EmitControlBarrier(const spvtools::opt::Instruction& inst);
+
   /// Returns an expression for a SPIR-V instruction that maps to a WGSL
   /// intrinsic function call.
   /// @param inst the SPIR-V instruction

src/reader/spirv/parser_impl_barrier_test.cc

@@ -0,0 +1,201 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "gmock/gmock.h"
+#include "src/ast/call_statement.h"
+#include "src/reader/spirv/function.h"
+#include "src/reader/spirv/parser_impl_test_helper.h"
+#include "src/reader/spirv/spirv_tools_helpers_test.h"
+#include "src/semantic/call.h"
+
+namespace tint {
+namespace reader {
+namespace spirv {
+namespace {
+
+using ::testing::Eq;
+using ::testing::HasSubstr;
+using ::testing::Not;
+using ::testing::StartsWith;
+
+Program ParseAndBuild(std::string spirv) {
+  const char* preamble = R"(OpCapability Shader
+            OpMemoryModel Logical GLSL450
+            OpEntryPoint GLCompute %main "main"
+            OpExecutionMode %main LocalSize 1 1 1
+            OpName %main "main"
+)";
+
+  auto p = std::make_unique<ParserImpl>(test::Assemble(preamble + spirv));
+  if (!p->BuildAndParseInternalModule()) {
+    ProgramBuilder builder;
+    builder.Diagnostics().add_error(p->error());
+    return Program(std::move(builder));
+  }
+  return p->program();
+}
+
+TEST_F(SpvParserTest, WorkgroupBarrier) {
+  auto program = ParseAndBuild(R"(
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpReturn
+               OpFunctionEnd
+  )");
+  ASSERT_TRUE(program.IsValid()) << program.Diagnostics().str();
+  auto* main = program.AST().Functions().Find(program.Symbols().Get("main"));
+  ASSERT_NE(main, nullptr);
+  ASSERT_GT(main->body()->size(), 0u);
+  auto* call = main->body()->get(0)->As<ast::CallStatement>();
+  ASSERT_NE(call, nullptr);
+  EXPECT_EQ(call->expr()->params().size(), 0u);
+  auto* sem_call = program.Sem().Get(call->expr());
+  ASSERT_NE(sem_call, nullptr);
+  auto* intrinsic = sem_call->Target()->As<semantic::Intrinsic>();
+  ASSERT_NE(intrinsic, nullptr);
+  EXPECT_EQ(intrinsic->Type(), semantic::IntrinsicType::kWorkgroupBarrier);
+}
+
+TEST_F(SpvParserTest, StorageBarrier) {
+  auto program = ParseAndBuild(R"(
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+    %uint_72 = OpConstant %uint 72
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpControlBarrier %uint_2 %uint_1 %uint_72
+               OpReturn
+               OpFunctionEnd
+  )");
+  ASSERT_TRUE(program.IsValid()) << program.Diagnostics().str();
+  auto* main = program.AST().Functions().Find(program.Symbols().Get("main"));
+  ASSERT_NE(main, nullptr);
+  ASSERT_GT(main->body()->size(), 0u);
+  auto* call = main->body()->get(0)->As<ast::CallStatement>();
+  ASSERT_NE(call, nullptr);
+  EXPECT_EQ(call->expr()->params().size(), 0u);
+  auto* sem_call = program.Sem().Get(call->expr());
+  ASSERT_NE(sem_call, nullptr);
+  auto* intrinsic = sem_call->Target()->As<semantic::Intrinsic>();
+  ASSERT_NE(intrinsic, nullptr);
+  EXPECT_EQ(intrinsic->Type(), semantic::IntrinsicType::kStorageBarrier);
+}
+
+TEST_F(SpvParserTest, ErrBarrierInvalidExecution) {
+  auto program = ParseAndBuild(R"(
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpControlBarrier %uint_0 %uint_2 %uint_264
+               OpReturn
+               OpFunctionEnd
+  )");
+  EXPECT_FALSE(program.IsValid());
+  EXPECT_THAT(program.Diagnostics().str(),
+              HasSubstr("unsupported control barrier execution scope"));
+}
+
+TEST_F(SpvParserTest, ErrBarrierSemanticsMissingAcquireRelease) {
+  auto program = ParseAndBuild(R"(
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_0
+               OpReturn
+               OpFunctionEnd
+  )");
+  EXPECT_FALSE(program.IsValid());
+  EXPECT_THAT(
+      program.Diagnostics().str(),
+      HasSubstr("control barrier semantics requires acquire and release"));
+}
+
+TEST_F(SpvParserTest, ErrBarrierInvalidSemantics) {
+  auto program = ParseAndBuild(R"(
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+     %uint_9 = OpConstant %uint 9
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_9
+               OpReturn
+               OpFunctionEnd
+  )");
+  EXPECT_FALSE(program.IsValid());
+  EXPECT_THAT(program.Diagnostics().str(),
+              HasSubstr("unsupported control barrier semantics"));
+}
+
+TEST_F(SpvParserTest, ErrWorkgroupBarrierInvalidMemory) {
+  auto program = ParseAndBuild(R"(
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+     %uint_8 = OpConstant %uint 8
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpControlBarrier %uint_2 %uint_8 %uint_264
+               OpReturn
+               OpFunctionEnd
+  )");
+  EXPECT_FALSE(program.IsValid());
+  EXPECT_THAT(program.Diagnostics().str(),
+              HasSubstr("workgroupBarrier requires workgroup memory scope"));
+}
+
+TEST_F(SpvParserTest, ErrStorageBarrierInvalidMemory) {
+  auto program = ParseAndBuild(R"(
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+     %uint_8 = OpConstant %uint 8
+    %uint_72 = OpConstant %uint 72
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpControlBarrier %uint_2 %uint_8 %uint_72
+               OpReturn
+               OpFunctionEnd
+  )");
+  EXPECT_FALSE(program.IsValid());
+  EXPECT_THAT(program.Diagnostics().str(),
+              HasSubstr("storageBarrier requires device memory scope"));
+}
+
+}  // namespace
+}  // namespace spirv
+}  // namespace reader
+}  // namespace tint

src/resolver/intrinsic_test.cc

@@ -538,6 +538,37 @@ inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
   return out;
 }
 
+using ResolverIntrinsicTest_Barrier = ResolverTestWithParam<IntrinsicData>;
+TEST_P(ResolverIntrinsicTest_Barrier, InferType) {
+  auto param = GetParam();
+
+  auto* call = Call(param.name);
+  WrapInFunction(call);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_NE(TypeOf(call), nullptr);
+  EXPECT_TRUE(TypeOf(call)->Is<type::Void>());
+}
+
+TEST_P(ResolverIntrinsicTest_Barrier, Error_TooManyParams) {
+  auto param = GetParam();
+
+  auto* call = Call(param.name, vec4<f32>(1.f, 2.f, 3.f, 4.f), 1.0f);
+  WrapInFunction(call);
+
+  EXPECT_FALSE(r()->Resolve());
+
+  EXPECT_THAT(r()->error(), HasSubstr("error: no matching call to " +
+                                      std::string(param.name)));
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    ResolverTest,
+    ResolverIntrinsicTest_Barrier,
+    testing::Values(
+        IntrinsicData{"storageBarrier", IntrinsicType::kStorageBarrier},
+        IntrinsicData{"workgroupBarrier", IntrinsicType::kWorkgroupBarrier}));
+
 using ResolverIntrinsicTest_DataPacking = ResolverTestWithParam<IntrinsicData>;
 TEST_P(ResolverIntrinsicTest_DataPacking, InferType) {
   auto param = GetParam();

src/semantic/intrinsic.h

@@ -88,6 +88,7 @@ enum class IntrinsicType {
   kSmoothStep,
   kSqrt,
   kStep,
+  kStorageBarrier,
   kTan,
   kTanh,
   kTextureDimensions,
@@ -107,6 +108,7 @@ enum class IntrinsicType {
   kUnpack2x16Unorm,
   kUnpack4x8Snorm,
   kUnpack4x8Unorm,
+  kWorkgroupBarrier,
 };
 
 /// Matches the IntrisicType by name
@@ -159,6 +161,11 @@ bool IsDataPackingIntrinsic(IntrinsicType i);
 /// @returns true if the given `i` is a data unpacking intrinsic
 bool IsDataUnpackingIntrinsic(IntrinsicType i);
 
+/// Determines if the given `i` is a barrier intrinsic
+/// @param i the intrinsic
+/// @returns true if the given `i` is a barrier intrinsic
+bool IsBarrierIntrinsic(IntrinsicType i);
+
 /// Intrinsic holds the semantic information for an intrinsic function.
 class Intrinsic : public Castable<Intrinsic, CallTarget> {
  public:
@@ -204,6 +211,9 @@ class Intrinsic : public Castable<Intrinsic, CallTarget> {
   /// @returns true if intrinsic is a data unpacking intrinsic
   bool IsDataUnpacking() const;
 
+  /// @returns true if intrinsic is a barrier intrinsic
+  bool IsBarrier() const;
+
  private:
   IntrinsicType const type_;
 };

src/semantic/sem_intrinsic.cc

@@ -94,6 +94,7 @@ const char* Intrinsic::str() const {
   INTRINSIC(IntrinsicType::kSmoothStep, "smoothStep")                     \
   INTRINSIC(IntrinsicType::kSqrt, "sqrt")                                 \
   INTRINSIC(IntrinsicType::kStep, "step")                                 \
+  INTRINSIC(IntrinsicType::kStorageBarrier, "storageBarrier")             \
   INTRINSIC(IntrinsicType::kTan, "tan")                                   \
   INTRINSIC(IntrinsicType::kTanh, "tanh")                                 \
   INTRINSIC(IntrinsicType::kTextureDimensions, "textureDimensions")       \
@@ -112,7 +113,8 @@ const char* Intrinsic::str() const {
   INTRINSIC(IntrinsicType::kUnpack2x16Snorm, "unpack2x16snorm")           \
   INTRINSIC(IntrinsicType::kUnpack2x16Unorm, "unpack2x16unorm")           \
   INTRINSIC(IntrinsicType::kUnpack4x8Snorm, "unpack4x8snorm")             \
-  INTRINSIC(IntrinsicType::kUnpack4x8Unorm, "unpack4x8unorm")
+  INTRINSIC(IntrinsicType::kUnpack4x8Unorm, "unpack4x8unorm")             \
+  INTRINSIC(IntrinsicType::kWorkgroupBarrier, "workgroupBarrier")
 
 IntrinsicType ParseIntrinsicType(const std::string& name) {
 #define INTRINSIC(ENUM, NAME) \
@@ -187,6 +189,11 @@ bool IsDataUnpackingIntrinsic(IntrinsicType i) {
          i == IntrinsicType::kUnpack2x16Float;
 }
 
+bool IsBarrierIntrinsic(IntrinsicType i) {
+  return i == IntrinsicType::kWorkgroupBarrier ||
+         i == IntrinsicType::kStorageBarrier;
+}
+
 Intrinsic::Intrinsic(IntrinsicType type,
                      type::Type* return_type,
                      const ParameterList& parameters)
@@ -226,5 +233,9 @@ bool Intrinsic::IsDataUnpacking() const {
   return IsDataUnpackingIntrinsic(type_);
 }
 
+bool Intrinsic::IsBarrier() const {
+  return IsBarrierIntrinsic(type_);
+}
+
 }  // namespace semantic
 }  // namespace tint

src/semantic/sem_intrinsic_test.cc

@@ -99,6 +99,7 @@ INSTANTIATE_TEST_SUITE_P(
         IntrinsicData{"smoothStep", IntrinsicType::kSmoothStep},
         IntrinsicData{"sqrt", IntrinsicType::kSqrt},
         IntrinsicData{"step", IntrinsicType::kStep},
+        IntrinsicData{"storageBarrier", IntrinsicType::kStorageBarrier},
         IntrinsicData{"tan", IntrinsicType::kTan},
         IntrinsicData{"tanh", IntrinsicType::kTanh},
         IntrinsicData{"textureDimensions", IntrinsicType::kTextureDimensions},
@@ -117,7 +118,8 @@ INSTANTIATE_TEST_SUITE_P(
         IntrinsicData{"unpack2x16snorm", IntrinsicType::kUnpack2x16Snorm},
         IntrinsicData{"unpack2x16unorm", IntrinsicType::kUnpack2x16Unorm},
         IntrinsicData{"unpack4x8snorm", IntrinsicType::kUnpack4x8Snorm},
-        IntrinsicData{"unpack4x8unorm", IntrinsicType::kUnpack4x8Unorm}));
+        IntrinsicData{"unpack4x8unorm", IntrinsicType::kUnpack4x8Unorm},
+        IntrinsicData{"workgroupBarrier", IntrinsicType::kWorkgroupBarrier}));
 
 TEST_F(IntrinsicTypeTest, ParseNoMatch) {
   EXPECT_EQ(ParseIntrinsicType("not_intrinsic"), IntrinsicType::kNone);

src/writer/hlsl/generator_impl.cc

@@ -505,6 +505,8 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
       return EmitDataPackingCall(pre, out, expr, intrinsic);
     } else if (intrinsic->IsDataUnpacking()) {
       return EmitDataUnpackingCall(pre, out, expr, intrinsic);
+    } else if (intrinsic->IsBarrier()) {
+      return EmitBarrierCall(pre, out, intrinsic);
     }
     auto name = generate_builtin_name(intrinsic);
     if (name.empty()) {
@@ -716,6 +718,23 @@ bool GeneratorImpl::EmitDataUnpackingCall(
   return true;
 }
 
+bool GeneratorImpl::EmitBarrierCall(std::ostream&,
+                                    std::ostream& out,
+                                    const semantic::Intrinsic* intrinsic) {
+  // TODO(crbug.com/tint/661): Combine sequential barriers to a single
+  // instruction.
+  if (intrinsic->Type() == semantic::IntrinsicType::kWorkgroupBarrier) {
+    out << "GroupMemoryBarrierWithGroupSync()";
+  } else if (intrinsic->Type() == semantic::IntrinsicType::kStorageBarrier) {
+    out << "DeviceMemoryBarrierWithGroupSync()";
+  } else {
+    TINT_UNREACHABLE(diagnostics_) << "unexpected barrier intrinsic type "
+                                   << semantic::str(intrinsic->Type());
+    return false;
+  }
+  return true;
+}
+
 bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
                                     std::ostream& out,
                                     ast::CallExpression* expr,

src/writer/hlsl/generator_impl.h

@@ -120,6 +120,14 @@ class GeneratorImpl : public TextGenerator {
   bool EmitCall(std::ostream& pre,
                 std::ostream& out,
                 ast::CallExpression* expr);
+  /// Handles generating a barrier intrinsic call
+  /// @param pre the preamble for the expression stream
+  /// @param out the output of the expression stream
+  /// @param intrinsic the semantic information for the barrier intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitBarrierCall(std::ostream& pre,
+                       std::ostream& out,
+                       const semantic::Intrinsic* intrinsic);
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param pre the preamble for the expression stream

src/writer/hlsl/generator_impl_intrinsic_test.cc

@@ -13,6 +13,8 @@
 // limitations under the License.
 
 #include "gmock/gmock.h"
+#include "src/ast/call_statement.h"
+#include "src/ast/stage_decoration.h"
 #include "src/semantic/call.h"
 #include "src/writer/hlsl/test_helper.h"
 
@@ -413,6 +415,44 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Float) {
       HasSubstr("f16tof32(uint2(_tint_tmp & 0xffff, _tint_tmp >> 16))"));
 }
 
+TEST_F(HlslGeneratorImplTest_Intrinsic, StorageBarrier) {
+  Func("main", {}, ty.void_(),
+       {create<ast::CallStatement>(Call("storageBarrier"))},
+       {create<ast::StageDecoration>(ast::PipelineStage::kCompute)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate(out)) << gen.error();
+  EXPECT_EQ(result(), R"([numthreads(1, 1, 1)]
+void main() {
+  DeviceMemoryBarrierWithGroupSync();
+  return;
+}
+
+)");
+
+  Validate();
+}
+
+TEST_F(HlslGeneratorImplTest_Intrinsic, WorkgroupBarrier) {
+  Func("main", {}, ty.void_(),
+       {create<ast::CallStatement>(Call("workgroupBarrier"))},
+       {create<ast::StageDecoration>(ast::PipelineStage::kCompute)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate(out)) << gen.error();
+  EXPECT_EQ(result(), R"([numthreads(1, 1, 1)]
+void main() {
+  GroupMemoryBarrierWithGroupSync();
+  return;
+}
+
+)");
+
+  Validate();
+}
+
 }  // namespace
 }  // namespace hlsl
 }  // namespace writer

src/writer/msl/generator_impl.cc

@@ -339,6 +339,18 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
       out_ << "))";
       return true;
     }
+    // TODO(crbug.com/tint/661): Combine sequential barriers to a single
+    // instruction.
+    if (intrinsic->Type() == semantic::IntrinsicType::kStorageBarrier) {
+      make_indent();
+      out_ << "threadgroup_barrier(mem_flags::mem_device)";
+      return true;
+    }
+    if (intrinsic->Type() == semantic::IntrinsicType::kWorkgroupBarrier) {
+      make_indent();
+      out_ << "threadgroup_barrier(mem_flags::mem_threadgroup)";
+      return true;
+    }
     auto name = generate_builtin_name(intrinsic);
     if (name.empty()) {
       return false;

src/writer/msl/generator_impl_intrinsic_test.cc

@@ -105,6 +105,8 @@ ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
     case IntrinsicType::kReflect:
     case IntrinsicType::kStep:
       return builder->Call(str.str(), "f2", "f2");
+    case IntrinsicType::kStorageBarrier:
+      return builder->Call(str.str());
     case IntrinsicType::kCross:
       return builder->Call(str.str(), "f3", "f3");
     case IntrinsicType::kFma:
@@ -152,6 +154,8 @@ ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
     case IntrinsicType::kUnpack2x16Snorm:
     case IntrinsicType::kUnpack2x16Unorm:
       return builder->Call(str.str(), "u1");
+    case IntrinsicType::kWorkgroupBarrier:
+      return builder->Call(str.str());
     default:
       break;
   }
@@ -286,6 +290,28 @@ TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
   EXPECT_EQ(gen.result(), "  dot(param1, param2)");
 }
 
+TEST_F(MslGeneratorImplTest, StorageBarrier) {
+  auto* call = Call("storageBarrier");
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  ASSERT_TRUE(gen.EmitExpression(call)) << gen.error();
+  EXPECT_EQ(gen.result(), "  threadgroup_barrier(mem_flags::mem_device)");
+}
+
+TEST_F(MslGeneratorImplTest, WorkgroupBarrier) {
+  auto* call = Call("workgroupBarrier");
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  ASSERT_TRUE(gen.EmitExpression(call)) << gen.error();
+  EXPECT_EQ(gen.result(), "  threadgroup_barrier(mem_flags::mem_threadgroup)");
+}
+
 TEST_F(MslGeneratorImplTest, Pack2x16Float) {
   auto* call = Call("pack2x16float", "p1");
   Global("p1", ty.vec2<f32>(), ast::StorageClass::kFunction);

src/writer/spirv/builder.cc

@@ -1921,6 +1921,13 @@ uint32_t Builder::GenerateIntrinsic(ast::CallExpression* call,
     return result_id;
   }
 
+  if (intrinsic->IsBarrier()) {
+    if (!GenerateControlBarrierIntrinsic(intrinsic)) {
+      return 0;
+    }
+    return result_id;
+  }
+
   OperandList params = {Operand::Int(result_type_id), result};
 
   spv::Op op = spv::Op::OpNop;
@@ -2455,6 +2462,47 @@ bool Builder::GenerateTextureIntrinsic(ast::CallExpression* call,
   return post_emission();
 }
 
+bool Builder::GenerateControlBarrierIntrinsic(
+    const semantic::Intrinsic* intrinsic) {
+  auto const op = spv::Op::OpControlBarrier;
+  uint32_t execution = 0;
+  uint32_t memory = 0;
+  uint32_t semantics = 0;
+
+  // TODO(crbug.com/tint/661): Combine sequential barriers to a single
+  // instruction.
+  if (intrinsic->Type() == semantic::IntrinsicType::kWorkgroupBarrier) {
+    execution = static_cast<uint32_t>(spv::Scope::Workgroup);
+    memory = static_cast<uint32_t>(spv::Scope::Workgroup);
+    semantics =
+        static_cast<uint32_t>(spv::MemorySemanticsMask::AcquireRelease) |
+        static_cast<uint32_t>(spv::MemorySemanticsMask::WorkgroupMemory);
+  } else if (intrinsic->Type() == semantic::IntrinsicType::kStorageBarrier) {
+    execution = static_cast<uint32_t>(spv::Scope::Workgroup);
+    memory = static_cast<uint32_t>(spv::Scope::Device);
+    semantics =
+        static_cast<uint32_t>(spv::MemorySemanticsMask::AcquireRelease) |
+        static_cast<uint32_t>(spv::MemorySemanticsMask::UniformMemory);
+  } else {
+    error_ = "unexpected barrier intrinsic type ";
+    error_ += semantic::str(intrinsic->Type());
+    return false;
+  }
+
+  auto execution_id = GenerateConstantIfNeeded(ScalarConstant::U32(execution));
+  auto memory_id = GenerateConstantIfNeeded(ScalarConstant::U32(memory));
+  auto semantics_id = GenerateConstantIfNeeded(ScalarConstant::U32(semantics));
+  if (execution_id == 0 || memory_id == 0 || semantics_id == 0) {
+    return false;
+  }
+
+  return push_function_inst(op, {
+                                    Operand::Int(execution_id),
+                                    Operand::Int(memory_id),
+                                    Operand::Int(semantics_id),
+                                });
+}
+
 uint32_t Builder::GenerateSampledImage(type::Type* texture_type,
                                        Operand texture_operand,
                                        Operand sampler_operand) {

src/writer/spirv/builder.h

@@ -358,6 +358,11 @@ class Builder {
                                 const semantic::Intrinsic* intrinsic,
                                 spirv::Operand result_type,
                                 spirv::Operand result_id);
+  /// Generates a control barrier statement.
+  /// @param intrinsic the semantic information for the barrier intrinsic
+  /// parameters
+  /// @returns true on success
+  bool GenerateControlBarrierIntrinsic(const semantic::Intrinsic* intrinsic);
   /// Generates a sampled image
   /// @param texture_type the texture type
   /// @param texture_operand the texture operand

src/writer/spirv/builder_intrinsic_test.cc

@@ -1596,6 +1596,71 @@ INSTANTIATE_TEST_SUITE_P(
                     IntrinsicData{"unpack2x16unorm", "UnpackUnorm2x16"},
                     IntrinsicData{"unpack2x16float", "UnpackHalf2x16"}));
 
+TEST_F(IntrinsicBuilderTest, Call_WorkgroupBarrier) {
+  Func("f", ast::VariableList{}, ty.void_(),
+       ast::StatementList{
+           create<ast::CallStatement>(Call("workgroupBarrier")),
+       },
+       ast::DecorationList{
+           create<ast::StageDecoration>(ast::PipelineStage::kCompute),
+       });
+
+  spirv::Builder& b = Build();
+
+  ASSERT_TRUE(b.Build()) << b.error();
+
+  ASSERT_EQ(b.functions().size(), 1u);
+
+  auto* expected_types = R"(%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%6 = OpTypeInt 32 0
+%7 = OpConstant %6 2
+%8 = OpConstant %6 264
+)";
+  auto got_types = DumpInstructions(b.types());
+  EXPECT_EQ(expected_types, got_types);
+
+  auto* expected_instructions = R"(OpControlBarrier %7 %7 %8
+)";
+  auto got_instructions = DumpInstructions(b.functions()[0].instructions());
+  EXPECT_EQ(expected_instructions, got_instructions);
+
+  Validate(b);
+}
+
+TEST_F(IntrinsicBuilderTest, Call_StorageBarrier) {
+  Func("f", ast::VariableList{}, ty.void_(),
+       ast::StatementList{
+           create<ast::CallStatement>(Call("storageBarrier")),
+       },
+       ast::DecorationList{
+           create<ast::StageDecoration>(ast::PipelineStage::kCompute),
+       });
+
+  spirv::Builder& b = Build();
+
+  ASSERT_TRUE(b.Build()) << b.error();
+
+  ASSERT_EQ(b.functions().size(), 1u);
+
+  auto* expected_types = R"(%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%6 = OpTypeInt 32 0
+%7 = OpConstant %6 2
+%8 = OpConstant %6 1
+%9 = OpConstant %6 72
+)";
+  auto got_types = DumpInstructions(b.types());
+  EXPECT_EQ(expected_types, got_types);
+
+  auto* expected_instructions = R"(OpControlBarrier %7 %8 %9
+)";
+  auto got_instructions = DumpInstructions(b.functions()[0].instructions());
+  EXPECT_EQ(expected_instructions, got_instructions);
+
+  Validate(b);
+}
+
 }  // namespace
 }  // namespace spirv
 }  // namespace writer

test/BUILD.gn

@@ -272,6 +272,7 @@ source_set("tint_unittests_spv_reader_src") {
     "../src/reader/spirv/function_misc_test.cc",
     "../src/reader/spirv/function_var_test.cc",
     "../src/reader/spirv/namer_test.cc",
+    "../src/reader/spirv/parser_impl_barrier_test.cc",
     "../src/reader/spirv/parser_impl_convert_member_decoration_test.cc",
     "../src/reader/spirv/parser_impl_convert_type_test.cc",
     "../src/reader/spirv/parser_impl_function_decl_test.cc",