diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index b223e03115..21353b4ff7 100644 --- a/src/CMakeLists.txt +++ b/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 diff --git a/src/intrinsic_table.cc b/src/intrinsic_table.cc index 30f18d7535..6728365176 100644 --- a/src/intrinsic_table.cc +++ b/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); diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc index 63a124e36b..0b03991cb9 100644 --- a/src/reader/spirv/function.cc +++ b/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(op)->literal(); + if (auto* int_lit = lit->As()) { + 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(builder_.Call("workgroupBarrier"))); + semantics &= ~SpvMemorySemanticsWorkgroupMemoryMask; + } + if (semantics & SpvMemorySemanticsUniformMemoryMask) { + if (memory != SpvScopeDevice) { + return Fail() << "storageBarrier requires device memory scope"; + } + AddStatement(create(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()); diff --git a/src/reader/spirv/function.h b/src/reader/spirv/function.h index 4e20c1e39c..052f056876 100644 --- a/src/reader/spirv/function.h +++ b/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 diff --git a/src/reader/spirv/parser_impl_barrier_test.cc b/src/reader/spirv/parser_impl_barrier_test.cc new file mode 100644 index 0000000000..751ce3f9a7 --- /dev/null +++ b/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(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(); + 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(); + 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(); + 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(); + 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 diff --git a/src/resolver/intrinsic_test.cc b/src/resolver/intrinsic_test.cc index b266b00734..22bb0eef80 100644 --- a/src/resolver/intrinsic_test.cc +++ b/src/resolver/intrinsic_test.cc @@ -538,6 +538,37 @@ inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) { return out; } +using ResolverIntrinsicTest_Barrier = ResolverTestWithParam; +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()); +} + +TEST_P(ResolverIntrinsicTest_Barrier, Error_TooManyParams) { + auto param = GetParam(); + + auto* call = Call(param.name, vec4(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; TEST_P(ResolverIntrinsicTest_DataPacking, InferType) { auto param = GetParam(); diff --git a/src/semantic/intrinsic.h b/src/semantic/intrinsic.h index 9b58135009..b1bd113d87 100644 --- a/src/semantic/intrinsic.h +++ b/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 { public: @@ -204,6 +211,9 @@ class Intrinsic : public Castable { /// @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_; }; diff --git a/src/semantic/sem_intrinsic.cc b/src/semantic/sem_intrinsic.cc index 4d9fc63c7a..0ff77a2102 100644 --- a/src/semantic/sem_intrinsic.cc +++ b/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 diff --git a/src/semantic/sem_intrinsic_test.cc b/src/semantic/sem_intrinsic_test.cc index 698fcd3210..227b29bfae 100644 --- a/src/semantic/sem_intrinsic_test.cc +++ b/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); diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 186ed20a8c..1b940e8580 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/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, diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index c967e553cb..48ef7e3fc9 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/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 diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc index e7945b1c8d..372e145691 100644 --- a/src/writer/hlsl/generator_impl_intrinsic_test.cc +++ b/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(Call("storageBarrier"))}, + {create(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(Call("workgroupBarrier"))}, + {create(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 diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 0a4e76881e..e971d50425 100644 --- a/src/writer/msl/generator_impl.cc +++ b/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; diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc index b918c1930f..154b0ac174 100644 --- a/src/writer/msl/generator_impl_intrinsic_test.cc +++ b/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(), ast::StorageClass::kFunction); diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index 7209aba4db..77401a8a09 100644 --- a/src/writer/spirv/builder.cc +++ b/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(spv::Scope::Workgroup); + memory = static_cast(spv::Scope::Workgroup); + semantics = + static_cast(spv::MemorySemanticsMask::AcquireRelease) | + static_cast(spv::MemorySemanticsMask::WorkgroupMemory); + } else if (intrinsic->Type() == semantic::IntrinsicType::kStorageBarrier) { + execution = static_cast(spv::Scope::Workgroup); + memory = static_cast(spv::Scope::Device); + semantics = + static_cast(spv::MemorySemanticsMask::AcquireRelease) | + static_cast(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) { diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h index 6147c9ff3f..c99749fccc 100644 --- a/src/writer/spirv/builder.h +++ b/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 diff --git a/src/writer/spirv/builder_intrinsic_test.cc b/src/writer/spirv/builder_intrinsic_test.cc index 4e2bb84f7c..ba07575d14 100644 --- a/src/writer/spirv/builder_intrinsic_test.cc +++ b/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(Call("workgroupBarrier")), + }, + ast::DecorationList{ + create(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(Call("storageBarrier")), + }, + ast::DecorationList{ + create(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 diff --git a/test/BUILD.gn b/test/BUILD.gn index 07ff399360..9d1d9d7de3 100644 --- a/test/BUILD.gn +++ b/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",