[dawn] IR Reflection for single entry and overrides
This adds 'reflection.h' as scaffolding for future API.
Bug: 380043635
Change-Id: I5021a1c30c2019630347b1a6982cd0043dd31d63
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/232335
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
diff --git a/include/tint/tint.h b/include/tint/tint.h
index ebfd5a5..be2bb51 100644
--- a/include/tint/tint.h
+++ b/include/tint/tint.h
@@ -36,6 +36,7 @@
#include "src/tint/api/common/vertex_pulling_config.h"
#include "src/tint/api/tint.h"
+#include "src/tint/lang/core/ir/reflection.h"
#include "src/tint/lang/core/ir/transform/single_entry_point.h"
#include "src/tint/lang/core/ir/transform/substitute_overrides.h"
#include "src/tint/lang/core/type/manager.h"
diff --git a/src/dawn/native/null/DeviceNull.cpp b/src/dawn/native/null/DeviceNull.cpp
index 4b0cd77..9399a89 100644
--- a/src/dawn/native/null/DeviceNull.cpp
+++ b/src/dawn/native/null/DeviceNull.cpp
@@ -479,35 +479,52 @@
MaybeError ComputePipeline::InitializeImpl() {
const ProgrammableStage& computeStage = GetStage(SingleShaderStage::Compute);
- tint::Program transformedProgram;
- tint::ast::transform::Manager transformManager;
- tint::ast::transform::DataMap transformInputs;
-
+ std::optional<tint::ast::transform::SubstituteOverride::Config> substituteOverrideConfig;
if (!computeStage.metadata->overrides.empty()) {
- transformManager.Add<tint::ast::transform::SingleEntryPoint>();
- transformInputs.Add<tint::ast::transform::SingleEntryPoint::Config>(
- computeStage.entryPoint.c_str());
-
- // This needs to run after SingleEntryPoint transform which removes unused overrides for
- // current entry point.
- transformManager.Add<tint::ast::transform::SubstituteOverride>();
- transformInputs.Add<tint::ast::transform::SubstituteOverride::Config>(
- BuildSubstituteOverridesTransformConfig(computeStage));
+ substituteOverrideConfig = BuildSubstituteOverridesTransformConfig(computeStage);
}
- auto tintProgram = computeStage.module->GetTintProgram();
- DAWN_TRY_ASSIGN(transformedProgram, RunTransforms(&transformManager, &(tintProgram->program),
- transformInputs, nullptr, nullptr));
+ // Convert the AST program to an IR module.
+ auto ir =
+ tint::wgsl::reader::ProgramToLoweredIR(computeStage.module->GetTintProgram()->program);
+ DAWN_INVALID_IF(ir != tint::Success, "An error occurred while generating Tint IR\n%s",
+ ir.Failure().reason);
- // Do the workgroup size validation.
+ auto singleEntryPointResult =
+ tint::core::ir::transform::SingleEntryPoint(ir.Get(), computeStage.entryPoint.c_str());
+ DAWN_INVALID_IF(singleEntryPointResult != tint::Success,
+ "Pipeline single entry point (IR) failed:\n%s",
+ singleEntryPointResult.Failure().reason);
+
+ if (substituteOverrideConfig) {
+ // this needs to run after SingleEntryPoint transform which removes unused
+ // overrides for the current entry point.
+ tint::core::ir::transform::SubstituteOverridesConfig cfg;
+ cfg.map = substituteOverrideConfig->map;
+ auto substituteOverridesResult =
+ tint::core::ir::transform::SubstituteOverrides(ir.Get(), cfg);
+ DAWN_INVALID_IF(substituteOverridesResult != tint::Success,
+ "Pipeline override substitution (IR) failed:\n%s",
+ substituteOverridesResult.Failure().reason);
+ }
+
+ auto limits = LimitsForCompilationRequest::Create(GetDevice()->GetLimits().v1);
+ auto adapterSupportedLimits =
+ LimitsForCompilationRequest::Create(GetDevice()->GetAdapter()->GetLimits().v1);
+ auto maxSubgroupSize = GetDevice()->GetAdapter()->GetPhysicalDevice()->GetSubgroupMaxSize();
+
+ // Workgroup validation has to come after overrides to have been substituted.
+ auto wgInfo = tint::core::ir::GetWorkgroupInfo(ir.Get());
+
+ DAWN_INVALID_IF(wgInfo != tint::Success, "Getting workgroup info has failed (IR):\n%s",
+ wgInfo.Failure().reason);
+
Extent3D _;
- DAWN_TRY_ASSIGN(
- _, ValidateComputeStageWorkgroupSize(
- transformedProgram, computeStage.entryPoint.c_str(),
- computeStage.metadata->usesSubgroupMatrix,
- GetDevice()->GetAdapter()->GetPhysicalDevice()->GetSubgroupMaxSize(),
- LimitsForCompilationRequest::Create(GetDevice()->GetLimits().v1),
- LimitsForCompilationRequest::Create(GetDevice()->GetAdapter()->GetLimits().v1)));
+ DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
+ wgInfo.Get().x, wgInfo.Get().y, wgInfo.Get().z,
+ wgInfo.Get().storage_size, computeStage.metadata->usesSubgroupMatrix,
+ maxSubgroupSize, limits, adapterSupportedLimits));
+
return {};
}
diff --git a/src/tint/lang/core/ir/BUILD.bazel b/src/tint/lang/core/ir/BUILD.bazel
index d45214d..c9e5f0c 100644
--- a/src/tint/lang/core/ir/BUILD.bazel
+++ b/src/tint/lang/core/ir/BUILD.bazel
@@ -82,6 +82,7 @@
"operand_instruction.cc",
"override.cc",
"phony.cc",
+ "reflection.cc",
"return.cc",
"store.cc",
"store_vector_element.cc",
@@ -145,6 +146,7 @@
"referenced_functions.h",
"referenced_module_decls.h",
"referenced_module_vars.h",
+ "reflection.h",
"return.h",
"store.h",
"store_vector_element.h",
@@ -224,6 +226,7 @@
"referenced_functions_test.cc",
"referenced_module_decls_test.cc",
"referenced_module_vars_test.cc",
+ "reflection_test.cc",
"return_test.cc",
"store_test.cc",
"store_vector_element_test.cc",
diff --git a/src/tint/lang/core/ir/BUILD.cmake b/src/tint/lang/core/ir/BUILD.cmake
index ce7d11a..6c92812 100644
--- a/src/tint/lang/core/ir/BUILD.cmake
+++ b/src/tint/lang/core/ir/BUILD.cmake
@@ -134,6 +134,8 @@
lang/core/ir/referenced_functions.h
lang/core/ir/referenced_module_decls.h
lang/core/ir/referenced_module_vars.h
+ lang/core/ir/reflection.cc
+ lang/core/ir/reflection.h
lang/core/ir/return.cc
lang/core/ir/return.h
lang/core/ir/store.cc
@@ -230,6 +232,7 @@
lang/core/ir/referenced_functions_test.cc
lang/core/ir/referenced_module_decls_test.cc
lang/core/ir/referenced_module_vars_test.cc
+ lang/core/ir/reflection_test.cc
lang/core/ir/return_test.cc
lang/core/ir/store_test.cc
lang/core/ir/store_vector_element_test.cc
diff --git a/src/tint/lang/core/ir/BUILD.gn b/src/tint/lang/core/ir/BUILD.gn
index 019dfc3..3cdec1d 100644
--- a/src/tint/lang/core/ir/BUILD.gn
+++ b/src/tint/lang/core/ir/BUILD.gn
@@ -135,6 +135,8 @@
"referenced_functions.h",
"referenced_module_decls.h",
"referenced_module_vars.h",
+ "reflection.cc",
+ "reflection.h",
"return.cc",
"return.h",
"store.cc",
@@ -225,6 +227,7 @@
"referenced_functions_test.cc",
"referenced_module_decls_test.cc",
"referenced_module_vars_test.cc",
+ "reflection_test.cc",
"return_test.cc",
"store_test.cc",
"store_vector_element_test.cc",
diff --git a/src/tint/lang/core/ir/reflection.cc b/src/tint/lang/core/ir/reflection.cc
new file mode 100644
index 0000000..824b0b6
--- /dev/null
+++ b/src/tint/lang/core/ir/reflection.cc
@@ -0,0 +1,71 @@
+// 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/core/ir/reflection.h"
+
+#include <string>
+#include <utility>
+
+#include "src/tint/lang/core/ir/var.h"
+#include "src/tint/lang/core/type/pointer.h"
+
+namespace tint::core::ir {
+
+Result<WorkgroupInfo> GetWorkgroupInfo(core::ir::Module& ir) {
+ std::optional<std::array<uint32_t, 3>> const_wg_size;
+ for (auto func : ir.functions) {
+ if (!func->IsEntryPoint()) {
+ continue;
+ }
+ const_wg_size = func->WorkgroupSizeAsConst();
+ }
+
+ if (!const_wg_size) {
+ return Failure{"IR GetWorkgroupInfo: Could not find workgroup size"};
+ }
+
+ size_t wg_storage_size = 0u;
+ for (auto* inst : *ir.root_block) {
+ if (auto* as_var = inst->As<core::ir::Var>()) {
+ auto* ptr = as_var->Result(0)->Type()->As<core::type::Pointer>();
+ if (ptr->AddressSpace() != core::AddressSpace::kWorkgroup) {
+ continue;
+ }
+ auto* ty = ptr->StoreType();
+ uint32_t align = ty->Align();
+ uint32_t size = ty->Size();
+
+ // This essentially matches std430 layout rules from GLSL, which are in
+ // turn specified as an upper bound for Vulkan layout sizing.
+ wg_storage_size += tint::RoundUp(16u, tint::RoundUp(align, size));
+ }
+ }
+ return WorkgroupInfo{(*const_wg_size)[0], (*const_wg_size)[1], (*const_wg_size)[2],
+ wg_storage_size};
+}
+
+} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/reflection.h b/src/tint/lang/core/ir/reflection.h
new file mode 100644
index 0000000..a44040f
--- /dev/null
+++ b/src/tint/lang/core/ir/reflection.h
@@ -0,0 +1,55 @@
+// 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.
+
+#ifndef SRC_TINT_LANG_CORE_IR_REFLECTION_H_
+#define SRC_TINT_LANG_CORE_IR_REFLECTION_H_
+
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/utils/result.h"
+
+namespace tint::core::ir {
+/// Workgroup size information
+struct WorkgroupInfo {
+ /// The x-component
+ uint32_t x = 0;
+ /// The y-component
+ uint32_t y = 0;
+ /// The z-component
+ uint32_t z = 0;
+
+ /// The needed workgroup storage size
+ size_t storage_size = 0;
+};
+
+/// Generate WorkgroupInfo for an IR module input.
+/// @param ir the workgroup info for the IR module
+/// @returns the resulting WorkgroupInfo for the IR, or failure.
+Result<WorkgroupInfo> GetWorkgroupInfo(core::ir::Module& ir);
+
+} // namespace tint::core::ir
+
+#endif // SRC_TINT_LANG_CORE_IR_REFLECTION_H_
diff --git a/src/tint/lang/core/ir/reflection_test.cc b/src/tint/lang/core/ir/reflection_test.cc
new file mode 100644
index 0000000..b43b1dc
--- /dev/null
+++ b/src/tint/lang/core/ir/reflection_test.cc
@@ -0,0 +1,175 @@
+// 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/core/ir/reflection.h"
+
+#include <string>
+
+#include "gmock/gmock.h"
+#include "src/tint/lang/core/ir/disassembler.h"
+#include "src/tint/lang/core/ir/ir_helper_test.h"
+#include "src/tint/utils/result.h"
+
+namespace tint::core::ir {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+class IR_ReflectionTest : public IRTestHelper {
+ protected:
+ /// @returns the module as a disassembled string
+ std::string Disassemble() const { return "\n" + ir::Disassembler(mod).Plain(); }
+};
+
+TEST_F(IR_ReflectionTest, GetWorkgroupInfoBasic) {
+ auto* var_a = mod.root_block->Append(b.Var<workgroup, u32>("a"));
+ auto* foo = b.ComputeFunction("foo", 3_u, 5_u, 7_u);
+ b.Append(foo->Block(), [&] { //
+ b.Load(var_a);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %a:ptr<workgroup, u32, read_write> = var undef
+}
+
+%foo = @compute @workgroup_size(3u, 5u, 7u) func():void {
+ $B2: {
+ %3:u32 = load %a
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, Disassemble());
+
+ auto res = GetWorkgroupInfo(mod);
+ EXPECT_TRUE(res == tint::Success);
+ EXPECT_EQ(res->x, 3u);
+ EXPECT_EQ(res->y, 5u);
+ EXPECT_EQ(res->z, 7u);
+ EXPECT_EQ(res->storage_size, 16u);
+}
+
+TEST_F(IR_ReflectionTest, GetWorkgroupInfoMultiVar) {
+ auto* var_a = mod.root_block->Append(b.Var<workgroup, u32>("a"));
+ auto* var_b = mod.root_block->Append(b.Var<workgroup, u32>("b"));
+ auto* var_c = mod.root_block->Append(b.Var<workgroup, u32>("c"));
+ auto* var_d = mod.root_block->Append(b.Var<workgroup, u32>("d"));
+ auto* foo = b.ComputeFunction("foo", 128_u, 1_u, 1_u);
+ b.Append(foo->Block(), [&] { //
+ b.Load(var_a);
+ b.Load(var_b);
+ b.Load(var_c);
+ b.Load(var_d);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %a:ptr<workgroup, u32, read_write> = var undef
+ %b:ptr<workgroup, u32, read_write> = var undef
+ %c:ptr<workgroup, u32, read_write> = var undef
+ %d:ptr<workgroup, u32, read_write> = var undef
+}
+
+%foo = @compute @workgroup_size(128u, 1u, 1u) func():void {
+ $B2: {
+ %6:u32 = load %a
+ %7:u32 = load %b
+ %8:u32 = load %c
+ %9:u32 = load %d
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, Disassemble());
+
+ auto res = GetWorkgroupInfo(mod);
+ EXPECT_TRUE(res == tint::Success);
+ EXPECT_EQ(res->x, 128u);
+ EXPECT_EQ(res->y, 1u);
+ EXPECT_EQ(res->z, 1u);
+ EXPECT_EQ(res->storage_size, 64u);
+}
+
+TEST_F(IR_ReflectionTest, GetWorkgroupInfoNoVar) {
+ auto* foo = b.ComputeFunction("foo", 128_u, 1_u, 1_u);
+ b.Append(foo->Block(), [&] { //
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(128u, 1u, 1u) func():void {
+ $B1: {
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, Disassemble());
+
+ auto res = GetWorkgroupInfo(mod);
+ EXPECT_TRUE(res == tint::Success);
+ EXPECT_EQ(res->x, 128u);
+ EXPECT_EQ(res->y, 1u);
+ EXPECT_EQ(res->z, 1u);
+ EXPECT_EQ(res->storage_size, 0u);
+}
+
+TEST_F(IR_ReflectionTest, GetWorkgroupInfoFailNoWorkgroupSize) {
+ // Referenced.
+ auto* var_a = mod.root_block->Append(b.Var<workgroup, u32>("a"));
+ auto* foo = b.Function("foo", ty.void_(), Function::PipelineStage::kCompute);
+ b.Append(foo->Block(), [&] { //
+ b.Load(var_a);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %a:ptr<workgroup, u32, read_write> = var undef
+}
+
+%foo = @compute func():void {
+ $B2: {
+ %3:u32 = load %a
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, Disassemble());
+
+ auto res = GetWorkgroupInfo(mod);
+ EXPECT_FALSE(res == tint::Success);
+
+ auto* failure_msg = R"(IR GetWorkgroupInfo: Could not find workgroup size)";
+ EXPECT_EQ(failure_msg, res.Failure().reason);
+}
+
+} // namespace
+} // namespace tint::core::ir
