[spirv-reader][ir] Emit OpControlBarrier
Add support for emitting the correct barrier types based on a control
barrier.
Bug: 401001910
Change-Id: Ic076f2374b6ea0cf0df6f4e3d774ca8258582498
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/232174
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 699fe19..bb3d2a2 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -81,6 +81,7 @@
alwayslink = True,
srcs = [
"atomics_test.cc",
+ "barrier_test.cc",
"binary_test.cc",
"bit_test.cc",
"branch_test.cc",
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index da6f3f8..02e0870 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -90,6 +90,7 @@
################################################################################
tint_add_target(tint_lang_spirv_reader_parser_test test
lang/spirv/reader/parser/atomics_test.cc
+ lang/spirv/reader/parser/barrier_test.cc
lang/spirv/reader/parser/binary_test.cc
lang/spirv/reader/parser/bit_test.cc
lang/spirv/reader/parser/branch_test.cc
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index 3e05963..a806fd4 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -89,6 +89,7 @@
tint_unittests_source_set("unittests") {
sources = [
"atomics_test.cc",
+ "barrier_test.cc",
"binary_test.cc",
"bit_test.cc",
"branch_test.cc",
diff --git a/src/tint/lang/spirv/reader/parser/barrier_test.cc b/src/tint/lang/spirv/reader/parser/barrier_test.cc
new file mode 100644
index 0000000..b6615e9
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/barrier_test.cc
@@ -0,0 +1,279 @@
+// Copyright 2025 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/spirv/reader/parser/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+TEST_F(SpirvParserTest, ControlBarrier_WorkgroupBarrier) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %helper "helper"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %helper = OpFunction %void None %1
+ %4 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %5 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%helper = func():void {
+ $B1: {
+ %2:void = workgroupBarrier
+ ret
+ }
+}
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ControlBarrier_StorageBarrier) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %helper "helper"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_1 = OpConstant %uint 1
+ %uint_72 = OpConstant %uint 72
+ %helper = OpFunction %void None %1
+ %4 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_72
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %5 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%helper = func():void {
+ $B1: {
+ %2:void = storageBarrier
+ ret
+ }
+}
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ControlBarrier_TextureBarrier) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %helper "helper"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_1 = OpConstant %uint 1
+ %uint_2056 = OpConstant %uint 2056
+ %helper = OpFunction %void None %1
+ %4 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_2056
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %5 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%helper = func():void {
+ $B1: {
+ %2:void = textureBarrier
+ ret
+ }
+}
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ControlBarrier_WorkgroupAndTextureAndStorageBarrier) {
+ // Check that we emit multiple adjacent barrier calls when the flags
+ // are combined.
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %helper "helper"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_x948 = OpConstant %uint 0x948
+ %helper = OpFunction %void None %1
+ %4 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_x948
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %5 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%helper = func():void {
+ $B1: {
+ %2:void = workgroupBarrier
+ %3:void = storageBarrier
+ %4:void = textureBarrier
+ ret
+ }
+}
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserDeathTest, ControlBarrier_ErrBarrierInvalidExecution) {
+ auto* src = R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_3 = OpConstant %uint 3
+ %uint_2 = OpConstant %uint 2
+ %uint_264 = OpConstant %uint 264
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ OpControlBarrier %uint_3 %uint_2 %uint_264
+ OpReturn
+ OpFunctionEnd
+ )";
+ EXPECT_DEATH_IF_SUPPORTED({ auto _ = Run(src); }, "internal compiler error");
+}
+
+TEST_F(SpirvParserDeathTest, ControlBarrier_ErrBarrierSemanticsMissingAcquireRelease) {
+ auto* src = R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %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_DEATH_IF_SUPPORTED({ auto _ = Run(src); }, "internal compiler error");
+}
+
+TEST_F(SpirvParserDeathTest, ControlBarrier_ErrStorageBarrierInvalidMemory) {
+ auto* src = R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_72 = OpConstant %uint 72
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ OpControlBarrier %uint_2 %uint_1 %uint_72
+ OpReturn
+ OpFunctionEnd
+ )";
+ EXPECT_DEATH_IF_SUPPORTED({ auto _ = Run(src); }, "internal compiler error");
+}
+
+TEST_F(SpirvParserDeathTest, ControlBarrier_ErrTextureBarrierInvalidMemory) {
+ auto* src = R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_2056 = OpConstant %uint 2056
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ OpControlBarrier %uint_2 %uint_1 %uint_2056
+ OpReturn
+ OpFunctionEnd
+ )";
+ EXPECT_DEATH_IF_SUPPORTED({ auto _ = Run(src); }, "internal compiler error");
+}
+
+} // namespace
+} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index cbd9a5b..0fd6b74 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -1190,6 +1190,9 @@
case spv::Op::OpAtomicIDecrement:
EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kAtomicIDecrement);
break;
+ case spv::Op::OpControlBarrier:
+ EmitControlBarrier(inst);
+ break;
default:
TINT_UNIMPLEMENTED()
<< "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());
@@ -1215,6 +1218,53 @@
}
}
+ void EmitControlBarrier(const spvtools::opt::Instruction& inst) {
+ auto get_constant = [&](uint32_t idx) {
+ uint32_t id = inst.GetSingleWordOperand(idx);
+ if (auto* constant = spirv_context_->get_constant_mgr()->FindDeclaredConstant(id)) {
+ return constant->GetU32();
+ }
+ TINT_ICE() << "invalid or missing operands for control barrier";
+ };
+
+ uint32_t execution = get_constant(0);
+ uint32_t memory = get_constant(1);
+ uint32_t semantics = get_constant(2);
+
+ if (execution != uint32_t(spv::Scope::Workgroup)) {
+ TINT_ICE() << "unsupported control barrier execution scope: "
+ << "expected Workgroup (2), got: " << execution;
+ }
+
+ if (semantics & uint32_t(spv::MemorySemanticsMask::AcquireRelease)) {
+ semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::AcquireRelease);
+ } else {
+ TINT_ICE() << "control barrier semantics requires acquire and release";
+ }
+ if (memory != uint32_t(spv::Scope::Workgroup)) {
+ TINT_ICE() << "control barrier requires workgroup memory scope";
+ }
+
+ if (semantics & uint32_t(spv::MemorySemanticsMask::WorkgroupMemory)) {
+ EmitWithoutSpvResult(b_.Call(ty_.void_(), core::BuiltinFn::kWorkgroupBarrier));
+ semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::WorkgroupMemory);
+ }
+
+ if (semantics & uint32_t(spv::MemorySemanticsMask::UniformMemory)) {
+ EmitWithoutSpvResult(b_.Call(ty_.void_(), core::BuiltinFn::kStorageBarrier));
+ semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::UniformMemory);
+ }
+
+ if (semantics & uint32_t(spv::MemorySemanticsMask::ImageMemory)) {
+ EmitWithoutSpvResult(b_.Call(ty_.void_(), core::BuiltinFn::kTextureBarrier));
+ semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::ImageMemory);
+ }
+
+ if (semantics) {
+ TINT_ICE() << "unsupported control barrier semantics: " << semantics;
+ }
+ }
+
void CheckAtomicNotFloat(const spvtools::opt::Instruction& inst) {
auto* ty = Type(inst.type_id());
if (ty->IsFloatScalar()) {
