[tint] Emit expressions in workgroup_size() to the root block
An instruction produced from an expression inside the workgroup_size
attribute needs to be emitted into the root block so that it is in
scope before the function declaration.
Fixed: 408034111
Change-Id: I9943c938818df66a2c5db209aa81e76c0dcb8f52
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/234655
Reviewed-by: Peter McNeeley <petermcneeley@google.com>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index 8047666..53f65f4 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -288,6 +288,8 @@
auto attr = ast::GetAttribute<ast::WorkgroupAttribute>(ast_func->attributes);
if (attr) {
+ TINT_SCOPED_ASSIGNMENT(current_block_, mod.root_block);
+
// The x size is always required (y, z are optional).
auto value_x = EmitValueExpression(attr->x);
bool is_unsigned = value_x->Type()->IsUnsignedIntegerScalar();
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
index 93e9fb8..35bab5b 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
@@ -1440,5 +1440,61 @@
)");
}
+TEST_F(IR_FromProgramTest, OverrideInExpressionInArraySize) {
+ auto* src = R"(
+override x : u32;
+var<workgroup> arr : array<u32, x*2>;
+
+@compute @workgroup_size(64)
+fn main() {
+ _ = arr[0];
+}
+)";
+
+ auto res = Build(src);
+ ASSERT_EQ(res, Success);
+
+ auto m = res.Move();
+ EXPECT_EQ(core::ir::Disassembler(m).Plain(), R"($B1: { # root
+ %x:u32 = override undef @id(0)
+ %2:u32 = mul %x, 2u
+ %arr:ptr<workgroup, array<u32, %2>, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(64i, 1i, 1i) func():void {
+ $B2: {
+ %5:ptr<workgroup, u32, read_write> = access %arr, 0i
+ %6:u32 = load %5
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_FromProgramTest, OverrideInExpressionInWorkgroupSizeAttribute) {
+ auto* src = R"(
+override x : u32;
+
+@compute @workgroup_size(x * 2)
+fn main() { }
+)";
+
+ auto res = Build(src);
+ ASSERT_EQ(res, Success);
+
+ auto m = res.Move();
+ EXPECT_EQ(core::ir::Disassembler(m).Plain(), R"($B1: { # root
+ %x:u32 = override undef @id(0)
+ %2:u32 = mul %x, 2u
+}
+
+%main = @compute @workgroup_size(%2, 1u, 1u) func():void {
+ $B2: {
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::wgsl::reader
