[tint][transform][hlsl] Fix bad materialization As reported in tint:1963, the PromoteSideEffectsToDecl transform could hoist a constant expression to an implicitly typed 'let' declaration, resulting in the expression type materializing early and resolve to a different type. Constant expressions don't need to be hoisted, but it seems generally safer to continue emitting these, but with an explicit type on the 'let' declaration. This is what this CL does. In doing this, I uncovered a second bug where the DecomposeMemoryAccess transform would create a new a struct for the result of the atomicCompareExchangeWeak() intrinsic, which would fail to assign to the new explicitly typed lets. The fix for this is simple - don't generate a new struct, but instead use the (fairly new) internal struct declarations for these. This keeps the resolver happy, and actually reduces a bunch of spew from the HLSL output. Fixed: tint:1963 Change-Id: I271a2130dc303da3b9f11876ec5c6e8b98a16ee7 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/137100 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: James Price <jrprice@google.com> Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/ast/transform/decompose_memory_access.cc b/src/tint/ast/transform/decompose_memory_access.cc index bd486063..c1b606c 100644 --- a/src/tint/ast/transform/decompose_memory_access.cc +++ b/src/tint/ast/transform/decompose_memory_access.cc
@@ -656,26 +656,7 @@ << el_ty->TypeInfo().name; } - Type ret_ty; - - // For intrinsics that return a struct, there is no AST node for it, so create one now. - if (intrinsic->Type() == builtin::Function::kAtomicCompareExchangeWeak) { - auto* str = intrinsic->ReturnType()->As<type::Struct>(); - TINT_ASSERT(Transform, str); - - utils::Vector<const StructMember*, 8> ast_members; - ast_members.Reserve(str->Members().Length()); - for (auto& m : str->Members()) { - ast_members.Push( - b.Member(ctx.Clone(m->Name()), CreateASTTypeFor(ctx, m->Type()))); - } - - auto name = b.Symbols().New("atomic_compare_exchange_weak_ret_type"); - auto* new_str = b.Structure(name, std::move(ast_members)); - ret_ty = b.ty.Of(new_str); - } else { - ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType()); - } + Type ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType()); auto name = b.Symbols().New(buffer.Name() + intrinsic->str()); b.Func(name, std::move(params), ret_ty, nullptr,
diff --git a/src/tint/ast/transform/decompose_memory_access_test.cc b/src/tint/ast/transform/decompose_memory_access_test.cc index 52f2ff2..3bff4a8 100644 --- a/src/tint/ast/transform/decompose_memory_access_test.cc +++ b/src/tint/ast/transform/decompose_memory_access_test.cc
@@ -3640,13 +3640,8 @@ @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body) fn sbatomicExchange(offset : u32, param_1 : i32) -> i32 -struct atomic_compare_exchange_weak_ret_type { - old_value : i32, - exchanged : bool, -} - @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body) -fn sbatomicCompareExchangeWeak(offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type +fn sbatomicCompareExchangeWeak(offset : u32, param_1 : i32, param_2 : i32) -> __atomic_compare_exchange_result_i32 @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body) fn sbatomicStore_1(offset : u32, param_1 : u32) @@ -3678,13 +3673,8 @@ @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body) fn sbatomicExchange_1(offset : u32, param_1 : u32) -> u32 -struct atomic_compare_exchange_weak_ret_type_1 { - old_value : u32, - exchanged : bool, -} - @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body) -fn sbatomicCompareExchangeWeak_1(offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 +fn sbatomicCompareExchangeWeak_1(offset : u32, param_1 : u32, param_2 : u32) -> __atomic_compare_exchange_result_u32 @compute @workgroup_size(1) fn main() { @@ -3787,13 +3777,8 @@ @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body) fn sbatomicExchange(offset : u32, param_1 : i32) -> i32 -struct atomic_compare_exchange_weak_ret_type { - old_value : i32, - exchanged : bool, -} - @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body) -fn sbatomicCompareExchangeWeak(offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type +fn sbatomicCompareExchangeWeak(offset : u32, param_1 : i32, param_2 : i32) -> __atomic_compare_exchange_result_i32 @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body) fn sbatomicStore_1(offset : u32, param_1 : u32) @@ -3825,13 +3810,8 @@ @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body) fn sbatomicExchange_1(offset : u32, param_1 : u32) -> u32 -struct atomic_compare_exchange_weak_ret_type_1 { - old_value : u32, - exchanged : bool, -} - @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body) -fn sbatomicCompareExchangeWeak_1(offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 +fn sbatomicCompareExchangeWeak_1(offset : u32, param_1 : u32, param_2 : u32) -> __atomic_compare_exchange_result_u32 @compute @workgroup_size(1) fn main() {
diff --git a/src/tint/ast/transform/promote_side_effects_to_decl.cc b/src/tint/ast/transform/promote_side_effects_to_decl.cc index 19136dd..fed8327 100644 --- a/src/tint/ast/transform/promote_side_effects_to_decl.cc +++ b/src/tint/ast/transform/promote_side_effects_to_decl.cc
@@ -406,7 +406,8 @@ auto clone_maybe_hoisted = [&](const Expression* e) -> const Expression* { if (to_hoist.count(e)) { auto name = b.Symbols().New(); - auto* v = b.Let(name, ctx.Clone(e)); + auto* ty = sem.GetVal(e)->Type(); + auto* v = b.Let(name, Transform::CreateASTTypeFor(ctx, ty), ctx.Clone(e)); auto* decl = b.Decl(v); curr_stmts->Push(decl); return b.Expr(name);
diff --git a/src/tint/ast/transform/promote_side_effects_to_decl_test.cc b/src/tint/ast/transform/promote_side_effects_to_decl_test.cc index f8ce860..ec54d24 100644 --- a/src/tint/ast/transform/promote_side_effects_to_decl_test.cc +++ b/src/tint/ast/transform/promote_side_effects_to_decl_test.cc
@@ -25,8 +25,7 @@ auto* src = ""; auto* expect = ""; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -44,8 +43,7 @@ auto* expect = src; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -75,14 +73,13 @@ } fn f() { - let tint_symbol = a(); - let tint_symbol_1 = b(); + let tint_symbol : i32 = a(); + let tint_symbol_1 : i32 = b(); let r = (tint_symbol + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -106,13 +103,12 @@ fn f() { var b = 1; - let tint_symbol = a(); + let tint_symbol : i32 = a(); let r = (tint_symbol + b); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -136,14 +132,13 @@ fn f() { var b = 1; - let tint_symbol = b; - let tint_symbol_1 = a(); + let tint_symbol : i32 = b; + let tint_symbol_1 : i32 = a(); let r = (tint_symbol + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -171,13 +166,12 @@ var b = 1; var c = 1; var d = 1; - let tint_symbol = a(); + let tint_symbol : i32 = a(); let r = (((tint_symbol + b) + c) + d); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -205,14 +199,13 @@ var b = 1; var c = 1; var d = 1; - let tint_symbol = ((b + c) + d); - let tint_symbol_1 = a(); + let tint_symbol : i32 = ((b + c) + d); + let tint_symbol_1 : i32 = a(); let r = (tint_symbol + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -242,14 +235,13 @@ var c = 1; var d = 1; var e = 1; - let tint_symbol = (b + c); - let tint_symbol_1 = a(); + let tint_symbol : i32 = (b + c); + let tint_symbol_1 : i32 = a(); let r = (((tint_symbol + tint_symbol_1) + d) + e); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -271,15 +263,14 @@ } fn f() { - let tint_symbol = a(0); - let tint_symbol_1 = a(1); - let tint_symbol_2 = a(2); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); + let tint_symbol_2 : i32 = a(2); let r = ((tint_symbol + tint_symbol_1) + tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -301,13 +292,12 @@ } fn f() { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = ((((1 + tint_symbol) - 2) + 3) - 4); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -331,15 +321,14 @@ fn f() { var b = 1; - let tint_symbol = a(0); - let tint_symbol_1 = b; - let tint_symbol_2 = a(1); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = b; + let tint_symbol_2 : i32 = a(1); let r = ((((tint_symbol + 1) + tint_symbol_1) + tint_symbol_2) + 2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -365,13 +354,12 @@ fn main() { var b = 1; var c = 1; - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = ((1 + tint_symbol) + b); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -395,14 +383,13 @@ fn main() { var b = 1; - let tint_symbol = b; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = b; + let tint_symbol_1 : i32 = a(0); let r = ((tint_symbol + tint_symbol_1) + 1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -428,13 +415,12 @@ fn main() { var b = 1; var c = 1; - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = (((tint_symbol + b) + 1) + c); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -462,13 +448,12 @@ fn f() { var b = 0; - let tint_symbol = atomicAdd(&(sb.a), 123); + let tint_symbol : i32 = atomicAdd(&(sb.a), 123); let r = (tint_symbol + b); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -489,8 +474,7 @@ auto* expect = src; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -526,14 +510,13 @@ fn f() { var b = 0; - let tint_symbol = atomicLoad(&(sb.a)); - let tint_symbol_1 = a(0); + let tint_symbol : i32 = atomicLoad(&(sb.a)); + let tint_symbol_1 : i32 = a(0); let r = (tint_symbol + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -559,15 +542,14 @@ fn f() { var b : vec3<i32>; var c : i32; - let tint_symbol = c; - let tint_symbol_1 = b[tint_symbol]; - let tint_symbol_2 = a(); + let tint_symbol : i32 = c; + let tint_symbol_1 : i32 = b[tint_symbol]; + let tint_symbol_2 : i32 = a(); let r = (tint_symbol_1 + tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -599,17 +581,16 @@ fn f() { var b = 1; - let tint_symbol = a(0); - let tint_symbol_1 = g(tint_symbol); - let tint_symbol_2 = b; - let tint_symbol_3 = a(1); - let tint_symbol_4 = g((tint_symbol_2 + tint_symbol_3)); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = g(tint_symbol); + let tint_symbol_2 : i32 = b; + let tint_symbol_3 : i32 = a(1); + let tint_symbol_4 : i32 = g((tint_symbol_2 + tint_symbol_3)); let r = (tint_symbol_1 - tint_symbol_4); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -634,19 +615,18 @@ fn f() { var b = 1; - let tint_symbol = a(0); - let tint_symbol_1 = i32(tint_symbol); - let tint_symbol_2 = a(1); - let tint_symbol_3 = i32((tint_symbol_2 + b)); - let tint_symbol_4 = a(2); - let tint_symbol_5 = a(3); - let tint_symbol_6 = i32((tint_symbol_4 - tint_symbol_5)); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = i32(tint_symbol); + let tint_symbol_2 : i32 = a(1); + let tint_symbol_3 : i32 = i32((tint_symbol_2 + b)); + let tint_symbol_4 : i32 = a(2); + let tint_symbol_5 : i32 = a(3); + let tint_symbol_6 : i32 = i32((tint_symbol_4 - tint_symbol_5)); let r = ((tint_symbol_1 + tint_symbol_3) - tint_symbol_6); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -671,16 +651,15 @@ fn f() { var b = 1u; - let tint_symbol = a(0); - let tint_symbol_1 = u32(tint_symbol); - let tint_symbol_2 = a(1); - let tint_symbol_3 = u32(tint_symbol_2); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : u32 = u32(tint_symbol); + let tint_symbol_2 : i32 = a(1); + let tint_symbol_3 : u32 = u32(tint_symbol_2); let r = ((tint_symbol_1 + tint_symbol_3) - b); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -705,19 +684,18 @@ fn f() { var b = 1; - let tint_symbol = a(0); - let tint_symbol_1 = abs(tint_symbol); - let tint_symbol_2 = a(1); - let tint_symbol_3 = abs((tint_symbol_2 + b)); - let tint_symbol_4 = a(2); - let tint_symbol_5 = a(3); - let tint_symbol_6 = abs((tint_symbol_4 + tint_symbol_5)); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = abs(tint_symbol); + let tint_symbol_2 : i32 = a(1); + let tint_symbol_3 : i32 = abs((tint_symbol_2 + b)); + let tint_symbol_4 : i32 = a(2); + let tint_symbol_5 : i32 = a(3); + let tint_symbol_6 : i32 = abs((tint_symbol_4 + tint_symbol_5)); let r = ((tint_symbol_1 + tint_symbol_3) - tint_symbol_6); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -750,15 +728,14 @@ fn f() { var b = 1; - let tint_symbol = a(0); - let tint_symbol_1 = b; - let tint_symbol_2 = a(1); + let tint_symbol : S = a(0); + let tint_symbol_1 : i32 = b; + let tint_symbol_2 : S = a(1); let r = ((tint_symbol.v + tint_symbol_1) + tint_symbol_2.v); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -782,15 +759,14 @@ fn f() { var b = 1; - let tint_symbol = -(a(0)); - let tint_symbol_1 = b; - let tint_symbol_2 = a(1); + let tint_symbol : i32 = -(a(0)); + let tint_symbol_1 : i32 = b; + let tint_symbol_2 : i32 = a(1); let r = (tint_symbol + -((tint_symbol_1 + tint_symbol_2))); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -814,14 +790,13 @@ fn f() { var b = 1; - let tint_symbol = a(0); - let tint_symbol_1 = a(1); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); let r = bitcast<u32>((tint_symbol + tint_symbol_1)); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -849,7 +824,7 @@ fn f() { var b = 1; { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); var r = (tint_symbol + b); loop { { @@ -861,8 +836,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -889,7 +863,7 @@ fn f() { var b = 1; loop { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); if (!(((tint_symbol + b) > 0))) { break; } @@ -900,8 +874,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -937,15 +910,14 @@ } continuing { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); r = (tint_symbol + b); } } } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -978,10 +950,10 @@ var d = 3; var r = 0; { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); var r = (tint_symbol + b); loop { - let tint_symbol_1 = a(1); + let tint_symbol_1 : i32 = a(1); if (!(((tint_symbol_1 + c) > 0))) { break; } @@ -990,7 +962,7 @@ } continuing { - let tint_symbol_2 = a(2); + let tint_symbol_2 : i32 = a(2); r = (tint_symbol_2 + d); } } @@ -998,8 +970,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1026,7 +997,7 @@ fn f() { var b = 1; loop { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); if (!(((tint_symbol + b) > 0))) { break; } @@ -1037,8 +1008,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1069,7 +1039,7 @@ if (true) { var marker = 0; } else { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); if (((tint_symbol + b) > 0)) { var marker = 1; } @@ -1077,8 +1047,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1119,12 +1088,12 @@ } else if (true) { var marker = 1; } else { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); if (((tint_symbol + b) > 0)) { var marker = 2; } else { - let tint_symbol_1 = a(1); - let tint_symbol_2 = a(2); + let tint_symbol_1 : i32 = a(1); + let tint_symbol_2 : i32 = a(2); if (((tint_symbol_1 + tint_symbol_2) > 0)) { var marker = 3; } else if (true) { @@ -1137,8 +1106,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1162,14 +1130,13 @@ fn f() -> i32 { var b = 1; - let tint_symbol = b; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = b; + let tint_symbol_1 : i32 = a(0); return (tint_symbol + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1196,8 +1163,8 @@ fn f() { var b = 1; - let tint_symbol = b; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = b; + let tint_symbol_1 : i32 = a(0); switch((tint_symbol + tint_symbol_1)) { default: { } @@ -1205,8 +1172,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1238,8 +1204,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1271,8 +1236,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1302,8 +1266,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1347,8 +1310,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1384,8 +1346,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1429,8 +1390,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1468,8 +1428,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1509,8 +1468,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1548,8 +1506,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1585,8 +1542,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1628,8 +1584,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1685,8 +1640,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1732,8 +1686,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1779,8 +1732,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1818,8 +1770,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1859,8 +1810,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1885,7 +1835,7 @@ fn f() { var b = true; - let tint_symbol_2 = a(0); + let tint_symbol_2 : bool = a(0); var tint_symbol_1 = bool(tint_symbol_2); if (tint_symbol_1) { var tint_symbol_3 = a(1); @@ -1906,8 +1856,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1932,10 +1881,10 @@ fn f() { var b = true; - let tint_symbol_2 = a(0); + let tint_symbol_2 : i32 = a(0); var tint_symbol_1 = bool(tint_symbol_2); if (tint_symbol_1) { - let tint_symbol_3 = a(1); + let tint_symbol_3 : i32 = a(1); tint_symbol_1 = bool(tint_symbol_3); } var tint_symbol = tint_symbol_1; @@ -1946,8 +1895,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -1974,18 +1922,17 @@ fn f() { var b = 1; - let tint_symbol_1 = a(0); + let tint_symbol_1 : i32 = a(0); var tint_symbol = bool((tint_symbol_1 == b)); if (tint_symbol) { - let tint_symbol_2 = a(1); + let tint_symbol_2 : i32 = a(1); tint_symbol = bool((tint_symbol_2 == b)); } let r = tint_symbol; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2010,7 +1957,7 @@ fn f() { var b = true; - let tint_symbol_2 = a(0); + let tint_symbol_2 : bool = a(0); var tint_symbol_1 = all(tint_symbol_2); if (tint_symbol_1) { var tint_symbol_3 = a(1); @@ -2031,8 +1978,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2065,22 +2011,21 @@ fn f() { var b = true; - let tint_symbol_2 = a(0); + let tint_symbol_2 : S = a(0); var tint_symbol_1 = tint_symbol_2.v; if (tint_symbol_1) { tint_symbol_1 = b; } var tint_symbol = tint_symbol_1; if (tint_symbol) { - let tint_symbol_3 = a(1); + let tint_symbol_3 : S = a(1); tint_symbol = tint_symbol_3.v; } let r = tint_symbol; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2113,8 +2058,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2146,8 +2090,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2190,8 +2133,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2232,8 +2174,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2279,8 +2220,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2342,8 +2282,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2384,8 +2323,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2427,8 +2365,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2492,8 +2429,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2517,8 +2453,7 @@ auto* expect = src; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2550,13 +2485,12 @@ fn f() { var b = 1; - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = g(tint_symbol, b, 3); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2586,15 +2520,14 @@ } fn f() { - let tint_symbol = a(0); - let tint_symbol_1 = a(1); - let tint_symbol_2 = a(2); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); + let tint_symbol_2 : i32 = a(2); let r = g(tint_symbol, tint_symbol_1, tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2627,15 +2560,14 @@ fn f() { var b = 1; - let tint_symbol = a(0); - let tint_symbol_1 = b; - let tint_symbol_2 = a(1); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = b; + let tint_symbol_2 : i32 = a(1); let r = g(tint_symbol, tint_symbol_1, tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2671,17 +2603,16 @@ var b = 0; var c = 0; var d = 0; - let tint_symbol = b; - let tint_symbol_1 = c; - let tint_symbol_2 = a(0); - let tint_symbol_3 = g(tint_symbol_1, tint_symbol_2, d); - let tint_symbol_4 = a(1); + let tint_symbol : i32 = b; + let tint_symbol_1 : i32 = c; + let tint_symbol_2 : i32 = a(0); + let tint_symbol_3 : i32 = g(tint_symbol_1, tint_symbol_2, d); + let tint_symbol_4 : i32 = a(1); let r = ((tint_symbol + tint_symbol_3) + tint_symbol_4); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2707,13 +2638,12 @@ fn f() { var b = array<array<i32, 10>, 10>(); var c = 1; - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); var r = b[tint_symbol][c]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2747,8 +2677,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2772,14 +2701,13 @@ fn f() { var b = array<array<i32, 10>, 10>(); - let tint_symbol = a(0); - let tint_symbol_1 = a(1); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); var r = b[tint_symbol][tint_symbol_1]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2805,8 +2733,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2830,13 +2757,12 @@ fn f() { var b = array<i32, 10>(); - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); b[tint_symbol] = a(1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2860,14 +2786,13 @@ fn f() { var b = array<array<i32, 10>, 10>(); - let tint_symbol = a(0); - let tint_symbol_1 = a(1); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); b[tint_symbol][tint_symbol_1] = a(2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2891,15 +2816,14 @@ fn f() { var b = array<array<array<i32, 10>, 10>, 10>(); - let tint_symbol = a(0); - let tint_symbol_1 = a(1); - let tint_symbol_2 = a(2); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); + let tint_symbol_2 : i32 = a(2); b[tint_symbol][tint_symbol_1][tint_symbol_2] = a(3); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2927,14 +2851,13 @@ var b = array<i32, 3>(); var d = array<array<i32, 3>, 3>(); var a_1 = 0; - let tint_symbol = a(2); - let tint_symbol_1 = a(0); + let tint_symbol : i32 = a(2); + let tint_symbol_1 : i32 = a(0); b[tint_symbol] = d[tint_symbol_1][a_1]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2958,13 +2881,12 @@ fn f() { var b = vec3<i32>(); - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); b[tint_symbol] = a(1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -2990,13 +2912,12 @@ fn f() { var b = vec3<i32>(); var c = 0; - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); b[tint_symbol] = c; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3022,13 +2943,12 @@ fn f() { var b = vec3<i32>(); var c = 0; - let tint_symbol = c; + let tint_symbol : i32 = c; b[tint_symbol] = a(0); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3062,15 +2982,14 @@ } fn f() { - let tint_symbol = a(0); - let tint_symbol_1 = a(1); - let tint_symbol_2 = a(2); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); + let tint_symbol_2 : i32 = a(2); var r = S(tint_symbol, tint_symbol_1, tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3092,15 +3011,14 @@ } fn f() { - let tint_symbol = a(0); - let tint_symbol_1 = a(1); - let tint_symbol_2 = a(2); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); + let tint_symbol_2 : i32 = a(2); var r = array<i32, 3>(tint_symbol, tint_symbol_1, tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3122,18 +3040,17 @@ } fn f() { - let tint_symbol = a(0); - let tint_symbol_1 = a(1); - let tint_symbol_2 = array<i32, 2>(tint_symbol, tint_symbol_1); - let tint_symbol_3 = a(2); - let tint_symbol_4 = a(3); - let tint_symbol_5 = array<i32, 2>(tint_symbol_3, tint_symbol_4); + let tint_symbol : i32 = a(0); + let tint_symbol_1 : i32 = a(1); + let tint_symbol_2 : array<i32, 2u> = array<i32, 2>(tint_symbol, tint_symbol_1); + let tint_symbol_3 : i32 = a(2); + let tint_symbol_4 : i32 = a(3); + let tint_symbol_5 : array<i32, 2u> = array<i32, 2>(tint_symbol_3, tint_symbol_4); var r = array<array<i32, 2>, 2>(tint_symbol_2, tint_symbol_5); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3155,14 +3072,13 @@ } fn f() { - let tint_symbol = a(0); - let tint_symbol_1 = a(1); + let tint_symbol : vec3<i32> = a(0); + let tint_symbol_1 : vec3<i32> = a(1); var r = (tint_symbol.x + tint_symbol_1.y); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3194,14 +3110,13 @@ } fn f() { - let tint_symbol = a(0); - let tint_symbol_1 = a(1); + let tint_symbol : S = a(0); + let tint_symbol_1 : S = a(1); var r = (tint_symbol.x + tint_symbol_1.y); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3253,19 +3168,18 @@ var k = 0; var l = 0; var m = 0; - let tint_symbol = a(0); - let tint_symbol_1 = i; - let tint_symbol_2 = a(1); - let tint_symbol_3 = j; - let tint_symbol_4 = a(2); - let tint_symbol_5 = k; - let tint_symbol_6 = b(3); + let tint_symbol : S = a(0); + let tint_symbol_1 : i32 = i; + let tint_symbol_2 : S = a(1); + let tint_symbol_3 : i32 = j; + let tint_symbol_4 : S = a(2); + let tint_symbol_5 : i32 = k; + let tint_symbol_6 : i32 = b(3); var r = (((((tint_symbol.x + tint_symbol_1) + tint_symbol_2.y) + tint_symbol_3) + tint_symbol_4.arr[((tint_symbol_5 + tint_symbol_6) + l)]) + m); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3289,14 +3203,13 @@ fn f() { var v = array<i32, 10>(); - let tint_symbol = v[0]; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = v[0]; + let tint_symbol_1 : i32 = a(0); let r = (tint_symbol + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3320,13 +3233,12 @@ fn f() { var v = array<i32, 10>(); - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = v[tint_symbol]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3350,13 +3262,12 @@ fn f() { var v = array<array<i32, 10>, 10>(); - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = v[tint_symbol][0]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3380,13 +3291,12 @@ fn f() { var v = array<array<i32, 10>, 10>(); - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = v[0][tint_symbol]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3412,13 +3322,12 @@ fn f() { var v = array<array<i32, 10>, 10>(); var b : i32; - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = v[tint_symbol][b]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3444,14 +3353,13 @@ fn f() { var v = array<array<i32, 10>, 10>(); var b : i32; - let tint_symbol = b; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = b; + let tint_symbol_1 : i32 = a(0); let r = v[tint_symbol][tint_symbol_1]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3477,15 +3385,14 @@ fn f() { var v = array<i32, 10>(); var b = 0; - let tint_symbol = b; - let tint_symbol_1 = v[tint_symbol]; - let tint_symbol_2 = a(0); + let tint_symbol : i32 = b; + let tint_symbol_1 : i32 = v[tint_symbol]; + let tint_symbol_2 : i32 = a(0); let r = (tint_symbol_1 + tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3509,14 +3416,13 @@ fn f() { var v = array<i32, 10>(); - let tint_symbol = v[0]; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = v[0]; + let tint_symbol_1 : i32 = a(0); let r = (tint_symbol + v[tint_symbol_1]); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3542,13 +3448,12 @@ fn f() { var v = array<i32, 10>(); var w = array<i32, 10>(); - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); v[w[tint_symbol]] = 1; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3574,14 +3479,13 @@ fn f() { var v = array<i32, 10>(); var w = array<i32, 10>(); - let tint_symbol = w[0]; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = w[0]; + let tint_symbol_1 : i32 = a(0); v[(tint_symbol + tint_symbol_1)] = 1; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3608,14 +3512,13 @@ fn f() { var v = array<i32, 10>(); var w = array<i32, 10>(); - let tint_symbol = w[0]; - let tint_symbol_1 = a(0); + let tint_symbol : i32 = w[0]; + let tint_symbol_1 : i32 = a(0); v[(tint_symbol + w[tint_symbol_1])] = 1; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3645,14 +3548,13 @@ } fn f() { - let tint_symbol = b(); - let tint_symbol_1 = a(0); + let tint_symbol : array<i32, 10u> = b(); + let tint_symbol_1 : i32 = a(0); let r = tint_symbol[tint_symbol_1]; } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3682,14 +3584,13 @@ } fn f() { - let tint_symbol = b(); - let tint_symbol_1 = a(0); + let tint_symbol : array<i32, 10u> = b(); + let tint_symbol_1 : i32 = a(0); let r = (tint_symbol[0] + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3718,14 +3619,13 @@ fn f() { var v = vec4<i32>(); - let tint_symbol = v.x; - let tint_symbol_1 = modify_vec(&(v)); + let tint_symbol : i32 = v.x; + let tint_symbol_1 : i32 = modify_vec(&(v)); let l = (tint_symbol + tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3755,13 +3655,12 @@ } fn f() { - let tint_symbol = get_uv(); + let tint_symbol : vec2<f32> = get_uv(); let r = textureGather(1, tex, samp, tint_symbol, 1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3801,13 +3700,12 @@ } fn f() { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); let r = g(&(b), tint_symbol); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3824,8 +3722,7 @@ auto* expect = src; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3863,18 +3760,17 @@ } fn f() { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); var tint_symbol_1 = b(1); if (tint_symbol_1) { tint_symbol_1 = b(2); } - let tint_symbol_2 = g(tint_symbol_1); + let tint_symbol_2 : i32 = g(tint_symbol_1); let r = (tint_symbol + tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3916,14 +3812,13 @@ if (tint_symbol) { tint_symbol = b(1); } - let tint_symbol_1 = g(tint_symbol); - let tint_symbol_2 = a(2); + let tint_symbol_1 : i32 = g(tint_symbol); + let tint_symbol_2 : i32 = a(2); let r = (tint_symbol_1 + tint_symbol_2); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -3961,7 +3856,7 @@ } fn f() { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); var tint_symbol_1 = b(1); if (tint_symbol_1) { tint_symbol_1 = b(2); @@ -3970,8 +3865,7 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -4013,13 +3907,12 @@ if (tint_symbol) { tint_symbol = b(1); } - let tint_symbol_1 = a(2); + let tint_symbol_1 : i32 = a(2); let r = g(tint_symbol, tint_symbol_1); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -4065,22 +3958,21 @@ } fn f() { - let tint_symbol = a(0); + let tint_symbol : i32 = a(0); var tint_symbol_1 = b(1); if (tint_symbol_1) { - let tint_symbol_2 = a(2); - let tint_symbol_3 = a(3); + let tint_symbol_2 : i32 = a(2); + let tint_symbol_3 : i32 = a(3); tint_symbol_1 = b((tint_symbol_2 + tint_symbol_3)); } - let tint_symbol_4 = a(4); - let tint_symbol_5 = g(tint_symbol_1, tint_symbol_4); - let tint_symbol_6 = a(5); + let tint_symbol_4 : i32 = a(4); + let tint_symbol_5 : i32 = g(tint_symbol_1, tint_symbol_4); + let tint_symbol_6 : i32 = a(5); let r = ((tint_symbol + tint_symbol_5) + tint_symbol_6); } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); } @@ -4118,7 +4010,7 @@ } fn f(t : texture_2d<f32>, s : sampler) -> vec4<f32> { - let tint_symbol = side_effects(); + let tint_symbol : vec2<f32> = side_effects(); return textureSample(t, s, tint_symbol); } @@ -4127,8 +4019,79 @@ } )"; - Transform::DataMap data; - auto got = Run<PromoteSideEffectsToDecl>(src, data); + auto got = Run<PromoteSideEffectsToDecl>(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PromoteSideEffectsToDeclTest, BuiltinReturnType) { + auto* src = R"( +fn X(a:vec2f, b:vec2f) { +} + +fn Y() -> vec2f { return vec2f(); } + +fn f() { + var v: vec2f; + X(vec2(), v); // okay + X(vec2(), Y()); // errors +} +)"; + + auto* expect = R"( +fn X(a : vec2f, b : vec2f) { +} + +fn Y() -> vec2f { + return vec2f(); +} + +fn f() { + var v : vec2f; + X(vec2(), v); + let tint_symbol : vec2<f32> = vec2(); + let tint_symbol_1 : vec2<f32> = Y(); + X(tint_symbol, tint_symbol_1); +} +)"; + + auto got = Run<PromoteSideEffectsToDecl>(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PromoteSideEffectsToDeclTest, Bug1963) { + auto* src = R"( +fn X(a:vec2f, b:vec2f) { +} + +fn Y() -> vec2f { return vec2f(); } + +fn f() { + var v: vec2f; + X(vec2(), v); // okay + X(vec2(), Y()); // errors +} +)"; + + auto* expect = R"( +fn X(a : vec2f, b : vec2f) { +} + +fn Y() -> vec2f { + return vec2f(); +} + +fn f() { + var v : vec2f; + X(vec2(), v); + let tint_symbol : vec2<f32> = vec2(); + let tint_symbol_1 : vec2<f32> = Y(); + X(tint_symbol, tint_symbol_1); +} +)"; + + auto got = Run<PromoteSideEffectsToDecl>(src); EXPECT_EQ(expect, str(got)); }
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc index 1b30d8d..def946a 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1738,6 +1738,10 @@ return true; } case Op::kAtomicCompareExchangeWeak: { + if (!EmitStructType(&helpers_, result_ty->As<type::Struct>())) { + return false; + } + auto* const value_ty = sem_func->Parameters()[1]->Type()->UnwrapRef(); // NOTE: We don't need to emit the return type struct here as DecomposeMemoryAccess // already added it to the AST, and it should have already been emitted by now.
diff --git a/test/tint/bug/tint/1573.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1573.wgsl.expected.dxc.hlsl index aa56552..31f56b4 100644 --- a/test/tint/bug/tint/1573.wgsl.expected.dxc.hlsl +++ b/test/tint/bug/tint/1573.wgsl.expected.dxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer a : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; +RWByteAddressBuffer a : register(u0); -atomic_compare_exchange_weak_ret_type aatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 aatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; a.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -16,6 +15,6 @@ [numthreads(16, 1, 1)] void main() { uint value = 42u; - const atomic_compare_exchange_weak_ret_type result = aatomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 result = aatomicCompareExchangeWeak(0u, 0u, value); return; }
diff --git a/test/tint/bug/tint/1573.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1573.wgsl.expected.fxc.hlsl index aa56552..31f56b4 100644 --- a/test/tint/bug/tint/1573.wgsl.expected.fxc.hlsl +++ b/test/tint/bug/tint/1573.wgsl.expected.fxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer a : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; +RWByteAddressBuffer a : register(u0); -atomic_compare_exchange_weak_ret_type aatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 aatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; a.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -16,6 +15,6 @@ [numthreads(16, 1, 1)] void main() { uint value = 42u; - const atomic_compare_exchange_weak_ret_type result = aatomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 result = aatomicCompareExchangeWeak(0u, 0u, value); return; }
diff --git a/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl index 8a4a1ba..1e5c723 100644 --- a/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl +++ b/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl
@@ -14,26 +14,17 @@ struct tint_symbol_1 { uint local_invocation_index : SV_GroupIndex; }; -struct atomic_compare_exchange_weak_ret_type { - uint old_value; - bool exchanged; -}; -atomic_compare_exchange_weak_ret_type a_u32atomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 a_u32atomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; a_u32.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; } -struct atomic_compare_exchange_weak_ret_type_1 { - int old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type_1 a_i32atomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type_1 result=(atomic_compare_exchange_weak_ret_type_1)0; +atomic_compare_exchange_result_i32 a_i32atomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; a_i32.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -50,15 +41,15 @@ GroupMemoryBarrierWithGroupSync(); { uint value = 42u; - const atomic_compare_exchange_weak_ret_type r1 = a_u32atomicCompareExchangeWeak(0u, 0u, value); - const atomic_compare_exchange_weak_ret_type r2 = a_u32atomicCompareExchangeWeak(0u, 0u, value); - const atomic_compare_exchange_weak_ret_type r3 = a_u32atomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 r1 = a_u32atomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 r2 = a_u32atomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 r3 = a_u32atomicCompareExchangeWeak(0u, 0u, value); } { int value = 42; - const atomic_compare_exchange_weak_ret_type_1 r1 = a_i32atomicCompareExchangeWeak(0u, 0, value); - const atomic_compare_exchange_weak_ret_type_1 r2 = a_i32atomicCompareExchangeWeak(0u, 0, value); - const atomic_compare_exchange_weak_ret_type_1 r3 = a_i32atomicCompareExchangeWeak(0u, 0, value); + const atomic_compare_exchange_result_i32 r1 = a_i32atomicCompareExchangeWeak(0u, 0, value); + const atomic_compare_exchange_result_i32 r2 = a_i32atomicCompareExchangeWeak(0u, 0, value); + const atomic_compare_exchange_result_i32 r3 = a_i32atomicCompareExchangeWeak(0u, 0, value); } { uint value = 42u;
diff --git a/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl index 8a4a1ba..1e5c723 100644 --- a/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl +++ b/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl
@@ -14,26 +14,17 @@ struct tint_symbol_1 { uint local_invocation_index : SV_GroupIndex; }; -struct atomic_compare_exchange_weak_ret_type { - uint old_value; - bool exchanged; -}; -atomic_compare_exchange_weak_ret_type a_u32atomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 a_u32atomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; a_u32.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; } -struct atomic_compare_exchange_weak_ret_type_1 { - int old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type_1 a_i32atomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type_1 result=(atomic_compare_exchange_weak_ret_type_1)0; +atomic_compare_exchange_result_i32 a_i32atomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; a_i32.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -50,15 +41,15 @@ GroupMemoryBarrierWithGroupSync(); { uint value = 42u; - const atomic_compare_exchange_weak_ret_type r1 = a_u32atomicCompareExchangeWeak(0u, 0u, value); - const atomic_compare_exchange_weak_ret_type r2 = a_u32atomicCompareExchangeWeak(0u, 0u, value); - const atomic_compare_exchange_weak_ret_type r3 = a_u32atomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 r1 = a_u32atomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 r2 = a_u32atomicCompareExchangeWeak(0u, 0u, value); + const atomic_compare_exchange_result_u32 r3 = a_u32atomicCompareExchangeWeak(0u, 0u, value); } { int value = 42; - const atomic_compare_exchange_weak_ret_type_1 r1 = a_i32atomicCompareExchangeWeak(0u, 0, value); - const atomic_compare_exchange_weak_ret_type_1 r2 = a_i32atomicCompareExchangeWeak(0u, 0, value); - const atomic_compare_exchange_weak_ret_type_1 r3 = a_i32atomicCompareExchangeWeak(0u, 0, value); + const atomic_compare_exchange_result_i32 r1 = a_i32atomicCompareExchangeWeak(0u, 0, value); + const atomic_compare_exchange_result_i32 r2 = a_i32atomicCompareExchangeWeak(0u, 0, value); + const atomic_compare_exchange_result_i32 r3 = a_i32atomicCompareExchangeWeak(0u, 0, value); } { uint value = 42u;
diff --git a/test/tint/bug/tint/1963_a.wgsl b/test/tint/bug/tint/1963_a.wgsl new file mode 100644 index 0000000..90b4ae8 --- /dev/null +++ b/test/tint/bug/tint/1963_a.wgsl
@@ -0,0 +1,10 @@ +fn X(a:vec2f, b:vec2f) { +} + +fn Y() -> vec2f { return vec2f(); } + +fn f() { + var v: vec2f; + X(vec2(), v); // okay + X(vec2(), Y()); // errors +}
diff --git a/test/tint/bug/tint/1963_a.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1963_a.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000..e1213df --- /dev/null +++ b/test/tint/bug/tint/1963_a.wgsl.expected.dxc.hlsl
@@ -0,0 +1,19 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void X(float2 a, float2 b) { +} + +float2 Y() { + return (0.0f).xx; +} + +void f() { + float2 v = float2(0.0f, 0.0f); + X((0.0f).xx, v); + const float2 tint_symbol = (0.0f).xx; + const float2 tint_symbol_1 = Y(); + X(tint_symbol, tint_symbol_1); +}
diff --git a/test/tint/bug/tint/1963_a.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1963_a.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000..e1213df --- /dev/null +++ b/test/tint/bug/tint/1963_a.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void X(float2 a, float2 b) { +} + +float2 Y() { + return (0.0f).xx; +} + +void f() { + float2 v = float2(0.0f, 0.0f); + X((0.0f).xx, v); + const float2 tint_symbol = (0.0f).xx; + const float2 tint_symbol_1 = Y(); + X(tint_symbol, tint_symbol_1); +}
diff --git a/test/tint/bug/tint/1963_a.wgsl.expected.glsl b/test/tint/bug/tint/1963_a.wgsl.expected.glsl new file mode 100644 index 0000000..8ac552f --- /dev/null +++ b/test/tint/bug/tint/1963_a.wgsl.expected.glsl
@@ -0,0 +1,21 @@ +#version 310 es + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void unused_entry_point() { + return; +} +void X(vec2 a, vec2 b) { +} + +vec2 Y() { + return vec2(0.0f); +} + +void f() { + vec2 v = vec2(0.0f, 0.0f); + X(vec2(0.0f), v); + vec2 tint_symbol = vec2(0.0f); + vec2 tint_symbol_1 = Y(); + X(tint_symbol, tint_symbol_1); +} +
diff --git a/test/tint/bug/tint/1963_a.wgsl.expected.msl b/test/tint/bug/tint/1963_a.wgsl.expected.msl new file mode 100644 index 0000000..7158f2f --- /dev/null +++ b/test/tint/bug/tint/1963_a.wgsl.expected.msl
@@ -0,0 +1,18 @@ +#include <metal_stdlib> + +using namespace metal; +void X(float2 a, float2 b) { +} + +float2 Y() { + return float2(0.0f); +} + +void f() { + float2 v = 0.0f; + X(float2(0.0f), v); + float2 const tint_symbol = float2(0.0f); + float2 const tint_symbol_1 = Y(); + X(tint_symbol, tint_symbol_1); +} +
diff --git a/test/tint/bug/tint/1963_a.wgsl.expected.spvasm b/test/tint/bug/tint/1963_a.wgsl.expected.spvasm new file mode 100644 index 0000000..610e6ec --- /dev/null +++ b/test/tint/bug/tint/1963_a.wgsl.expected.spvasm
@@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %X "X" + OpName %a "a" + OpName %b "b" + OpName %Y "Y" + OpName %f "f" + OpName %v "v" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 + %5 = OpTypeFunction %void %v2float %v2float + %12 = OpTypeFunction %v2float + %15 = OpConstantNull %v2float +%_ptr_Function_v2float = OpTypePointer Function %v2float +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %X = OpFunction %void None %5 + %a = OpFunctionParameter %v2float + %b = OpFunctionParameter %v2float + %11 = OpLabel + OpReturn + OpFunctionEnd + %Y = OpFunction %v2float None %12 + %14 = OpLabel + OpReturnValue %15 + OpFunctionEnd + %f = OpFunction %void None %1 + %17 = OpLabel + %v = OpVariable %_ptr_Function_v2float Function %15 + %21 = OpLoad %v2float %v + %20 = OpFunctionCall %void %X %15 %21 + %22 = OpFunctionCall %v2float %Y + %23 = OpFunctionCall %void %X %15 %22 + OpReturn + OpFunctionEnd
diff --git a/test/tint/bug/tint/1963_a.wgsl.expected.wgsl b/test/tint/bug/tint/1963_a.wgsl.expected.wgsl new file mode 100644 index 0000000..2941e83 --- /dev/null +++ b/test/tint/bug/tint/1963_a.wgsl.expected.wgsl
@@ -0,0 +1,12 @@ +fn X(a : vec2f, b : vec2f) { +} + +fn Y() -> vec2f { + return vec2f(); +} + +fn f() { + var v : vec2f; + X(vec2(), v); + X(vec2(), Y()); +}
diff --git a/test/tint/bug/tint/1963_b.wgsl b/test/tint/bug/tint/1963_b.wgsl new file mode 100644 index 0000000..29c0e25 --- /dev/null +++ b/test/tint/bug/tint/1963_b.wgsl
@@ -0,0 +1,13 @@ +@group(0) @binding(0) var<storage, read_write> a : atomic<i32>; + +@compute @workgroup_size(1) +fn compute_main() { + // Secondary bug encountered while fixing crbug.com/tint/1963 + // With the fix, PromoteSideEffectsToDecl will emit a type for 'v', which is the un-typable + // '__atomic_compare_exchange_result_i32' result. Later, DecomposeMemoryAccess would try to + // construct its own struct to assign the intrinsic result of atomicCompareExchangeWeak() to + // that let. The two types wouldn't match, and so this would fail to compile. The fix is to have + // DecomposeMemoryAccess use the internal '__atomic_compare_exchange_result_i32' result instead + // of building its own. + let v = atomicCompareExchangeWeak(&a, 1, 1).old_value; +}
diff --git a/test/tint/bug/tint/1963_b.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1963_b.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000..9c7b402 --- /dev/null +++ b/test/tint/bug/tint/1963_b.wgsl.expected.dxc.hlsl
@@ -0,0 +1,20 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; +RWByteAddressBuffer a : register(u0); + +atomic_compare_exchange_result_i32 aatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; + a.InterlockedCompareExchange(offset, compare, value, result.old_value); + result.exchanged = result.old_value == compare; + return result; +} + + +[numthreads(1, 1, 1)] +void compute_main() { + const atomic_compare_exchange_result_i32 tint_symbol = aatomicCompareExchangeWeak(0u, 1, 1); + const int v = tint_symbol.old_value; + return; +}
diff --git a/test/tint/bug/tint/1963_b.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1963_b.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000..9c7b402 --- /dev/null +++ b/test/tint/bug/tint/1963_b.wgsl.expected.fxc.hlsl
@@ -0,0 +1,20 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; +RWByteAddressBuffer a : register(u0); + +atomic_compare_exchange_result_i32 aatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; + a.InterlockedCompareExchange(offset, compare, value, result.old_value); + result.exchanged = result.old_value == compare; + return result; +} + + +[numthreads(1, 1, 1)] +void compute_main() { + const atomic_compare_exchange_result_i32 tint_symbol = aatomicCompareExchangeWeak(0u, 1, 1); + const int v = tint_symbol.old_value; + return; +}
diff --git a/test/tint/bug/tint/1963_b.wgsl.expected.glsl b/test/tint/bug/tint/1963_b.wgsl.expected.glsl new file mode 100644 index 0000000..9565468 --- /dev/null +++ b/test/tint/bug/tint/1963_b.wgsl.expected.glsl
@@ -0,0 +1,25 @@ +#version 310 es + +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; + + +layout(binding = 0, std430) buffer a_block_ssbo { + int inner; +} a; + +void compute_main() { + atomic_compare_exchange_result_i32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(a.inner, 1, 1); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; + atomic_compare_exchange_result_i32 tint_symbol = atomic_compare_result; + int v = tint_symbol.old_value; +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +}
diff --git a/test/tint/bug/tint/1963_b.wgsl.expected.msl b/test/tint/bug/tint/1963_b.wgsl.expected.msl new file mode 100644 index 0000000..d3fa9e7 --- /dev/null +++ b/test/tint/bug/tint/1963_b.wgsl.expected.msl
@@ -0,0 +1,20 @@ +#include <metal_stdlib> + +using namespace metal; + +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; +atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) { + int old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +kernel void compute_main(device atomic_int* tint_symbol_1 [[buffer(0)]]) { + atomic_compare_exchange_result_i32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_1, 1, 1); + int const v = tint_symbol.old_value; + return; +} +
diff --git a/test/tint/bug/tint/1963_b.wgsl.expected.spvasm b/test/tint/bug/tint/1963_b.wgsl.expected.spvasm new file mode 100644 index 0000000..ae05780 --- /dev/null +++ b/test/tint/bug/tint/1963_b.wgsl.expected.spvasm
@@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 22 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %a_block "a_block" + OpMemberName %a_block 0 "inner" + OpName %a "a" + OpName %compute_main "compute_main" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" + OpDecorate %a_block Block + OpMemberDecorate %a_block 0 Offset 0 + OpDecorate %a DescriptorSet 0 + OpDecorate %a Binding 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 + %int = OpTypeInt 32 1 + %a_block = OpTypeStruct %int +%_ptr_StorageBuffer_a_block = OpTypePointer StorageBuffer %a_block + %a = OpVariable %_ptr_StorageBuffer_a_block StorageBuffer + %void = OpTypeVoid + %5 = OpTypeFunction %void + %bool = OpTypeBool +%__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 +%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int + %int_1 = OpConstant %int 1 +%compute_main = OpFunction %void None %5 + %8 = OpLabel + %17 = OpAccessChain %_ptr_StorageBuffer_int %a %uint_0 + %19 = OpAtomicCompareExchange %int %17 %uint_1 %uint_0 %uint_0 %int_1 %int_1 + %20 = OpIEqual %bool %19 %int_1 + %9 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %19 %20 + %21 = OpCompositeExtract %int %9 0 + OpReturn + OpFunctionEnd
diff --git a/test/tint/bug/tint/1963_b.wgsl.expected.wgsl b/test/tint/bug/tint/1963_b.wgsl.expected.wgsl new file mode 100644 index 0000000..4955fd0 --- /dev/null +++ b/test/tint/bug/tint/1963_b.wgsl.expected.wgsl
@@ -0,0 +1,6 @@ +@group(0) @binding(0) var<storage, read_write> a : atomic<i32>; + +@compute @workgroup_size(1) +fn compute_main() { + let v = atomicCompareExchangeWeak(&(a), 1, 1).old_value; +}
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl index 994456a..0a188ef 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resulti32 { int old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - int old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -20,7 +19,7 @@ void atomicCompareExchangeWeak_1bd40a() { x__atomic_compare_exchange_resulti32 res = (x__atomic_compare_exchange_resulti32)0; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1, 1); + const atomic_compare_exchange_result_i32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1, 1); const int old_value_1 = tint_symbol.old_value; const int x_19 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_1 = {x_19, (x_19 == 1)};
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl index 994456a..0a188ef 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resulti32 { int old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - int old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -20,7 +19,7 @@ void atomicCompareExchangeWeak_1bd40a() { x__atomic_compare_exchange_resulti32 res = (x__atomic_compare_exchange_resulti32)0; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1, 1); + const atomic_compare_exchange_result_i32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1, 1); const int old_value_1 = tint_symbol.old_value; const int x_19 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_1 = {x_19, (x_19 == 1)};
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl index 41c8e6f..de0fce4 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_u32 { + uint old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resultu32 { uint old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - uint old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -20,7 +19,7 @@ void atomicCompareExchangeWeak_63d8e6() { x__atomic_compare_exchange_resultu32 res = (x__atomic_compare_exchange_resultu32)0; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); + const atomic_compare_exchange_result_u32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); const uint old_value_1 = tint_symbol.old_value; const uint x_17 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_1 = {x_17, (x_17 == 1u)};
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl index 41c8e6f..de0fce4 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_u32 { + uint old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resultu32 { uint old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - uint old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -20,7 +19,7 @@ void atomicCompareExchangeWeak_63d8e6() { x__atomic_compare_exchange_resultu32 res = (x__atomic_compare_exchange_resultu32)0; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); + const atomic_compare_exchange_result_u32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); const uint old_value_1 = tint_symbol.old_value; const uint x_17 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_1 = {x_17, (x_17 == 1u)};
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl index e9bd8f6..97da599 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.dxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resulti32 { int old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - int old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -26,7 +25,7 @@ arg_2 = 1; const int x_23 = arg_2; const int x_24 = arg_1; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_24, x_23); + const atomic_compare_exchange_result_i32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_24, x_23); const int old_value_1 = tint_symbol.old_value; const int x_25 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_1 = {x_25, (x_25 == x_23)};
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl index e9bd8f6..97da599 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.fxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resulti32 { int old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - int old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -26,7 +25,7 @@ arg_2 = 1; const int x_23 = arg_2; const int x_24 = arg_1; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_24, x_23); + const atomic_compare_exchange_result_i32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_24, x_23); const int old_value_1 = tint_symbol.old_value; const int x_25 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_1 = {x_25, (x_25 == x_23)};
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl index ec3eff1..368a192 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.dxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_u32 { + uint old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resultu32 { uint old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - uint old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -26,7 +25,7 @@ arg_2 = 1u; const uint x_21 = arg_2; const uint x_22 = arg_1; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_22, x_21); + const atomic_compare_exchange_result_u32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_22, x_21); const uint old_value_1 = tint_symbol.old_value; const uint x_23 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_1 = {x_23, (x_23 == x_21)};
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl index ec3eff1..368a192 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.fxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_u32 { + uint old_value; + bool exchanged; +}; struct x__atomic_compare_exchange_resultu32 { uint old_value; bool exchanged; @@ -5,13 +9,8 @@ RWByteAddressBuffer sb_rw : register(u0); -struct atomic_compare_exchange_weak_ret_type { - uint old_value; - bool exchanged; -}; - -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -26,7 +25,7 @@ arg_2 = 1u; const uint x_21 = arg_2; const uint x_22 = arg_1; - const atomic_compare_exchange_weak_ret_type tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_22, x_21); + const atomic_compare_exchange_result_u32 tint_symbol = sb_rwatomicCompareExchangeWeak(0u, x_22, x_21); const uint old_value_1 = tint_symbol.old_value; const uint x_23 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_1 = {x_23, (x_23 == x_21)};
diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl index 0d3c442..53eff72 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -14,7 +13,7 @@ void atomicCompareExchangeWeak_1bd40a() { - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, 1, 1); + atomic_compare_exchange_result_i32 res = sb_rwatomicCompareExchangeWeak(0u, 1, 1); } void fragment_main() {
diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl index 0d3c442..53eff72 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -14,7 +13,7 @@ void atomicCompareExchangeWeak_1bd40a() { - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, 1, 1); + atomic_compare_exchange_result_i32 res = sb_rwatomicCompareExchangeWeak(0u, 1, 1); } void fragment_main() {
diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl index 4a8edfc..904b54e 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -14,7 +13,7 @@ void atomicCompareExchangeWeak_63d8e6() { - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); + atomic_compare_exchange_result_u32 res = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); } void fragment_main() {
diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl index 4a8edfc..904b54e 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -14,7 +13,7 @@ void atomicCompareExchangeWeak_63d8e6() { - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); + atomic_compare_exchange_result_u32 res = sb_rwatomicCompareExchangeWeak(0u, 1u, 1u); } void fragment_main() {
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl index 6d7060e..72ad802 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.dxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -16,7 +15,7 @@ void atomicCompareExchangeWeak_1bd40a() { int arg_1 = 1; int arg_2 = 1; - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); + atomic_compare_exchange_result_i32 res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); } void fragment_main() {
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl index 6d7060e..72ad802 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.fxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 sb_rwatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -16,7 +15,7 @@ void atomicCompareExchangeWeak_1bd40a() { int arg_1 = 1; int arg_2 = 1; - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); + atomic_compare_exchange_result_i32 res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); } void fragment_main() {
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl index 25ba9b8..c1804c6 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.dxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -16,7 +15,7 @@ void atomicCompareExchangeWeak_63d8e6() { uint arg_1 = 1u; uint arg_2 = 1u; - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); + atomic_compare_exchange_result_u32 res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); } void fragment_main() {
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl index 25ba9b8..c1804c6 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.fxc.hlsl
@@ -1,12 +1,11 @@ -RWByteAddressBuffer sb_rw : register(u0); - -struct atomic_compare_exchange_weak_ret_type { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; +RWByteAddressBuffer sb_rw : register(u0); -atomic_compare_exchange_weak_ret_type sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_u32 sb_rwatomicCompareExchangeWeak(uint offset, uint compare, uint value) { + atomic_compare_exchange_result_u32 result=(atomic_compare_exchange_result_u32)0; sb_rw.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -16,7 +15,7 @@ void atomicCompareExchangeWeak_63d8e6() { uint arg_1 = 1u; uint arg_2 = 1u; - atomic_compare_exchange_weak_ret_type res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); + atomic_compare_exchange_result_u32 res = sb_rwatomicCompareExchangeWeak(0u, arg_1, arg_2); } void fragment_main() {
diff --git a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.dxc.hlsl b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.dxc.hlsl index 248ab255..2f5f807 100644 --- a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.dxc.hlsl +++ b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.dxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; static bool tint_discarded = false; struct tint_symbol_2 { @@ -10,13 +14,9 @@ struct tint_symbol { int value : SV_Target0; }; -struct atomic_compare_exchange_weak_ret_type { - int old_value; - bool exchanged; -}; -atomic_compare_exchange_weak_ret_type aatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 aatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; a.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -28,7 +28,7 @@ int x = 0; tint_symbol_2 tint_symbol_1 = (tint_symbol_2)0; if (!(tint_discarded)) { - const atomic_compare_exchange_weak_ret_type tint_symbol_3 = aatomicCompareExchangeWeak(0u, 0, 1); + const atomic_compare_exchange_result_i32 tint_symbol_3 = aatomicCompareExchangeWeak(0u, 0, 1); tint_symbol_1.old_value = tint_symbol_3.old_value; tint_symbol_1.exchanged = tint_symbol_3.exchanged; }
diff --git a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.fxc.hlsl b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.fxc.hlsl index 248ab255..2f5f807 100644 --- a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.fxc.hlsl +++ b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.fxc.hlsl
@@ -1,3 +1,7 @@ +struct atomic_compare_exchange_result_i32 { + int old_value; + bool exchanged; +}; static bool tint_discarded = false; struct tint_symbol_2 { @@ -10,13 +14,9 @@ struct tint_symbol { int value : SV_Target0; }; -struct atomic_compare_exchange_weak_ret_type { - int old_value; - bool exchanged; -}; -atomic_compare_exchange_weak_ret_type aatomicCompareExchangeWeak(uint offset, int compare, int value) { - atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; +atomic_compare_exchange_result_i32 aatomicCompareExchangeWeak(uint offset, int compare, int value) { + atomic_compare_exchange_result_i32 result=(atomic_compare_exchange_result_i32)0; a.InterlockedCompareExchange(offset, compare, value, result.old_value); result.exchanged = result.old_value == compare; return result; @@ -28,7 +28,7 @@ int x = 0; tint_symbol_2 tint_symbol_1 = (tint_symbol_2)0; if (!(tint_discarded)) { - const atomic_compare_exchange_weak_ret_type tint_symbol_3 = aatomicCompareExchangeWeak(0u, 0, 1); + const atomic_compare_exchange_result_i32 tint_symbol_3 = aatomicCompareExchangeWeak(0u, 0, 1); tint_symbol_1.old_value = tint_symbol_3.old_value; tint_symbol_1.exchanged = tint_symbol_3.exchanged; }