HLSL-IR: implement NumWorkgroups builtin support
Bug: 357896924
Bug: 42251045
Change-Id: Icd426b9a62cd077affae59ccacdbe51b5aa8276f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/206514
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/lang/core/ir/transform/helper_test.h b/src/tint/lang/core/ir/transform/helper_test.h
index e0e2583..22ed109 100644
--- a/src/tint/lang/core/ir/transform/helper_test.h
+++ b/src/tint/lang/core/ir/transform/helper_test.h
@@ -51,7 +51,7 @@
template <typename TRANSFORM, typename... ARGS>
void Run(TRANSFORM&& transform_func, ARGS&&... args) {
// Run the transform.
- auto result = transform_func(mod, args...);
+ auto result = transform_func(mod, std::forward<ARGS>(args)...);
EXPECT_EQ(result, Success);
if (result != Success) {
return;
diff --git a/src/tint/lang/hlsl/writer/raise/raise.cc b/src/tint/lang/hlsl/writer/raise/raise.cc
index 607fcd43..3e49693 100644
--- a/src/tint/lang/hlsl/writer/raise/raise.cc
+++ b/src/tint/lang/hlsl/writer/raise/raise.cc
@@ -153,6 +153,19 @@
RUN_TRANSFORM(core::ir::transform::Robustness, module, config);
}
+ if (!options.disable_workgroup_init) {
+ // Must run before ShaderIO as it may introduce a builtin parameter (local_invocation_index)
+ RUN_TRANSFORM(core::ir::transform::ZeroInitWorkgroupMemory, module);
+ }
+
+ // ShaderIO must be run before DecomposeUniformAccess because it might
+ // introduce a uniform buffer for kNumWorkgroups.
+ {
+ raise::ShaderIOConfig config;
+ config.num_workgroups_binding = options.root_constant_binding_point;
+ RUN_TRANSFORM(raise::ShaderIO, module, config);
+ }
+
RUN_TRANSFORM(core::ir::transform::DirectVariableAccess, module,
core::ir::transform::DirectVariableAccessOptions{});
// DecomposeStorageAccess must come after Robustness and DirectVariableAccess
@@ -160,18 +173,12 @@
// Comes after DecomposeStorageAccess.
RUN_TRANSFORM(raise::DecomposeUniformAccess, module);
- if (!options.disable_workgroup_init) {
- RUN_TRANSFORM(core::ir::transform::ZeroInitWorkgroupMemory, module);
- }
-
// TODO(dsinclair): LocalizeStructArrayAssignment
// TODO(dsinclair): PixelLocal transform
// TODO(dsinclair): TruncateInterstageVariables
- // TODO(dsinclair): NumWorkgroupsFromUniform
// TODO(dsinclair): CalculateArrayLength
// TODO(dsinclair): RemoveContinueInSwitch
- RUN_TRANSFORM(raise::ShaderIO, module);
// DemoteToHelper must come before any transform that introduces non-core instructions.
// Run after ShaderIO to ensure the discards are added to the entry point it introduces.
RUN_TRANSFORM(core::ir::transform::DemoteToHelper, module);
diff --git a/src/tint/lang/hlsl/writer/raise/shader_io.cc b/src/tint/lang/hlsl/writer/raise/shader_io.cc
index df7e3dc..0da0c57 100644
--- a/src/tint/lang/hlsl/writer/raise/shader_io.cc
+++ b/src/tint/lang/hlsl/writer/raise/shader_io.cc
@@ -49,6 +49,9 @@
/// For HLSL, move all inputs to a struct passed as an entry point parameter, and wrap outputs in
/// a structure returned by the entry point.
struct StateImpl : core::ir::transform::ShaderIOBackendState {
+ /// The config
+ const ShaderIOConfig& config;
+
/// The input parameter
core::ir::FunctionParam* input_param = nullptr;
@@ -61,12 +64,14 @@
/// The output values to return from the entry point.
Vector<core::ir::Value*, 4> output_values;
- // Indices of subgroup invocation id and size, if set
+ // Indices of inputs that require special handling
std::optional<uint32_t> subgroup_invocation_id_index;
std::optional<uint32_t> subgroup_size_index;
+ std::optional<uint32_t> num_workgroups_index;
/// Constructor
- StateImpl(core::ir::Module& mod, core::ir::Function* f) : ShaderIOBackendState(mod, f) {}
+ StateImpl(core::ir::Module& mod, core::ir::Function* f, const ShaderIOConfig& c)
+ : ShaderIOBackendState(mod, f), config(c) {}
/// Destructor
~StateImpl() override {}
@@ -183,6 +188,9 @@
} else if (*builtin == core::BuiltinValue::kSubgroupSize) {
subgroup_size_index = i;
continue;
+ } else if (*builtin == core::BuiltinValue::kNumWorkgroups) {
+ num_workgroups_index = i;
+ continue;
}
}
@@ -257,12 +265,45 @@
output_struct->AddUsage(core::type::PipelineStageUsage::kVertexOutput);
break;
case core::ir::Function::PipelineStage::kCompute:
+ output_struct->AddUsage(core::type::PipelineStageUsage::kComputeOutput);
+ break;
case core::ir::Function::PipelineStage::kUndefined:
TINT_UNREACHABLE();
}
return output_struct;
}
+ /// Handles kNumWorkgroups builtin by emitting a UBO to hold the num_workgroups value,
+ /// along with the load of the value. Returns the loaded value.
+ core::ir::Value* GetInputForNumWorkgroups(core::ir::Builder& builder) {
+ // Create uniform var that will receive the number of workgroups
+ core::ir::Var* num_wg_var = nullptr;
+ builder.Append(ir.root_block, [&] {
+ num_wg_var = builder.Var("tint_num_workgroups", ty.ptr(uniform, ty.vec3<u32>()));
+ });
+ if (config.num_workgroups_binding.has_value()) {
+ // If config.num_workgroups_binding holds a value, use it.
+ auto bp = *config.num_workgroups_binding;
+ num_wg_var->SetBindingPoint(bp.group, bp.binding);
+ } else {
+ // Otherwise, use the binding 0 of the largest used group plus 1, or group 0 if no
+ // resources are bound.
+ uint32_t group = 0;
+ for (auto* inst : *ir.root_block.Get()) {
+ if (auto* var = inst->As<core::ir::Var>()) {
+ if (const auto& bp = var->BindingPoint()) {
+ if (bp->group >= group) {
+ group = bp->group + 1;
+ }
+ }
+ }
+ }
+ num_wg_var->SetBindingPoint(group, 0);
+ }
+ auto* load = builder.Load(num_wg_var);
+ return load->Result(0);
+ }
+
/// @copydoc ShaderIO::BackendState::GetInput
core::ir::Value* GetInput(core::ir::Builder& builder, uint32_t idx) override {
if (subgroup_invocation_id_index == idx) {
@@ -275,6 +316,9 @@
.Call<hlsl::ir::BuiltinCall>(ty.u32(), hlsl::BuiltinFn::kWaveGetLaneCount)
->Result(0);
}
+ if (num_workgroups_index == idx) {
+ return GetInputForNumWorkgroups(builder);
+ }
auto index = input_indices[idx];
@@ -307,14 +351,14 @@
};
} // namespace
-Result<SuccessType> ShaderIO(core::ir::Module& ir) {
+Result<SuccessType> ShaderIO(core::ir::Module& ir, const ShaderIOConfig& config) {
auto result = ValidateAndDumpIfNeeded(ir, "ShaderIO transform");
if (result != Success) {
return result;
}
core::ir::transform::RunShaderIOBase(ir, [&](core::ir::Module& mod, core::ir::Function* func) {
- return std::make_unique<StateImpl>(mod, func);
+ return std::make_unique<StateImpl>(mod, func, config);
});
return Success;
diff --git a/src/tint/lang/hlsl/writer/raise/shader_io.h b/src/tint/lang/hlsl/writer/raise/shader_io.h
index 8bc3b86..af0e3e5 100644
--- a/src/tint/lang/hlsl/writer/raise/shader_io.h
+++ b/src/tint/lang/hlsl/writer/raise/shader_io.h
@@ -28,6 +28,9 @@
#ifndef SRC_TINT_LANG_HLSL_WRITER_RAISE_SHADER_IO_H_
#define SRC_TINT_LANG_HLSL_WRITER_RAISE_SHADER_IO_H_
+#include <optional>
+
+#include "src/tint/api/common/binding_point.h"
#include "src/tint/utils/result/result.h"
// Forward declarations.
@@ -37,10 +40,17 @@
namespace tint::hlsl::writer::raise {
+struct ShaderIOConfig {
+ /// The binding point to use for the num_workgroups generated uniform buffer. If it contains
+ /// no value, a free binding point will be used. Specifically, binding 0 of the largest used
+ /// group plus 1 is used if at least one resource is bound, otherwise group 0 binding 0 is used.
+ std::optional<BindingPoint> num_workgroups_binding;
+};
+
/// ShaderIO is a transform that prepares entry point inputs and outputs for HLSL codegen.
/// @param module the module to transform
/// @returns success or failure
-Result<SuccessType> ShaderIO(core::ir::Module& module);
+Result<SuccessType> ShaderIO(core::ir::Module& module, const ShaderIOConfig& config);
} // namespace tint::hlsl::writer::raise
diff --git a/src/tint/lang/hlsl/writer/raise/shader_io_test.cc b/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
index bd0d91d..f7cedb5 100644
--- a/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
@@ -55,7 +55,7 @@
auto* expect = src;
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -140,7 +140,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -294,7 +294,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -419,7 +419,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -461,7 +461,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -502,7 +502,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -607,7 +607,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -687,7 +687,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -835,7 +835,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -929,7 +929,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -1016,7 +1016,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -1064,7 +1064,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -1116,7 +1116,7 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
@@ -1209,7 +1209,278 @@
}
)";
- Run(ShaderIO);
+ Run(ShaderIO, ShaderIOConfig{});
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOParameters_NumWorkgroups_NonStruct) {
+ auto* num_workgroups = b.FunctionParam("num_wgs", ty.vec3<u32>());
+ num_workgroups->SetBuiltin(core::BuiltinValue::kNumWorkgroups);
+
+ auto* ep = b.Function("foo", ty.vec3<u32>(), core::ir::Function::PipelineStage::kCompute);
+ ep->SetParams({num_workgroups});
+ ep->SetWorkgroupSize(1, 1, 1);
+
+ b.Append(ep->Block(), [&] {
+ auto* r = b.Multiply(ty.vec3<u32>(), num_workgroups, num_workgroups);
+ b.Return(ep, r);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1, 1, 1) func(%num_wgs:vec3<u32> [@num_workgroups]):vec3<u32> {
+ $B1: {
+ %3:vec3<u32> = mul %num_wgs, %num_wgs
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+foo_outputs = struct @align(16) {
+ tint_symbol:vec3<u32> @offset(0)
+}
+
+$B1: { # root
+ %tint_num_workgroups:ptr<uniform, vec3<u32>, read> = var @binding_point(0, 0)
+}
+
+%foo_inner = func(%num_wgs:vec3<u32>):vec3<u32> {
+ $B2: {
+ %4:vec3<u32> = mul %num_wgs, %num_wgs
+ ret %4
+ }
+}
+%foo = @compute @workgroup_size(1, 1, 1) func():foo_outputs {
+ $B3: {
+ %6:vec3<u32> = load %tint_num_workgroups
+ %7:vec3<u32> = call %foo_inner, %6
+ %8:foo_outputs = construct %7
+ ret %8
+ }
+}
+)";
+
+ Run(ShaderIO, ShaderIOConfig{});
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOParameters_NumWorkgroups_Struct) {
+ auto* str_ty = ty.Struct(mod.symbols.New("Inputs"),
+ {
+ {
+ mod.symbols.New("num_wgs"),
+ ty.vec3<u32>(),
+ core::IOAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kNumWorkgroups,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* str_param = b.FunctionParam("inputs", str_ty);
+
+ auto* ep = b.Function("foo", ty.vec3<u32>(), core::ir::Function::PipelineStage::kCompute);
+ ep->SetParams({str_param});
+ ep->SetWorkgroupSize(1, 1, 1);
+
+ b.Append(ep->Block(), [&] {
+ auto* num_workgroups = b.Access(ty.vec3<u32>(), str_param, 0_i);
+ auto* r = b.Multiply(ty.vec3<u32>(), num_workgroups, num_workgroups);
+ b.Return(ep, r);
+ });
+
+ auto* src = R"(
+Inputs = struct @align(16) {
+ num_wgs:vec3<u32> @offset(0), @builtin(num_workgroups)
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func(%inputs:Inputs):vec3<u32> {
+ $B1: {
+ %3:vec3<u32> = access %inputs, 0i
+ %4:vec3<u32> = mul %3, %3
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Inputs = struct @align(16) {
+ num_wgs:vec3<u32> @offset(0)
+}
+
+foo_outputs = struct @align(16) {
+ tint_symbol:vec3<u32> @offset(0)
+}
+
+$B1: { # root
+ %tint_num_workgroups:ptr<uniform, vec3<u32>, read> = var @binding_point(0, 0)
+}
+
+%foo_inner = func(%inputs:Inputs):vec3<u32> {
+ $B2: {
+ %4:vec3<u32> = access %inputs, 0i
+ %5:vec3<u32> = mul %4, %4
+ ret %5
+ }
+}
+%foo = @compute @workgroup_size(1, 1, 1) func():foo_outputs {
+ $B3: {
+ %7:vec3<u32> = load %tint_num_workgroups
+ %8:Inputs = construct %7
+ %9:vec3<u32> = call %foo_inner, %8
+ %10:foo_outputs = construct %9
+ ret %10
+ }
+}
+)";
+
+ Run(ShaderIO, ShaderIOConfig{});
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOParameters_NumWorkgroups_ExplicitBinding) {
+ auto* num_workgroups = b.FunctionParam("num_wgs", ty.vec3<u32>());
+ num_workgroups->SetBuiltin(core::BuiltinValue::kNumWorkgroups);
+
+ auto* ep = b.Function("foo", ty.vec3<u32>(), core::ir::Function::PipelineStage::kCompute);
+ ep->SetParams({num_workgroups});
+ ep->SetWorkgroupSize(1, 1, 1);
+
+ b.Append(ep->Block(), [&] {
+ auto* r = b.Multiply(ty.vec3<u32>(), num_workgroups, num_workgroups);
+ b.Return(ep, r);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1, 1, 1) func(%num_wgs:vec3<u32> [@num_workgroups]):vec3<u32> {
+ $B1: {
+ %3:vec3<u32> = mul %num_wgs, %num_wgs
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+foo_outputs = struct @align(16) {
+ tint_symbol:vec3<u32> @offset(0)
+}
+
+$B1: { # root
+ %tint_num_workgroups:ptr<uniform, vec3<u32>, read> = var @binding_point(1, 23)
+}
+
+%foo_inner = func(%num_wgs:vec3<u32>):vec3<u32> {
+ $B2: {
+ %4:vec3<u32> = mul %num_wgs, %num_wgs
+ ret %4
+ }
+}
+%foo = @compute @workgroup_size(1, 1, 1) func():foo_outputs {
+ $B3: {
+ %6:vec3<u32> = load %tint_num_workgroups
+ %7:vec3<u32> = call %foo_inner, %6
+ %8:foo_outputs = construct %7
+ ret %8
+ }
+}
+)";
+
+ ShaderIOConfig config;
+ config.num_workgroups_binding = {1u, 23u};
+ Run(ShaderIO, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOParameters_NumWorkgroups_AutoBinding) {
+ auto* num_workgroups = b.FunctionParam("num_wgs", ty.vec3<u32>());
+ num_workgroups->SetBuiltin(core::BuiltinValue::kNumWorkgroups);
+
+ auto* ep = b.Function("foo", ty.vec3<u32>(), core::ir::Function::PipelineStage::kCompute);
+ ep->SetParams({num_workgroups});
+ ep->SetWorkgroupSize(1, 1, 1);
+
+ b.Append(ep->Block(), [&] {
+ auto* r = b.Multiply(ty.vec3<u32>(), num_workgroups, num_workgroups);
+ b.Return(ep, r);
+ });
+
+ b.Append(mod.root_block, [&] {
+ for (uint32_t group = 0; group < 10; ++group) {
+ auto* v = b.Var<core::AddressSpace::kStorage, i32>();
+ v->SetBindingPoint(group, group + 1u);
+ }
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %1:ptr<storage, i32, read_write> = var @binding_point(0, 1)
+ %2:ptr<storage, i32, read_write> = var @binding_point(1, 2)
+ %3:ptr<storage, i32, read_write> = var @binding_point(2, 3)
+ %4:ptr<storage, i32, read_write> = var @binding_point(3, 4)
+ %5:ptr<storage, i32, read_write> = var @binding_point(4, 5)
+ %6:ptr<storage, i32, read_write> = var @binding_point(5, 6)
+ %7:ptr<storage, i32, read_write> = var @binding_point(6, 7)
+ %8:ptr<storage, i32, read_write> = var @binding_point(7, 8)
+ %9:ptr<storage, i32, read_write> = var @binding_point(8, 9)
+ %10:ptr<storage, i32, read_write> = var @binding_point(9, 10)
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func(%num_wgs:vec3<u32> [@num_workgroups]):vec3<u32> {
+ $B2: {
+ %13:vec3<u32> = mul %num_wgs, %num_wgs
+ ret %13
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+foo_outputs = struct @align(16) {
+ tint_symbol:vec3<u32> @offset(0)
+}
+
+$B1: { # root
+ %1:ptr<storage, i32, read_write> = var @binding_point(0, 1)
+ %2:ptr<storage, i32, read_write> = var @binding_point(1, 2)
+ %3:ptr<storage, i32, read_write> = var @binding_point(2, 3)
+ %4:ptr<storage, i32, read_write> = var @binding_point(3, 4)
+ %5:ptr<storage, i32, read_write> = var @binding_point(4, 5)
+ %6:ptr<storage, i32, read_write> = var @binding_point(5, 6)
+ %7:ptr<storage, i32, read_write> = var @binding_point(6, 7)
+ %8:ptr<storage, i32, read_write> = var @binding_point(7, 8)
+ %9:ptr<storage, i32, read_write> = var @binding_point(8, 9)
+ %10:ptr<storage, i32, read_write> = var @binding_point(9, 10)
+ %tint_num_workgroups:ptr<uniform, vec3<u32>, read> = var @binding_point(10, 0)
+}
+
+%foo_inner = func(%num_wgs:vec3<u32>):vec3<u32> {
+ $B2: {
+ %14:vec3<u32> = mul %num_wgs, %num_wgs
+ ret %14
+ }
+}
+%foo = @compute @workgroup_size(1, 1, 1) func():foo_outputs {
+ $B3: {
+ %16:vec3<u32> = load %tint_num_workgroups
+ %17:vec3<u32> = call %foo_inner, %16
+ %18:foo_outputs = construct %17
+ ret %18
+ }
+}
+)";
+
+ Run(ShaderIO, ShaderIOConfig{});
EXPECT_EQ(expect, str());
}
diff --git a/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.dxc.hlsl b/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.dxc.hlsl
index 7568472..127819c 100644
--- a/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.dxc.hlsl
@@ -1,11 +1,20 @@
-SKIP: FAILED
+struct main_inputs {
+ uint3 local_invocation_id : SV_GroupThreadID;
+ uint local_invocation_index : SV_GroupIndex;
+ uint3 global_invocation_id : SV_DispatchThreadID;
+ uint3 workgroup_id : SV_GroupID;
+};
-..\..\src\tint\lang\hlsl\writer\printer\printer.cc:1505 internal compiler error: TINT_ASSERT(!name.empty())
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
-tint executable returned error: exit status 0xc000001d
+cbuffer cbuffer_tint_num_workgroups : register(b0) {
+ uint4 tint_num_workgroups[1];
+};
+void main_inner(uint3 local_invocation_id, uint local_invocation_index, uint3 global_invocation_id, uint3 workgroup_id, uint3 num_workgroups) {
+ uint foo = ((((local_invocation_id[0u] + local_invocation_index) + global_invocation_id[0u]) + workgroup_id[0u]) + num_workgroups[0u]);
+}
+
+[numthreads(1, 1, 1)]
+void main(main_inputs inputs) {
+ main_inner(inputs.local_invocation_id, inputs.local_invocation_index, inputs.global_invocation_id, inputs.workgroup_id, tint_num_workgroups[0u].xyz);
+}
+
diff --git a/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.fxc.hlsl b/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.fxc.hlsl
index 7568472..127819c 100644
--- a/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.fxc.hlsl
@@ -1,11 +1,20 @@
-SKIP: FAILED
+struct main_inputs {
+ uint3 local_invocation_id : SV_GroupThreadID;
+ uint local_invocation_index : SV_GroupIndex;
+ uint3 global_invocation_id : SV_DispatchThreadID;
+ uint3 workgroup_id : SV_GroupID;
+};
-..\..\src\tint\lang\hlsl\writer\printer\printer.cc:1505 internal compiler error: TINT_ASSERT(!name.empty())
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
-tint executable returned error: exit status 0xc000001d
+cbuffer cbuffer_tint_num_workgroups : register(b0) {
+ uint4 tint_num_workgroups[1];
+};
+void main_inner(uint3 local_invocation_id, uint local_invocation_index, uint3 global_invocation_id, uint3 workgroup_id, uint3 num_workgroups) {
+ uint foo = ((((local_invocation_id[0u] + local_invocation_index) + global_invocation_id[0u]) + workgroup_id[0u]) + num_workgroups[0u]);
+}
+
+[numthreads(1, 1, 1)]
+void main(main_inputs inputs) {
+ main_inner(inputs.local_invocation_id, inputs.local_invocation_index, inputs.global_invocation_id, inputs.workgroup_id, tint_num_workgroups[0u].xyz);
+}
+
diff --git a/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.dxc.hlsl b/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.dxc.hlsl
index 7568472..8c73304 100644
--- a/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.dxc.hlsl
@@ -1,11 +1,29 @@
-SKIP: FAILED
+struct ComputeInputs {
+ uint3 local_invocation_id;
+ uint local_invocation_index;
+ uint3 global_invocation_id;
+ uint3 workgroup_id;
+ uint3 num_workgroups;
+};
-..\..\src\tint\lang\hlsl\writer\printer\printer.cc:1505 internal compiler error: TINT_ASSERT(!name.empty())
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct main_inputs {
+ uint3 ComputeInputs_local_invocation_id : SV_GroupThreadID;
+ uint ComputeInputs_local_invocation_index : SV_GroupIndex;
+ uint3 ComputeInputs_global_invocation_id : SV_DispatchThreadID;
+ uint3 ComputeInputs_workgroup_id : SV_GroupID;
+};
-tint executable returned error: exit status 0xc000001d
+
+cbuffer cbuffer_tint_num_workgroups : register(b0) {
+ uint4 tint_num_workgroups[1];
+};
+void main_inner(ComputeInputs inputs) {
+ uint foo = ((((inputs.local_invocation_id[0u] + inputs.local_invocation_index) + inputs.global_invocation_id[0u]) + inputs.workgroup_id[0u]) + inputs.num_workgroups[0u]);
+}
+
+[numthreads(1, 1, 1)]
+void main(main_inputs inputs) {
+ ComputeInputs v = {inputs.ComputeInputs_local_invocation_id, inputs.ComputeInputs_local_invocation_index, inputs.ComputeInputs_global_invocation_id, inputs.ComputeInputs_workgroup_id, tint_num_workgroups[0u].xyz};
+ main_inner(v);
+}
+
diff --git a/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.fxc.hlsl b/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.fxc.hlsl
index 7568472..8c73304 100644
--- a/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/types/functions/shader_io/compute_input_builtins_struct.wgsl.expected.ir.fxc.hlsl
@@ -1,11 +1,29 @@
-SKIP: FAILED
+struct ComputeInputs {
+ uint3 local_invocation_id;
+ uint local_invocation_index;
+ uint3 global_invocation_id;
+ uint3 workgroup_id;
+ uint3 num_workgroups;
+};
-..\..\src\tint\lang\hlsl\writer\printer\printer.cc:1505 internal compiler error: TINT_ASSERT(!name.empty())
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct main_inputs {
+ uint3 ComputeInputs_local_invocation_id : SV_GroupThreadID;
+ uint ComputeInputs_local_invocation_index : SV_GroupIndex;
+ uint3 ComputeInputs_global_invocation_id : SV_DispatchThreadID;
+ uint3 ComputeInputs_workgroup_id : SV_GroupID;
+};
-tint executable returned error: exit status 0xc000001d
+
+cbuffer cbuffer_tint_num_workgroups : register(b0) {
+ uint4 tint_num_workgroups[1];
+};
+void main_inner(ComputeInputs inputs) {
+ uint foo = ((((inputs.local_invocation_id[0u] + inputs.local_invocation_index) + inputs.global_invocation_id[0u]) + inputs.workgroup_id[0u]) + inputs.num_workgroups[0u]);
+}
+
+[numthreads(1, 1, 1)]
+void main(main_inputs inputs) {
+ ComputeInputs v = {inputs.ComputeInputs_local_invocation_id, inputs.ComputeInputs_local_invocation_index, inputs.ComputeInputs_global_invocation_id, inputs.ComputeInputs_workgroup_id, tint_num_workgroups[0u].xyz};
+ main_inner(v);
+}
+
