[tint][ir] Always emit ir::Let for WGSL lets
Instead of only emitting them when a value cannot be named.
Preserves variable naming when going through WGSL -> textual writers.
Change-Id: Ifbd7d994f8e3830bc363f3c89fe8989fd4817e55
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/163240
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
index f72262b..7b8a407 100644
--- a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
+++ b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
@@ -2168,7 +2168,7 @@
}
fn f() {
- len_S();
+ let n = len_S();
}
)";
@@ -2200,7 +2200,7 @@
}
fn f() {
- load_W();
+ let v = load_W();
}
)";
@@ -2369,33 +2369,33 @@
fn b() {
let I = 3i;
let J = 4i;
- fn_u_U();
- fn_u_U_str_i();
- fn_u_U_arr_X(array<u32, 1u>(u32(0i)));
- fn_u_U_arr_X(array<u32, 1u>(u32(1i)));
- fn_u_U_arr_X(array<u32, 1u>(u32(I)));
- fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(1i), u32(0i)));
- fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(2i), u32(I)));
- fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(I), u32(2i)));
- fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(I), u32(J)));
- fn_s_S();
- fn_s_S_str_i();
- fn_s_S_arr_X(array<u32, 1u>(u32(0i)));
- fn_s_S_arr_X(array<u32, 1u>(u32(1i)));
- fn_s_S_arr_X(array<u32, 1u>(u32(I)));
- fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(1i), u32(0i)));
- fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(2i), u32(I)));
- fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(I), u32(2i)));
- fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(I), u32(J)));
- fn_w_W();
- fn_w_W_str_i();
- fn_w_W_arr_X(array<u32, 1u>(u32(0i)));
- fn_w_W_arr_X(array<u32, 1u>(u32(1i)));
- fn_w_W_arr_X(array<u32, 1u>(u32(I)));
- fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(1i), u32(0i)));
- fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(2i), u32(I)));
- fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(I), u32(2i)));
- fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(I), u32(J)));
+ let u = fn_u_U();
+ let u_str = fn_u_U_str_i();
+ let u_arr0 = fn_u_U_arr_X(array<u32, 1u>(u32(0i)));
+ let u_arr1 = fn_u_U_arr_X(array<u32, 1u>(u32(1i)));
+ let u_arrI = fn_u_U_arr_X(array<u32, 1u>(u32(I)));
+ let u_arr1_arr0 = fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(1i), u32(0i)));
+ let u_arr2_arrI = fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(2i), u32(I)));
+ let u_arrI_arr2 = fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(I), u32(2i)));
+ let u_arrI_arrJ = fn_u_U_arr_arr_X_X(array<u32, 2u>(u32(I), u32(J)));
+ let s = fn_s_S();
+ let s_str = fn_s_S_str_i();
+ let s_arr0 = fn_s_S_arr_X(array<u32, 1u>(u32(0i)));
+ let s_arr1 = fn_s_S_arr_X(array<u32, 1u>(u32(1i)));
+ let s_arrI = fn_s_S_arr_X(array<u32, 1u>(u32(I)));
+ let s_arr1_arr0 = fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(1i), u32(0i)));
+ let s_arr2_arrI = fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(2i), u32(I)));
+ let s_arrI_arr2 = fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(I), u32(2i)));
+ let s_arrI_arrJ = fn_s_S_arr_arr_X_X(array<u32, 2u>(u32(I), u32(J)));
+ let w = fn_w_W();
+ let w_str = fn_w_W_str_i();
+ let w_arr0 = fn_w_W_arr_X(array<u32, 1u>(u32(0i)));
+ let w_arr1 = fn_w_W_arr_X(array<u32, 1u>(u32(1i)));
+ let w_arrI = fn_w_W_arr_X(array<u32, 1u>(u32(I)));
+ let w_arr1_arr0 = fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(1i), u32(0i)));
+ let w_arr2_arrI = fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(2i), u32(I)));
+ let w_arrI_arr2 = fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(I), u32(2i)));
+ let w_arrI_arrJ = fn_w_W_arr_arr_X_X(array<u32, 2u>(u32(I), u32(J)));
}
)";
@@ -2436,7 +2436,7 @@
}
fn c() {
- b_S_X(array<u32, 1u>(u32(42i)));
+ let v = b_S_X(array<u32, 1u>(u32(42i)));
}
)";
@@ -2481,7 +2481,7 @@
}
fn c() {
- b_S_X(array<u32, 1u>(u32(42i)));
+ let v = b_S_X(array<u32, 1u>(u32(42i)));
}
)";
@@ -2525,7 +2525,7 @@
}
fn c() {
- b_S_X_U_X(array<u32, 1u>(u32(42i)), array<u32, 1u>(u32(24i)));
+ let v = b_S_X_U_X(array<u32, 1u>(u32(42i)), array<u32, 1u>(u32(24i)));
}
)";
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
index 3f939fa..9e666f9 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
@@ -58,7 +58,8 @@
%b1 = block {
%a:ptr<function, array<u32, 3>, read_write> = var
%3:ptr<function, u32, read_write> = access %a, 2u
- %b:u32 = load %3
+ %4:u32 = load %3
+ %b:u32 = let %4
ret
}
}
@@ -83,8 +84,10 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:vec3<u32> = let vec3<u32>(0u)
- %b:u32 = access %a, 2u
- %c:u32 = access %a, 1u
+ %3:u32 = access %a, 2u
+ %b:u32 = let %3
+ %5:u32 = access %a, 1u
+ %c:u32 = let %5
ret
}
}
@@ -106,7 +109,8 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, vec3<u32>, read_write> = var
- %b:u32 = load_vector_element %a, 2u
+ %3:u32 = load_vector_element %a, 2u
+ %b:u32 = let %3
ret
}
}
@@ -129,7 +133,8 @@
%b1 = block {
%a:ptr<function, array<array<f32, 4>, 3>, read_write> = var
%3:ptr<function, f32, read_write> = access %a, 2u, 3u
- %b:f32 = load %3
+ %4:f32 = load %3
+ %b:f32 = let %4
ret
}
}
@@ -152,7 +157,8 @@
%b1 = block {
%a:ptr<function, mat3x4<f32>, read_write> = var
%3:ptr<function, vec4<f32>, read_write> = access %a, 2u
- %b:f32 = load_vector_element %3, 3u
+ %4:f32 = load_vector_element %3, 3u
+ %b:f32 = let %4
ret
}
}
@@ -183,7 +189,8 @@
%b1 = block {
%a:ptr<function, MyStruct, read_write> = var
%3:ptr<function, i32, read_write> = access %a, 0u
- %b:i32 = load %3
+ %4:i32 = load %3
+ %b:i32 = let %4
ret
}
}
@@ -224,7 +231,8 @@
%b1 = block {
%a:ptr<function, Outer, read_write> = var
%3:ptr<function, f32, read_write> = access %a, 1u, 0u
- %b:f32 = load %3
+ %4:f32 = load %3
+ %b:f32 = let %4
ret
}
}
@@ -271,7 +279,8 @@
%b1 = block {
%a:ptr<function, array<Outer, 4>, read_write> = var
%3:ptr<function, vec4<f32>, read_write> = access %a, 0u, 1u, 1u, 2u
- %b:vec4<f32> = load %3
+ %4:vec4<f32> = load %3
+ %b:vec4<f32> = let %4
ret
}
}
@@ -316,7 +325,8 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, vec2<f32>, read_write> = var
- %b:f32 = load_vector_element %a, 1u
+ %3:f32 = load_vector_element %a, 1u
+ %b:f32 = let %3
ret
}
}
@@ -339,7 +349,8 @@
%b1 = block {
%a:ptr<function, vec3<f32>, read_write> = var
%3:vec3<f32> = load %a
- %b:vec4<f32> = swizzle %3, zyxz
+ %4:vec4<f32> = swizzle %3, zyxz
+ %b:vec4<f32> = let %4
ret
}
}
@@ -363,7 +374,8 @@
%a:ptr<function, vec3<f32>, read_write> = var
%3:vec3<f32> = load %a
%4:vec3<f32> = swizzle %3, zyx
- %b:vec2<f32> = swizzle %4, yy
+ %5:vec2<f32> = swizzle %4, yy
+ %b:vec2<f32> = let %5
ret
}
}
@@ -401,7 +413,8 @@
%4:vec4<f32> = load %3
%5:vec3<f32> = swizzle %4, zyx
%6:vec2<f32> = swizzle %5, yx
- %b:f32 = access %6, 0u
+ %7:f32 = access %6, 0u
+ %b:f32 = let %7
ret
}
}
@@ -422,7 +435,8 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:vec3<u32> = let vec3<u32>(0u)
- %b:u32 = access %a, 2u
+ %3:u32 = access %a, 2u
+ %b:u32 = let %3
ret
}
}
@@ -444,7 +458,8 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:mat3x4<f32> = let mat3x4<f32>(vec4<f32>(0.0f))
- %b:f32 = access %a, 2u, 3u
+ %3:f32 = access %a, 2u, 3u
+ %b:f32 = let %3
ret
}
}
@@ -474,7 +489,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:MyStruct = let MyStruct(0i)
- %b:i32 = access %a, 0u
+ %3:i32 = access %a, 0u
+ %b:i32 = let %3
ret
}
}
@@ -514,7 +530,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:Outer = let Outer(0i, Inner(0.0f))
- %b:f32 = access %a, 1u, 0u
+ %3:f32 = access %a, 1u, 0u
+ %b:f32 = let %3
ret
}
}
@@ -560,7 +577,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:array<Outer, 4> = let array<Outer, 4>(Outer(0i, array<Inner, 4>(Inner(0i, 0.0f, vec4<f32>(0.0f)))))
- %b:vec4<f32> = access %a, 0u, 1u, 1u, 2u
+ %3:vec4<f32> = access %a, 0u, 1u, 1u, 2u
+ %b:vec4<f32> = let %3
ret
}
}
@@ -582,7 +600,8 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:vec2<f32> = let vec2<f32>(0.0f)
- %b:f32 = access %a, 1u
+ %3:f32 = access %a, 1u
+ %b:f32 = let %3
ret
}
}
@@ -604,7 +623,8 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:vec3<f32> = let vec3<f32>(0.0f)
- %b:vec4<f32> = swizzle %a, zyxz
+ %3:vec4<f32> = swizzle %a, zyxz
+ %b:vec4<f32> = let %3
ret
}
}
@@ -627,7 +647,8 @@
%b1 = block {
%a:vec3<f32> = let vec3<f32>(0.0f)
%3:vec3<f32> = swizzle %a, zyx
- %b:vec2<f32> = swizzle %3, yy
+ %4:vec2<f32> = swizzle %3, yy
+ %b:vec2<f32> = let %4
ret
}
}
@@ -664,7 +685,8 @@
%3:vec4<f32> = access %a, 1u
%4:vec3<f32> = swizzle %3, zyx
%5:vec2<f32> = swizzle %4, yx
- %b:f32 = access %5, 0u
+ %6:f32 = access %5, 0u
+ %b:f32 = let %6
ret
}
}
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc
index 31d31ef..11df6fa 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc
@@ -54,7 +54,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = add %3, 4u
+ %4:u32 = add %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -123,7 +124,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = sub %3, 4u
+ %4:u32 = sub %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -192,7 +194,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = mul %3, 4u
+ %4:u32 = mul %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -238,7 +241,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = div %3, 4u
+ %4:u32 = div %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -284,7 +288,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = mod %3, 4u
+ %4:u32 = mod %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -330,7 +335,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = and %3, 4u
+ %4:u32 = and %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -376,7 +382,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = or %3, 4u
+ %4:u32 = or %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -422,7 +429,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = xor %3, 4u
+ %4:u32 = xor %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -469,7 +477,7 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
- %logical_and:bool = if %3 [t: %b3, f: %b4] { # if_1
+ %4:bool = if %3 [t: %b3, f: %b4] { # if_1
%b3 = block { # true
exit_if false # if_1
}
@@ -477,6 +485,7 @@
exit_if false # if_1
}
}
+ %logical_and:bool = let %4
if %logical_and [t: %b5] { # if_2
%b5 = block { # true
exit_if # if_2
@@ -505,7 +514,7 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
- %logical_or:bool = if %3 [t: %b3, f: %b4] { # if_1
+ %4:bool = if %3 [t: %b3, f: %b4] { # if_1
%b3 = block { # true
exit_if true # if_1
}
@@ -513,6 +522,7 @@
exit_if true # if_1
}
}
+ %logical_or:bool = let %4
if %logical_or [t: %b5] { # if_2
%b5 = block { # true
exit_if # if_2
@@ -540,7 +550,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:bool = eq %3, 4u
+ %4:bool = eq %3, 4u
+ %tint_symbol:bool = let %4
ret
}
}
@@ -563,7 +574,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:bool = neq %3, 4u
+ %4:bool = neq %3, 4u
+ %tint_symbol:bool = let %4
ret
}
}
@@ -586,7 +598,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:bool = lt %3, 4u
+ %4:bool = lt %3, 4u
+ %tint_symbol:bool = let %4
ret
}
}
@@ -609,7 +622,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:bool = gt %3, 4u
+ %4:bool = gt %3, 4u
+ %tint_symbol:bool = let %4
ret
}
}
@@ -632,7 +646,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:bool = lte %3, 4u
+ %4:bool = lte %3, 4u
+ %tint_symbol:bool = let %4
ret
}
}
@@ -655,7 +670,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:bool = gte %3, 4u
+ %4:bool = gte %3, 4u
+ %tint_symbol:bool = let %4
ret
}
}
@@ -678,7 +694,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = shiftl %3, 4u
+ %4:u32 = shiftl %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -724,7 +741,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = shiftr %3, 4u
+ %4:u32 = shiftr %3, 4u
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -773,7 +791,7 @@
%b2 = block {
%3:f32 = call %my_func
%4:bool = lt %3, 2.0f
- %tint_symbol:bool = if %4 [t: %b3, f: %b4] { # if_1
+ %5:bool = if %4 [t: %b3, f: %b4] { # if_1
%b3 = block { # true
%6:f32 = call %my_func
%7:f32 = call %my_func
@@ -786,6 +804,7 @@
exit_if false # if_1
}
}
+ %tint_symbol:bool = let %5
ret
}
}
@@ -808,7 +827,8 @@
}
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
- %tint_symbol:bool = call %my_func, false
+ %4:bool = call %my_func, false
+ %tint_symbol:bool = let %4
ret
}
}
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc
index fd6f77e..b92e960 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc
@@ -53,7 +53,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = load %i
- %tint_symbol:f32 = asin %3
+ %4:f32 = asin %3
+ %tint_symbol:f32 = let %4
ret
}
}
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc
index 4c1dbfc..52cd00c 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc
@@ -56,7 +56,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = call %my_func
- %tint_symbol:f32 = bitcast %3
+ %4:f32 = bitcast %3
+ %tint_symbol:f32 = let %4
ret
}
}
@@ -119,7 +120,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = load %i
- %tint_symbol:f32 = convert %3
+ %4:f32 = convert %3
+ %tint_symbol:f32 = let %4
ret
}
}
@@ -155,7 +157,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = load %i
- %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
+ %4:vec3<f32> = construct 2.0f, 3.0f, %3
+ %tint_symbol:vec3<f32> = let %4
ret
}
}
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 fa33289..4644964 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
@@ -1334,28 +1334,15 @@
builder_.ir.SetName(val, v->name->symbol.Name());
},
[&](const ast::Let* l) {
- auto* last_stmt = current_block_->Back();
auto init = EmitValueExpression(l->initializer);
if (!init) {
return;
}
- auto* value = init;
- if (current_block_->Back() == last_stmt) {
- // Emitting the let's initializer didn't create an instruction.
- // Create an core::ir::Let to give the let an instruction. This gives the let a
- // place of declaration and name, which preserves runtime semantics of the
- // let, and can be used by consumers of the IR to produce a variable or
- // debug info.
- auto* let = current_block_->Append(builder_.Let(l->name->symbol.Name(), value));
- value = let->Result(0);
- } else {
- // Record the original name of the let
- builder_.ir.SetName(value, l->name->symbol.Name());
- }
+ auto* let = current_block_->Append(builder_.Let(l->name->symbol.Name(), init));
// Store the results of the initialization
- scopes_.Set(l->name->symbol, value);
+ scopes_.Set(l->name->symbol, let->Result(0));
},
[&](const ast::Override*) {
add_error(var->source,
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc
index 1536188..09af944 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc
@@ -59,8 +59,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-S = struct @align(4) {
+ EXPECT_EQ(Disassemble(m.Get()), R"(S = struct @align(4) {
i:i32 @offset(0)
}
@@ -88,8 +87,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-S = struct @align(4) {
+ EXPECT_EQ(Disassemble(m.Get()), R"(S = struct @align(4) {
i:i32 @offset(0)
}
@@ -115,8 +113,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%b1 = block { # root
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%i:ptr<private, i32, read_write> = var, 1i
}
@@ -148,8 +145,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%b1 = block { # root
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%i:ptr<private, i32, read_write> = var, 1i
}
@@ -159,7 +155,8 @@
%4:i32 = add %3, 1i
store %i, %4
%5:i32 = load %i
- %i_1:i32 = add %5, 1i # %i_1: 'i'
+ %6:i32 = add %5, 1i
+ %i_1:i32 = let %6 # %i_1: 'i'
ret %i_1
}
}
@@ -181,8 +178,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
if true [t: %b2] { # if_1
@@ -221,8 +217,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
if true [t: %b2] { # if_1
@@ -231,12 +226,13 @@
%4:i32 = add %3, 1i
store %i, %4
%5:i32 = load %i
- %i_1:i32 = add %5, 1i # %i_1: 'i'
+ %6:i32 = add %5, 1i
+ %i_1:i32 = let %6 # %i_1: 'i'
ret %i_1
}
}
- %7:i32 = load %i
- ret %7
+ %8:i32 = load %i
+ ret %8
}
}
)");
@@ -256,8 +252,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [b: %b2, c: %b3] { # loop_1
@@ -303,8 +298,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [b: %b2, c: %b3] { # loop_1
@@ -320,15 +314,16 @@
}
}
%5:i32 = load %i
- %i_1:i32 = add %5, 1i # %i_1: 'i'
+ %6:i32 = add %5, 1i
+ %i_1:i32 = let %6 # %i_1: 'i'
ret %i_1
}
%b3 = block { # continuing
next_iteration %b2
}
}
- %7:i32 = load %i
- ret %7
+ %8:i32 = load %i
+ ret %8
}
}
)");
@@ -347,8 +342,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [i: %b2, b: %b3] { # loop_1
@@ -367,12 +361,13 @@
exit_loop # loop_1
}
}
- %j:f32 = load %i_1
+ %6:f32 = load %i_1
+ %j:f32 = let %6
continue %b6
}
}
- %7:i32 = load %i
- ret %7
+ %8:i32 = load %i
+ ret %8
}
}
)");
@@ -391,8 +386,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [i: %b2, b: %b3] { # loop_1
@@ -435,8 +429,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [i: %b2, b: %b3] { # loop_1
@@ -483,8 +476,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [i: %b2, b: %b3] { # loop_1
@@ -504,12 +496,13 @@
}
}
%6:i32 = load %i
- %i_1:i32 = add %6, 1i # %i_1: 'i'
+ %7:i32 = add %6, 1i
+ %i_1:i32 = let %7 # %i_1: 'i'
ret %i_1
}
}
- %8:i32 = load %i
- ret %8
+ %9:i32 = load %i
+ ret %9
}
}
)");
@@ -534,8 +527,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [b: %b2, c: %b3] { # loop_1
@@ -589,8 +581,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [b: %b2, c: %b3] { # loop_1
@@ -603,9 +594,10 @@
}
}
%5:i32 = load %i
- %i_1:i32 = add %5, 1i # %i_1: 'i'
- %7:bool = eq %i_1, 3i
- if %7 [t: %b5] { # if_2
+ %6:i32 = add %5, 1i
+ %i_1:i32 = let %6 # %i_1: 'i'
+ %8:bool = eq %i_1, 3i
+ if %8 [t: %b5] { # if_2
%b5 = block { # true
exit_loop # loop_1
}
@@ -616,8 +608,8 @@
next_iteration %b2
}
}
- %8:i32 = load %i
- ret %8
+ %9:i32 = load %i
+ ret %9
}
}
)");
@@ -643,8 +635,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [b: %b2, c: %b3] { # loop_1
@@ -694,8 +685,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
loop [b: %b2, c: %b3] { # loop_1
@@ -711,13 +701,14 @@
}
%b3 = block { # continuing
%5:i32 = load %i
- %i_1:i32 = add %5, 1i # %i_1: 'i'
- %7:bool = gt %i_1, 2i
- break_if %7 %b2
+ %6:i32 = add %5, 1i
+ %i_1:i32 = let %6 # %i_1: 'i'
+ %8:bool = gt %i_1, 2i
+ break_if %8 %b2
}
}
- %8:i32 = load %i
- ret %8
+ %9:i32 = load %i
+ ret %9
}
}
)");
@@ -744,8 +735,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
%3:i32 = load %i
@@ -793,8 +783,7 @@
ASSERT_TRUE(m) << m;
- EXPECT_EQ("\n" + Disassemble(m.Get()), R"(
-%f = func():i32 -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 -> %b1 {
%b1 = block {
%i:ptr<function, i32, read_write> = var
%3:i32 = load %i
@@ -805,12 +794,13 @@
}
%b3 = block { # case
%5:i32 = load %i
- %i_1:i32 = add %5, 1i # %i_1: 'i'
+ %6:i32 = add %5, 1i
+ %i_1:i32 = let %6 # %i_1: 'i'
ret %i_1
}
%b4 = block { # case
- %7:i32 = load %i
- ret %7
+ %8:i32 = load %i
+ ret %8
}
}
unreachable
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc
index 0db5a17..98ff492 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc
@@ -54,7 +54,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
- %tint_symbol:bool = eq %3, false
+ %4:bool = eq %3, false
+ %tint_symbol:bool = let %4
ret
}
}
@@ -77,7 +78,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:vec4<bool> = call %my_func
- %tint_symbol:vec4<bool> = eq %3, vec4<bool>(false)
+ %4:vec4<bool> = eq %3, vec4<bool>(false)
+ %tint_symbol:vec4<bool> = let %4
ret
}
}
@@ -100,7 +102,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
- %tint_symbol:u32 = complement %3
+ %4:u32 = complement %3
+ %tint_symbol:u32 = let %4
ret
}
}
@@ -123,7 +126,8 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = call %my_func
- %tint_symbol:i32 = negation %3
+ %4:i32 = negation %3
+ %tint_symbol:i32 = let %4
ret
}
}
