From 75db82c96bbfadd8e28cc2ecda0fe635324be988 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Fri, 18 Jun 2021 22:44:31 +0000 Subject: [PATCH] sanitizers: Use the ZeroInitWorkgroupMemory transform Zero the workgroup memory for all backends. We can probably disable this for the backends that support workgroup zeroing, but that's an optimization we can perform later. Fixed: tint:280 Change-Id: I9cad919ba3a15b8cedfe6939317d1f6b95425453 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55244 Kokoro: Kokoro Reviewed-by: David Neto Reviewed-by: James Price --- src/transform/hlsl.cc | 4 + src/transform/msl.cc | 10 +- src/transform/msl_test.cc | 44 +++--- src/transform/spirv.cc | 2 + src/writer/hlsl/generator_impl.cc | 2 + src/writer/msl/generator_impl.cc | 4 +- ...rator_impl_variable_decl_statement_test.cc | 3 +- .../gen/atomicAdd/794055.wgsl.expected.hlsl | 12 +- .../gen/atomicAdd/794055.wgsl.expected.spvasm | 41 ++++-- .../gen/atomicAdd/d5db1d.wgsl.expected.hlsl | 12 +- .../gen/atomicAdd/d5db1d.wgsl.expected.spvasm | 39 ++++-- .../gen/atomicAnd/34edd3.wgsl.expected.hlsl | 12 +- .../gen/atomicAnd/34edd3.wgsl.expected.spvasm | 39 ++++-- .../gen/atomicAnd/45a819.wgsl.expected.hlsl | 12 +- .../gen/atomicAnd/45a819.wgsl.expected.spvasm | 41 ++++-- .../89ea3b.wgsl.expected.hlsl | 12 +- .../89ea3b.wgsl.expected.spvasm | 49 ++++--- .../b2ab2c.wgsl.expected.hlsl | 12 +- .../b2ab2c.wgsl.expected.spvasm | 45 ++++--- .../atomicExchange/0a5dca.wgsl.expected.hlsl | 12 +- .../0a5dca.wgsl.expected.spvasm | 39 ++++-- .../atomicExchange/e114ba.wgsl.expected.hlsl | 12 +- .../e114ba.wgsl.expected.spvasm | 41 ++++-- .../gen/atomicLoad/361bf1.wgsl.expected.hlsl | 12 +- .../atomicLoad/361bf1.wgsl.expected.spvasm | 39 ++++-- .../gen/atomicLoad/afcc03.wgsl.expected.hlsl | 12 +- .../atomicLoad/afcc03.wgsl.expected.spvasm | 41 ++++-- .../gen/atomicMax/a89cc3.wgsl.expected.hlsl | 12 +- .../gen/atomicMax/a89cc3.wgsl.expected.spvasm | 41 ++++-- .../gen/atomicMax/beccfc.wgsl.expected.hlsl | 12 +- .../gen/atomicMax/beccfc.wgsl.expected.spvasm | 39 ++++-- .../gen/atomicMin/278235.wgsl.expected.hlsl | 12 +- .../gen/atomicMin/278235.wgsl.expected.spvasm | 41 ++++-- .../gen/atomicMin/69d383.wgsl.expected.hlsl | 12 +- .../gen/atomicMin/69d383.wgsl.expected.spvasm | 39 ++++-- .../gen/atomicOr/5e3d61.wgsl.expected.hlsl | 12 +- .../gen/atomicOr/5e3d61.wgsl.expected.spvasm | 39 ++++-- .../gen/atomicOr/d09248.wgsl.expected.hlsl | 12 +- .../gen/atomicOr/d09248.wgsl.expected.spvasm | 41 ++++-- .../gen/atomicStore/726882.wgsl.expected.hlsl | 12 +- .../atomicStore/726882.wgsl.expected.spvasm | 32 +++-- .../gen/atomicStore/8bea94.wgsl.expected.hlsl | 12 +- .../atomicStore/8bea94.wgsl.expected.spvasm | 34 +++-- .../gen/atomicXor/75dc95.wgsl.expected.hlsl | 12 +- .../gen/atomicXor/75dc95.wgsl.expected.spvasm | 41 ++++-- .../gen/atomicXor/c8e6be.wgsl.expected.hlsl | 12 +- .../gen/atomicXor/c8e6be.wgsl.expected.spvasm | 39 ++++-- .../gen/frexp/0da285.wgsl.expected.hlsl | 11 +- .../gen/frexp/0da285.wgsl.expected.spvasm | 45 +++++-- .../gen/frexp/40fc9b.wgsl.expected.hlsl | 11 +- .../gen/frexp/40fc9b.wgsl.expected.spvasm | 45 +++++-- .../gen/frexp/6be229.wgsl.expected.hlsl | 11 +- .../gen/frexp/6be229.wgsl.expected.spvasm | 44 ++++-- .../gen/frexp/841515.wgsl.expected.hlsl | 11 +- .../gen/frexp/841515.wgsl.expected.spvasm | 44 ++++-- .../gen/frexp/8b3191.wgsl.expected.hlsl | 11 +- .../gen/frexp/8b3191.wgsl.expected.spvasm | 44 ++++-- .../gen/frexp/a3f940.wgsl.expected.hlsl | 11 +- .../gen/frexp/a3f940.wgsl.expected.spvasm | 45 +++++-- .../gen/frexp/b87f4e.wgsl.expected.hlsl | 11 +- .../gen/frexp/b87f4e.wgsl.expected.spvasm | 45 +++++-- .../gen/frexp/e8b4d1.wgsl.expected.hlsl | 11 +- .../gen/frexp/e8b4d1.wgsl.expected.spvasm | 44 ++++-- .../gen/modf/1d59e5.wgsl.expected.hlsl | 11 +- .../gen/modf/1d59e5.wgsl.expected.spvasm | 44 ++++-- .../gen/modf/a128ab.wgsl.expected.hlsl | 11 +- .../gen/modf/a128ab.wgsl.expected.spvasm | 44 ++++-- .../gen/modf/bb9088.wgsl.expected.hlsl | 11 +- .../gen/modf/bb9088.wgsl.expected.spvasm | 44 ++++-- .../gen/modf/e38ae6.wgsl.expected.hlsl | 11 +- .../gen/modf/e38ae6.wgsl.expected.spvasm | 44 ++++-- .../local/ptr_workgroup.wgsl.expected.hlsl | 11 +- .../local/ptr_workgroup.wgsl.expected.msl | 12 +- .../local/ptr_workgroup.wgsl.expected.spvasm | 33 ++++- test/var/workgroup.wgsl.expected.hlsl | 34 ++++- test/var/workgroup.wgsl.expected.msl | 59 ++++---- test/var/workgroup.wgsl.expected.spvasm | 126 ++++++++++++------ 77 files changed, 1509 insertions(+), 539 deletions(-) diff --git a/src/transform/hlsl.cc b/src/transform/hlsl.cc index bacc7f637a..1009f053e4 100644 --- a/src/transform/hlsl.cc +++ b/src/transform/hlsl.cc @@ -27,6 +27,7 @@ #include "src/transform/promote_initializers_to_const_var.h" #include "src/transform/simplify.h" #include "src/transform/wrap_arrays_in_structs.h" +#include "src/transform/zero_init_workgroup_memory.h" namespace tint { namespace transform { @@ -37,6 +38,9 @@ Hlsl::~Hlsl() = default; Output Hlsl::Run(const Program* in, const DataMap&) { Manager manager; DataMap data; + // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as + // ZeroInitWorkgroupMemory may inject new builtin parameters. + manager.Add(); manager.Add(); manager.Add(); // Simplify cleans up messy `*(&(expr))` expressions from InlinePointerLets. diff --git a/src/transform/msl.cc b/src/transform/msl.cc index 10e3ea27ec..f32c01aabb 100644 --- a/src/transform/msl.cc +++ b/src/transform/msl.cc @@ -32,6 +32,7 @@ #include "src/transform/promote_initializers_to_const_var.h" #include "src/transform/simplify.h" #include "src/transform/wrap_arrays_in_structs.h" +#include "src/transform/zero_init_workgroup_memory.h" namespace tint { namespace transform { @@ -42,6 +43,9 @@ Msl::~Msl() = default; Output Msl::Run(const Program* in, const DataMap&) { Manager manager; DataMap data; + // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as + // ZeroInitWorkgroupMemory may inject new builtin parameters. + manager.Add(); manager.Add(); manager.Add(); manager.Add(); @@ -173,9 +177,9 @@ void Msl::HandleModuleScopeVariables(CloneContext& ctx) const { ctx.dst->ID(), ast::DisabledValidation::kFunctionVarStorageClass); auto* constructor = ctx.Clone(var->Declaration()->constructor()); - auto* local_var = ctx.dst->Var(new_var_symbol, store_type, - var->StorageClass(), constructor, - ast::DecorationList{disable_validation}); + auto* local_var = ctx.dst->Var( + new_var_symbol, store_type, var->StorageClass(), constructor, + ast::DecorationList{disable_validation}); ctx.InsertFront(func_ast->body()->statements(), ctx.dst->Decl(local_var)); } diff --git a/src/transform/msl_test.cc b/src/transform/msl_test.cc index 6e19c44ef7..6b0968c501 100644 --- a/src/transform/msl_test.cc +++ b/src/transform/msl_test.cc @@ -35,10 +35,14 @@ fn main() { auto* expect = R"( [[stage(compute)]] -fn main() { - [[internal(disable_validation__function_var_storage_class)]] var tint_symbol : f32; - [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_1 : f32; - tint_symbol = tint_symbol_1; +fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) { + [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_1 : f32; + [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_2 : f32; + if ((local_invocation_index == 0u)) { + tint_symbol_1 = f32(); + } + workgroupBarrier(); + tint_symbol_1 = tint_symbol_2; } )"; @@ -76,22 +80,26 @@ fn main() { fn no_uses() { } -fn bar(a : f32, b : f32, tint_symbol : ptr, tint_symbol_1 : ptr) { - *(tint_symbol) = a; - *(tint_symbol_1) = b; +fn bar(a : f32, b : f32, tint_symbol_1 : ptr, tint_symbol_2 : ptr) { + *(tint_symbol_1) = a; + *(tint_symbol_2) = b; } -fn foo(a : f32, tint_symbol_2 : ptr, tint_symbol_3 : ptr) { +fn foo(a : f32, tint_symbol_3 : ptr, tint_symbol_4 : ptr) { let b : f32 = 2.0; - bar(a, b, tint_symbol_2, tint_symbol_3); + bar(a, b, tint_symbol_3, tint_symbol_4); no_uses(); } [[stage(compute)]] -fn main() { - [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_4 : f32; +fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) { [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_5 : f32; - foo(1.0, &(tint_symbol_4), &(tint_symbol_5)); + [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_6 : f32; + if ((local_invocation_index == 0u)) { + tint_symbol_5 = f32(); + } + workgroupBarrier(); + foo(1.0, &(tint_symbol_6), &(tint_symbol_5)); } )"; @@ -141,11 +149,15 @@ fn main() { auto* expect = R"( [[stage(compute)]] -fn main() { - [[internal(disable_validation__function_var_storage_class)]] var tint_symbol : f32; +fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) { [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_1 : f32; - let x : f32 = (tint_symbol + tint_symbol_1); - tint_symbol = x; + [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_2 : f32; + if ((local_invocation_index == 0u)) { + tint_symbol_1 = f32(); + } + workgroupBarrier(); + let x : f32 = (tint_symbol_2 + tint_symbol_1); + tint_symbol_2 = x; } )"; diff --git a/src/transform/spirv.cc b/src/transform/spirv.cc index 456ee08ca8..81fdbb9208 100644 --- a/src/transform/spirv.cc +++ b/src/transform/spirv.cc @@ -31,6 +31,7 @@ #include "src/transform/inline_pointer_lets.h" #include "src/transform/manager.h" #include "src/transform/simplify.h" +#include "src/transform/zero_init_workgroup_memory.h" TINT_INSTANTIATE_TYPEINFO(tint::transform::Spirv::Config); @@ -42,6 +43,7 @@ Spirv::~Spirv() = default; Output Spirv::Run(const Program* in, const DataMap& data) { Manager manager; + manager.Add(); manager.Add(); // Required for arrayLength() manager.Add(); // Required for arrayLength() manager.Add(); diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index a8b5fca05c..810985118c 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -2293,6 +2293,8 @@ bool GeneratorImpl::EmitEntryPointFunction(std::ostream& out, auto* sem = builder_.Sem().Get(var); auto* type = sem->Type(); if (!type->Is()) { + // ICE likely indicates that the CanonicalizeEntryPointIO transform was + // not run, or a builtin parameter was added after it was run. TINT_ICE(diagnostics_) << "Unsupported non-struct entry point parameter"; } diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 274174555b..460b059b46 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -2356,16 +2356,16 @@ bool GeneratorImpl::EmitVariable(const sem::Variable* var, } if (!skip_constructor) { - out_ << " = "; if (decl->constructor() != nullptr) { + out_ << " = "; if (!EmitExpression(decl->constructor())) { return false; } } else if (var->StorageClass() == ast::StorageClass::kPrivate || var->StorageClass() == ast::StorageClass::kFunction || - var->StorageClass() == ast::StorageClass::kWorkgroup || var->StorageClass() == ast::StorageClass::kNone || var->StorageClass() == ast::StorageClass::kOutput) { + out_ << " = "; if (!EmitZeroValue(type)) { return false; } diff --git a/src/writer/msl/generator_impl_variable_decl_statement_test.cc b/src/writer/msl/generator_impl_variable_decl_statement_test.cc index d3eea3ead2..61c8b66c52 100644 --- a/src/writer/msl/generator_impl_variable_decl_statement_test.cc +++ b/src/writer/msl/generator_impl_variable_decl_statement_test.cc @@ -146,8 +146,7 @@ TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Workgroup) { gen.increment_indent(); ASSERT_TRUE(gen.Generate()) << gen.error(); - EXPECT_THAT(gen.result(), - HasSubstr("threadgroup float tint_symbol_1 = 0.0f;\n")); + EXPECT_THAT(gen.result(), HasSubstr("threadgroup float tint_symbol_2;\n")); } TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Initializer_ZeroVec) { diff --git a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.hlsl b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.hlsl index 58d211c7b2..0bf78633e0 100644 --- a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicAdd_794055() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicAdd_794055(); return; } diff --git a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.spvasm b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.spvasm index e2b9026b7d..a3b6ce2b5d 100644 --- a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.spvasm @@ -1,36 +1,51 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 31 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicAdd_794055 "atomicAdd_794055" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %_ptr_Function_int = OpTypePointer Function %int - %16 = OpConstantNull %int -%atomicAdd_794055 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %16 - %8 = OpAtomicIAdd %int %arg_0 %uint_2 %uint_0 %int_1 - OpStore %res %8 + %18 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicAdd_794055 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %18 + %11 = OpAtomicIAdd %int %arg_0 %uint_2 %uint_0 %int_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %18 = OpLabel - %19 = OpFunctionCall %void %atomicAdd_794055 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %18 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %atomicAdd_794055 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.hlsl b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.hlsl index b14673c092..a7edfc6cbd 100644 --- a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicAdd_d5db1d() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicAdd_d5db1d(); return; } diff --git a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.spvasm b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.spvasm index 33ae4883d8..f174f0c905 100644 --- a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicAdd_d5db1d "atomicAdd_d5db1d" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %_ptr_Function_uint = OpTypePointer Function %uint - %15 = OpConstantNull %uint -%atomicAdd_d5db1d = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %15 - %8 = OpAtomicIAdd %uint %arg_0 %uint_2 %uint_0 %uint_1 - OpStore %res %8 + %17 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicAdd_d5db1d = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %17 + %10 = OpAtomicIAdd %uint %arg_0 %uint_2 %uint_0 %uint_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicAdd_d5db1d +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicAdd_d5db1d OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.hlsl b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.hlsl index c05630abc4..538a0b59a4 100644 --- a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicAnd_34edd3() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicAnd_34edd3(); return; } diff --git a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.spvasm b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.spvasm index 9ddb9cdc1f..a46263b599 100644 --- a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicAnd_34edd3 "atomicAnd_34edd3" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %_ptr_Function_uint = OpTypePointer Function %uint - %15 = OpConstantNull %uint -%atomicAnd_34edd3 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %15 - %8 = OpAtomicAnd %uint %arg_0 %uint_2 %uint_0 %uint_1 - OpStore %res %8 + %17 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicAnd_34edd3 = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %17 + %10 = OpAtomicAnd %uint %arg_0 %uint_2 %uint_0 %uint_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicAnd_34edd3 +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicAnd_34edd3 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.hlsl b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.hlsl index 76085c8d3f..f6f3699d0f 100644 --- a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicAnd_45a819() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicAnd_45a819(); return; } diff --git a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.spvasm b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.spvasm index 5f56565b58..9f15c78a53 100644 --- a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.spvasm @@ -1,36 +1,51 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 31 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicAnd_45a819 "atomicAnd_45a819" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %_ptr_Function_int = OpTypePointer Function %int - %16 = OpConstantNull %int -%atomicAnd_45a819 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %16 - %8 = OpAtomicAnd %int %arg_0 %uint_2 %uint_0 %int_1 - OpStore %res %8 + %18 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicAnd_45a819 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %18 + %11 = OpAtomicAnd %int %arg_0 %uint_2 %uint_0 %int_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %18 = OpLabel - %19 = OpFunctionCall %void %atomicAnd_45a819 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %18 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %atomicAnd_45a819 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl index fbc219f979..82998fdff1 100644 --- a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl @@ -8,8 +8,18 @@ void atomicCompareExchangeWeak_89ea3b() { int2 res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicCompareExchangeWeak_89ea3b(); return; } diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm index 537dbb6aba..e868537cc0 100644 --- a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm @@ -1,42 +1,57 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 26 +; Bound: 37 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicCompareExchangeWeak_89ea3b "atomicCompareExchangeWeak_89ea3b" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void - %v2int = OpTypeVector %int 2 %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void + %v2int = OpTypeVector %int 2 %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %bool = OpTypeBool %int_0 = OpConstant %int 0 %_ptr_Function_v2int = OpTypePointer Function %v2int - %22 = OpConstantNull %v2int -%atomicCompareExchangeWeak_89ea3b = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_v2int Function %22 - %16 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1 - %17 = OpIEqual %bool %16 %int_1 - %19 = OpSelect %int %17 %int_1 %int_0 - %8 = OpCompositeConstruct %v2int %16 %19 - OpStore %res %8 + %24 = OpConstantNull %v2int + %33 = OpConstantNull %int + %uint_264 = OpConstant %uint 264 +%atomicCompareExchangeWeak_89ea3b = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_v2int Function %24 + %18 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1 + %19 = OpIEqual %bool %18 %int_1 + %21 = OpSelect %int %19 %int_1 %int_0 + %11 = OpCompositeConstruct %v2int %18 %21 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %24 = OpLabel - %25 = OpFunctionCall %void %atomicCompareExchangeWeak_89ea3b +%compute_main = OpFunction %void None %7 + %26 = OpLabel + %27 = OpLoad %uint %tint_symbol + %28 = OpIEqual %bool %27 %uint_0 + OpSelectionMerge %29 None + OpBranchConditional %28 %30 %29 + %30 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %33 + OpBranch %29 + %29 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %36 = OpFunctionCall %void %atomicCompareExchangeWeak_89ea3b OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl index 6342525293..a572f3bc0b 100644 --- a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl @@ -8,8 +8,18 @@ void atomicCompareExchangeWeak_b2ab2c() { uint2 res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicCompareExchangeWeak_b2ab2c(); return; } diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm index 9492ec65e6..de6c4a3406 100644 --- a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm @@ -1,40 +1,55 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 24 +; Bound: 35 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicCompareExchangeWeak_b2ab2c "atomicCompareExchangeWeak_b2ab2c" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %v2uint = OpTypeVector %uint 2 %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %bool = OpTypeBool %_ptr_Function_v2uint = OpTypePointer Function %v2uint - %20 = OpConstantNull %v2uint -%atomicCompareExchangeWeak_b2ab2c = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_v2uint Function %20 - %15 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1 - %16 = OpIEqual %bool %15 %uint_1 - %17 = OpSelect %uint %16 %uint_1 %uint_0 - %8 = OpCompositeConstruct %v2uint %15 %17 - OpStore %res %8 + %22 = OpConstantNull %v2uint + %31 = OpConstantNull %uint + %uint_264 = OpConstant %uint 264 +%atomicCompareExchangeWeak_b2ab2c = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_v2uint Function %22 + %17 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1 + %18 = OpIEqual %bool %17 %uint_1 + %19 = OpSelect %uint %18 %uint_1 %uint_0 + %10 = OpCompositeConstruct %v2uint %17 %19 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %22 = OpLabel - %23 = OpFunctionCall %void %atomicCompareExchangeWeak_b2ab2c +%compute_main = OpFunction %void None %6 + %24 = OpLabel + %25 = OpLoad %uint %tint_symbol + %26 = OpIEqual %bool %25 %uint_0 + OpSelectionMerge %27 None + OpBranchConditional %26 %28 %27 + %28 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %31 + OpBranch %27 + %27 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %34 = OpFunctionCall %void %atomicCompareExchangeWeak_b2ab2c OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.hlsl b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.hlsl index 7878fe1a01..d22eb78bc0 100644 --- a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicExchange_0a5dca() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicExchange_0a5dca(); return; } diff --git a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.spvasm b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.spvasm index 680a4a879a..d77463e0ae 100644 --- a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicExchange_0a5dca "atomicExchange_0a5dca" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %_ptr_Function_uint = OpTypePointer Function %uint - %15 = OpConstantNull %uint -%atomicExchange_0a5dca = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %15 - %8 = OpAtomicExchange %uint %arg_0 %uint_2 %uint_0 %uint_1 - OpStore %res %8 + %17 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicExchange_0a5dca = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %17 + %10 = OpAtomicExchange %uint %arg_0 %uint_2 %uint_0 %uint_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicExchange_0a5dca +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicExchange_0a5dca OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.hlsl b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.hlsl index 289f8b9d84..2f6a06e4c4 100644 --- a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicExchange_e114ba() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicExchange_e114ba(); return; } diff --git a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.spvasm b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.spvasm index 9232da2002..4a39abf2d4 100644 --- a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.spvasm @@ -1,36 +1,51 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 31 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicExchange_e114ba "atomicExchange_e114ba" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %_ptr_Function_int = OpTypePointer Function %int - %16 = OpConstantNull %int -%atomicExchange_e114ba = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %16 - %8 = OpAtomicExchange %int %arg_0 %uint_2 %uint_0 %int_1 - OpStore %res %8 + %18 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicExchange_e114ba = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %18 + %11 = OpAtomicExchange %int %arg_0 %uint_2 %uint_0 %int_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %18 = OpLabel - %19 = OpFunctionCall %void %atomicExchange_e114ba +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %18 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %atomicExchange_e114ba OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.hlsl b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.hlsl index 08e144602f..ffea6caee9 100644 --- a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicLoad_361bf1() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicLoad_361bf1(); return; } diff --git a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.spvasm b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.spvasm index 583063ae2e..60dd135543 100644 --- a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.spvasm @@ -1,34 +1,49 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 18 +; Bound: 29 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicLoad_361bf1 "atomicLoad_361bf1" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %_ptr_Function_uint = OpTypePointer Function %uint - %14 = OpConstantNull %uint -%atomicLoad_361bf1 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %14 - %8 = OpAtomicLoad %uint %arg_0 %uint_2 %uint_0 - OpStore %res %8 + %16 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicLoad_361bf1 = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %16 + %10 = OpAtomicLoad %uint %arg_0 %uint_2 %uint_0 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %16 = OpLabel - %17 = OpFunctionCall %void %atomicLoad_361bf1 +%compute_main = OpFunction %void None %6 + %18 = OpLabel + %19 = OpLoad %uint %tint_symbol + %20 = OpIEqual %bool %19 %uint_0 + OpSelectionMerge %22 None + OpBranchConditional %20 %23 %22 + %23 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %16 + OpBranch %22 + %22 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %28 = OpFunctionCall %void %atomicLoad_361bf1 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.hlsl b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.hlsl index bc4f09c3c9..5c7323e9bf 100644 --- a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicLoad_afcc03() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicLoad_afcc03(); return; } diff --git a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.spvasm b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.spvasm index 1b599d0966..09d5bb595f 100644 --- a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicLoad_afcc03 "atomicLoad_afcc03" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %_ptr_Function_int = OpTypePointer Function %int - %15 = OpConstantNull %int -%atomicLoad_afcc03 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %15 - %8 = OpAtomicLoad %int %arg_0 %uint_2 %uint_0 - OpStore %res %8 + %17 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicLoad_afcc03 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %17 + %11 = OpAtomicLoad %int %arg_0 %uint_2 %uint_0 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicLoad_afcc03 +%compute_main = OpFunction %void None %7 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicLoad_afcc03 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.hlsl b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.hlsl index f6a9e658fb..78d5fc37ae 100644 --- a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicMax_a89cc3() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicMax_a89cc3(); return; } diff --git a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.spvasm b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.spvasm index 810c250c3e..7a95b2378f 100644 --- a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.spvasm @@ -1,36 +1,51 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 31 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicMax_a89cc3 "atomicMax_a89cc3" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %_ptr_Function_int = OpTypePointer Function %int - %16 = OpConstantNull %int -%atomicMax_a89cc3 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %16 - %8 = OpAtomicSMax %int %arg_0 %uint_2 %uint_0 %int_1 - OpStore %res %8 + %18 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicMax_a89cc3 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %18 + %11 = OpAtomicSMax %int %arg_0 %uint_2 %uint_0 %int_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %18 = OpLabel - %19 = OpFunctionCall %void %atomicMax_a89cc3 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %18 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %atomicMax_a89cc3 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.hlsl b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.hlsl index 0183b3d340..9912c5833a 100644 --- a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicMax_beccfc() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicMax_beccfc(); return; } diff --git a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.spvasm b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.spvasm index a2dce6b3cd..9ae4cbb239 100644 --- a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicMax_beccfc "atomicMax_beccfc" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %_ptr_Function_uint = OpTypePointer Function %uint - %15 = OpConstantNull %uint -%atomicMax_beccfc = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %15 - %8 = OpAtomicUMax %uint %arg_0 %uint_2 %uint_0 %uint_1 - OpStore %res %8 + %17 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicMax_beccfc = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %17 + %10 = OpAtomicUMax %uint %arg_0 %uint_2 %uint_0 %uint_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicMax_beccfc +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicMax_beccfc OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.hlsl b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.hlsl index b6c6f08ff5..d0240e56c1 100644 --- a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicMin_278235() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicMin_278235(); return; } diff --git a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.spvasm b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.spvasm index 3eaaf73083..2ccd485f9c 100644 --- a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.spvasm @@ -1,36 +1,51 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 31 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicMin_278235 "atomicMin_278235" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %_ptr_Function_int = OpTypePointer Function %int - %16 = OpConstantNull %int -%atomicMin_278235 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %16 - %8 = OpAtomicSMin %int %arg_0 %uint_2 %uint_0 %int_1 - OpStore %res %8 + %18 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicMin_278235 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %18 + %11 = OpAtomicSMin %int %arg_0 %uint_2 %uint_0 %int_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %18 = OpLabel - %19 = OpFunctionCall %void %atomicMin_278235 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %18 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %atomicMin_278235 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.hlsl b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.hlsl index 73abc32f3a..96ed9d468a 100644 --- a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicMin_69d383() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicMin_69d383(); return; } diff --git a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.spvasm b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.spvasm index 590c53253b..8abdaf2ceb 100644 --- a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicMin_69d383 "atomicMin_69d383" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %_ptr_Function_uint = OpTypePointer Function %uint - %15 = OpConstantNull %uint -%atomicMin_69d383 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %15 - %8 = OpAtomicUMin %uint %arg_0 %uint_2 %uint_0 %uint_1 - OpStore %res %8 + %17 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicMin_69d383 = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %17 + %10 = OpAtomicUMin %uint %arg_0 %uint_2 %uint_0 %uint_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicMin_69d383 +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicMin_69d383 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.hlsl b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.hlsl index 8b6e92c5b1..55eae7e4ee 100644 --- a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicOr_5e3d61() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicOr_5e3d61(); return; } diff --git a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.spvasm b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.spvasm index 4376949538..e6afd84ff2 100644 --- a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicOr_5e3d61 "atomicOr_5e3d61" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %_ptr_Function_uint = OpTypePointer Function %uint - %15 = OpConstantNull %uint -%atomicOr_5e3d61 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %15 - %8 = OpAtomicOr %uint %arg_0 %uint_2 %uint_0 %uint_1 - OpStore %res %8 + %17 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicOr_5e3d61 = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %17 + %10 = OpAtomicOr %uint %arg_0 %uint_2 %uint_0 %uint_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicOr_5e3d61 +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicOr_5e3d61 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.hlsl b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.hlsl index 889ce1cba5..2037316955 100644 --- a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicOr_d09248() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicOr_d09248(); return; } diff --git a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.spvasm b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.spvasm index 46fbd2b16e..495683dd25 100644 --- a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.spvasm @@ -1,36 +1,51 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 31 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicOr_d09248 "atomicOr_d09248" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %_ptr_Function_int = OpTypePointer Function %int - %16 = OpConstantNull %int -%atomicOr_d09248 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %16 - %8 = OpAtomicOr %int %arg_0 %uint_2 %uint_0 %int_1 - OpStore %res %8 + %18 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicOr_d09248 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %18 + %11 = OpAtomicOr %int %arg_0 %uint_2 %uint_0 %int_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %18 = OpLabel - %19 = OpFunctionCall %void %atomicOr_d09248 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %18 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %atomicOr_d09248 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.hlsl b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.hlsl index 65a71a8971..6890fa3e6e 100644 --- a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.hlsl @@ -5,8 +5,18 @@ void atomicStore_726882() { InterlockedExchange(arg_0, 1u, atomic_result); } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicStore_726882(); return; } diff --git a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.spvasm b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.spvasm index 58aa6c6e8b..aeeebac399 100644 --- a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.spvasm @@ -1,30 +1,46 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 16 +; Bound: 28 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicStore_726882 "atomicStore_726882" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 -%atomicStore_726882 = OpFunction %void None %4 - %7 = OpLabel + %bool = OpTypeBool + %24 = OpConstantNull %uint + %uint_264 = OpConstant %uint 264 +%atomicStore_726882 = OpFunction %void None %6 + %9 = OpLabel OpAtomicStore %arg_0 %uint_2 %uint_0 %uint_1 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %14 = OpLabel - %15 = OpFunctionCall %void %atomicStore_726882 +%compute_main = OpFunction %void None %6 + %16 = OpLabel + %17 = OpLoad %uint %tint_symbol + %18 = OpIEqual %bool %17 %uint_0 + OpSelectionMerge %20 None + OpBranchConditional %18 %21 %20 + %21 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %24 + OpBranch %20 + %20 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %27 = OpFunctionCall %void %atomicStore_726882 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.hlsl b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.hlsl index 31257e0254..b0c26d25f6 100644 --- a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.hlsl @@ -5,8 +5,18 @@ void atomicStore_8bea94() { InterlockedExchange(arg_0, 1, atomic_result); } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicStore_8bea94(); return; } diff --git a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.spvasm b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.spvasm index c8a60c0c2c..04d4ba9f81 100644 --- a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.spvasm @@ -1,31 +1,47 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 17 +; Bound: 29 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicStore_8bea94 "atomicStore_8bea94" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 -%atomicStore_8bea94 = OpFunction %void None %4 - %7 = OpLabel + %bool = OpTypeBool + %25 = OpConstantNull %int + %uint_264 = OpConstant %uint 264 +%atomicStore_8bea94 = OpFunction %void None %7 + %10 = OpLabel OpAtomicStore %arg_0 %uint_2 %uint_0 %int_1 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %15 = OpLabel - %16 = OpFunctionCall %void %atomicStore_8bea94 +%compute_main = OpFunction %void None %7 + %17 = OpLabel + %18 = OpLoad %uint %tint_symbol + %19 = OpIEqual %bool %18 %uint_0 + OpSelectionMerge %21 None + OpBranchConditional %19 %22 %21 + %22 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %25 + OpBranch %21 + %21 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %28 = OpFunctionCall %void %atomicStore_8bea94 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.hlsl b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.hlsl index a2f200590c..1e096e1e40 100644 --- a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicXor_75dc95() { int res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + int atomic_result_1 = 0; + InterlockedExchange(arg_0, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicXor_75dc95(); return; } diff --git a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.spvasm b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.spvasm index 110db2b13e..3da5305c36 100644 --- a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.spvasm @@ -1,36 +1,51 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 31 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicXor_75dc95 "atomicXor_75dc95" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup - %void = OpTypeVoid - %4 = OpTypeFunction %void %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input + %void = OpTypeVoid + %7 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 %_ptr_Function_int = OpTypePointer Function %int - %16 = OpConstantNull %int -%atomicXor_75dc95 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_int Function %16 - %8 = OpAtomicXor %int %arg_0 %uint_2 %uint_0 %int_1 - OpStore %res %8 + %18 = OpConstantNull %int + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicXor_75dc95 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_int Function %18 + %11 = OpAtomicXor %int %arg_0 %uint_2 %uint_0 %int_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %18 = OpLabel - %19 = OpFunctionCall %void %atomicXor_75dc95 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %18 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %atomicXor_75dc95 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.hlsl b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.hlsl index afa7fd5127..d2c96fdecc 100644 --- a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.hlsl +++ b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.hlsl @@ -6,8 +6,18 @@ void atomicXor_c8e6be() { uint res = atomic_result; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + uint atomic_result_1 = 0u; + InterlockedExchange(arg_0, 0u, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); atomicXor_c8e6be(); return; } diff --git a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.spvasm b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.spvasm index cfd20e6cc1..9a88f16b40 100644 --- a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.spvasm +++ b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.spvasm @@ -1,35 +1,50 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_0 "arg_0" + OpName %tint_symbol "tint_symbol" OpName %atomicXor_c8e6be "atomicXor_c8e6be" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 %_ptr_Function_uint = OpTypePointer Function %uint - %15 = OpConstantNull %uint -%atomicXor_c8e6be = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_uint Function %15 - %8 = OpAtomicXor %uint %arg_0 %uint_2 %uint_0 %uint_1 - OpStore %res %8 + %17 = OpConstantNull %uint + %bool = OpTypeBool + %uint_264 = OpConstant %uint 264 +%atomicXor_c8e6be = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %17 + %10 = OpAtomicXor %uint %arg_0 %uint_2 %uint_0 %uint_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %atomicXor_c8e6be +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %21 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %17 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %atomicXor_c8e6be OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl index 91ae845b6a..b841839f57 100644 --- a/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_0da285() { float res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = 0; + } + GroupMemoryBarrierWithGroupSync(); frexp_0da285(); return; } diff --git a/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm index dcb100b106..3ea3539422 100644 --- a/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm @@ -1,35 +1,54 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 32 ; Schema: 0 OpCapability Shader - %10 = OpExtInstImport "GLSL.std.450" + %13 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_0da285 "frexp_0da285" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %arg_1 = OpVariable %_ptr_Workgroup_int Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %7 = OpTypeFunction %void %float = OpTypeFloat 32 %float_1 = OpConstant %float 1 %_ptr_Function_float = OpTypePointer Function %float - %15 = OpConstantNull %float -%frexp_0da285 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_float Function %15 - %8 = OpExtInst %float %10 Frexp %float_1 %arg_1 - OpStore %res %8 + %18 = OpConstantNull %float + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %27 = OpConstantNull %int + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_0da285 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_float Function %18 + %11 = OpExtInst %float %13 Frexp %float_1 %arg_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %frexp_0da285 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %23 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %25 None + OpBranchConditional %23 %26 %25 + %26 = OpLabel + OpStore %arg_1 %27 + OpBranch %25 + %25 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %31 = OpFunctionCall %void %frexp_0da285 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl index ab6efc6991..f7f2ecdf09 100644 --- a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_40fc9b() { float3 res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = int3(0, 0, 0); + } + GroupMemoryBarrierWithGroupSync(); frexp_40fc9b(); return; } diff --git a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm index 32d27faf97..2facf1be33 100644 --- a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm @@ -1,36 +1,55 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 33 ; Schema: 0 OpCapability Shader - %12 = OpExtInstImport "GLSL.std.450" + %15 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_40fc9b "frexp_40fc9b" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %v3int = OpTypeVector %int 3 %_ptr_Workgroup_v3int = OpTypePointer Workgroup %v3int %arg_1 = OpVariable %_ptr_Workgroup_v3int Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void + %8 = OpTypeFunction %void %float = OpTypeFloat 32 %v3float = OpTypeVector %float 3 - %13 = OpConstantNull %v3float + %16 = OpConstantNull %v3float %_ptr_Function_v3float = OpTypePointer Function %v3float -%frexp_40fc9b = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v3float Function %13 - %9 = OpExtInst %v3float %12 Frexp %13 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %28 = OpConstantNull %v3int + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_40fc9b = OpFunction %void None %8 + %11 = OpLabel + %res = OpVariable %_ptr_Function_v3float Function %16 + %12 = OpExtInst %v3float %15 Frexp %16 %arg_1 + OpStore %res %12 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %18 = OpLabel - %19 = OpFunctionCall %void %frexp_40fc9b +%compute_main = OpFunction %void None %8 + %21 = OpLabel + %22 = OpLoad %uint %tint_symbol + %24 = OpIEqual %bool %22 %uint_0 + OpSelectionMerge %26 None + OpBranchConditional %24 %27 %26 + %27 = OpLabel + OpStore %arg_1 %28 + OpBranch %26 + %26 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %32 = OpFunctionCall %void %frexp_40fc9b OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/6be229.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/6be229.wgsl.expected.hlsl index 7c0540905e..aa0759b551 100644 --- a/test/intrinsics/gen/frexp/6be229.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/6be229.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_6be229() { float3 res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = uint3(0u, 0u, 0u); + } + GroupMemoryBarrierWithGroupSync(); frexp_6be229(); return; } diff --git a/test/intrinsics/gen/frexp/6be229.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/6be229.wgsl.expected.spvasm index dad453258c..059ff2d468 100644 --- a/test/intrinsics/gen/frexp/6be229.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/6be229.wgsl.expected.spvasm @@ -1,36 +1,54 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 32 ; Schema: 0 OpCapability Shader - %12 = OpExtInstImport "GLSL.std.450" + %14 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_6be229 "frexp_6be229" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %v3uint = OpTypeVector %uint 3 %_ptr_Workgroup_v3uint = OpTypePointer Workgroup %v3uint %arg_1 = OpVariable %_ptr_Workgroup_v3uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void + %7 = OpTypeFunction %void %float = OpTypeFloat 32 %v3float = OpTypeVector %float 3 - %13 = OpConstantNull %v3float + %15 = OpConstantNull %v3float %_ptr_Function_v3float = OpTypePointer Function %v3float -%frexp_6be229 = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v3float Function %13 - %9 = OpExtInst %v3float %12 Frexp %13 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %27 = OpConstantNull %v3uint + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_6be229 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_v3float Function %15 + %11 = OpExtInst %v3float %14 Frexp %15 %arg_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %18 = OpLabel - %19 = OpFunctionCall %void %frexp_6be229 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %23 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %25 None + OpBranchConditional %23 %26 %25 + %26 = OpLabel + OpStore %arg_1 %27 + OpBranch %25 + %25 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %31 = OpFunctionCall %void %frexp_6be229 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/841515.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/841515.wgsl.expected.hlsl index 2ea0a1efd1..70551bf960 100644 --- a/test/intrinsics/gen/frexp/841515.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/841515.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_841515() { float4 res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = uint4(0u, 0u, 0u, 0u); + } + GroupMemoryBarrierWithGroupSync(); frexp_841515(); return; } diff --git a/test/intrinsics/gen/frexp/841515.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/841515.wgsl.expected.spvasm index a0e108756e..e023805836 100644 --- a/test/intrinsics/gen/frexp/841515.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/841515.wgsl.expected.spvasm @@ -1,36 +1,54 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 32 ; Schema: 0 OpCapability Shader - %12 = OpExtInstImport "GLSL.std.450" + %14 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_841515 "frexp_841515" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %v4uint = OpTypeVector %uint 4 %_ptr_Workgroup_v4uint = OpTypePointer Workgroup %v4uint %arg_1 = OpVariable %_ptr_Workgroup_v4uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void + %7 = OpTypeFunction %void %float = OpTypeFloat 32 %v4float = OpTypeVector %float 4 - %13 = OpConstantNull %v4float + %15 = OpConstantNull %v4float %_ptr_Function_v4float = OpTypePointer Function %v4float -%frexp_841515 = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v4float Function %13 - %9 = OpExtInst %v4float %12 Frexp %13 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %27 = OpConstantNull %v4uint + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_841515 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_v4float Function %15 + %11 = OpExtInst %v4float %14 Frexp %15 %arg_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %18 = OpLabel - %19 = OpFunctionCall %void %frexp_841515 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %23 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %25 None + OpBranchConditional %23 %26 %25 + %26 = OpLabel + OpStore %arg_1 %27 + OpBranch %25 + %25 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %31 = OpFunctionCall %void %frexp_841515 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/8b3191.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/8b3191.wgsl.expected.hlsl index 38564b5bf0..2f3d67b2c0 100644 --- a/test/intrinsics/gen/frexp/8b3191.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/8b3191.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_8b3191() { float res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = 0u; + } + GroupMemoryBarrierWithGroupSync(); frexp_8b3191(); return; } diff --git a/test/intrinsics/gen/frexp/8b3191.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/8b3191.wgsl.expected.spvasm index c398a3e8bc..921aa6a281 100644 --- a/test/intrinsics/gen/frexp/8b3191.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/8b3191.wgsl.expected.spvasm @@ -1,35 +1,53 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 19 +; Bound: 31 ; Schema: 0 OpCapability Shader - %10 = OpExtInstImport "GLSL.std.450" + %12 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_8b3191 "frexp_8b3191" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %arg_1 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %6 = OpTypeFunction %void %float = OpTypeFloat 32 %float_1 = OpConstant %float 1 %_ptr_Function_float = OpTypePointer Function %float - %15 = OpConstantNull %float -%frexp_8b3191 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_float Function %15 - %8 = OpExtInst %float %10 Frexp %float_1 %arg_1 - OpStore %res %8 + %17 = OpConstantNull %float + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %26 = OpConstantNull %uint + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_8b3191 = OpFunction %void None %6 + %9 = OpLabel + %res = OpVariable %_ptr_Function_float Function %17 + %10 = OpExtInst %float %12 Frexp %float_1 %arg_1 + OpStore %res %10 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %17 = OpLabel - %18 = OpFunctionCall %void %frexp_8b3191 +%compute_main = OpFunction %void None %6 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpStore %arg_1 %26 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %30 = OpFunctionCall %void %frexp_8b3191 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl index e96a3e8cb1..5f1a007c54 100644 --- a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_a3f940() { float2 res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = int2(0, 0); + } + GroupMemoryBarrierWithGroupSync(); frexp_a3f940(); return; } diff --git a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm index 71c1378e0d..11801b6871 100644 --- a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm @@ -1,36 +1,55 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 33 ; Schema: 0 OpCapability Shader - %12 = OpExtInstImport "GLSL.std.450" + %15 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_a3f940 "frexp_a3f940" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %v2int = OpTypeVector %int 2 %_ptr_Workgroup_v2int = OpTypePointer Workgroup %v2int %arg_1 = OpVariable %_ptr_Workgroup_v2int Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void + %8 = OpTypeFunction %void %float = OpTypeFloat 32 %v2float = OpTypeVector %float 2 - %13 = OpConstantNull %v2float + %16 = OpConstantNull %v2float %_ptr_Function_v2float = OpTypePointer Function %v2float -%frexp_a3f940 = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v2float Function %13 - %9 = OpExtInst %v2float %12 Frexp %13 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %28 = OpConstantNull %v2int + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_a3f940 = OpFunction %void None %8 + %11 = OpLabel + %res = OpVariable %_ptr_Function_v2float Function %16 + %12 = OpExtInst %v2float %15 Frexp %16 %arg_1 + OpStore %res %12 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %18 = OpLabel - %19 = OpFunctionCall %void %frexp_a3f940 +%compute_main = OpFunction %void None %8 + %21 = OpLabel + %22 = OpLoad %uint %tint_symbol + %24 = OpIEqual %bool %22 %uint_0 + OpSelectionMerge %26 None + OpBranchConditional %24 %27 %26 + %27 = OpLabel + OpStore %arg_1 %28 + OpBranch %26 + %26 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %32 = OpFunctionCall %void %frexp_a3f940 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl index 908e0b893f..a1abbebe60 100644 --- a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_b87f4e() { float4 res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = int4(0, 0, 0, 0); + } + GroupMemoryBarrierWithGroupSync(); frexp_b87f4e(); return; } diff --git a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm index 46bdb0d83d..f8c3322dec 100644 --- a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm @@ -1,36 +1,55 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 33 ; Schema: 0 OpCapability Shader - %12 = OpExtInstImport "GLSL.std.450" + %15 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_b87f4e "frexp_b87f4e" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %v4int = OpTypeVector %int 4 %_ptr_Workgroup_v4int = OpTypePointer Workgroup %v4int %arg_1 = OpVariable %_ptr_Workgroup_v4int Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void + %8 = OpTypeFunction %void %float = OpTypeFloat 32 %v4float = OpTypeVector %float 4 - %13 = OpConstantNull %v4float + %16 = OpConstantNull %v4float %_ptr_Function_v4float = OpTypePointer Function %v4float -%frexp_b87f4e = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v4float Function %13 - %9 = OpExtInst %v4float %12 Frexp %13 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %28 = OpConstantNull %v4int + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_b87f4e = OpFunction %void None %8 + %11 = OpLabel + %res = OpVariable %_ptr_Function_v4float Function %16 + %12 = OpExtInst %v4float %15 Frexp %16 %arg_1 + OpStore %res %12 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %18 = OpLabel - %19 = OpFunctionCall %void %frexp_b87f4e +%compute_main = OpFunction %void None %8 + %21 = OpLabel + %22 = OpLoad %uint %tint_symbol + %24 = OpIEqual %bool %22 %uint_0 + OpSelectionMerge %26 None + OpBranchConditional %24 %27 %26 + %27 = OpLabel + OpStore %arg_1 %28 + OpBranch %26 + %26 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %32 = OpFunctionCall %void %frexp_b87f4e OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.hlsl index 54beba1184..6fe28bbdd3 100644 --- a/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.hlsl +++ b/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.hlsl @@ -7,8 +7,17 @@ void frexp_e8b4d1() { float2 res = tint_tmp_1; } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = uint2(0u, 0u); + } + GroupMemoryBarrierWithGroupSync(); frexp_e8b4d1(); return; } diff --git a/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.spvasm index 1abbb51409..013fb0f4b3 100644 --- a/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.spvasm +++ b/test/intrinsics/gen/frexp/e8b4d1.wgsl.expected.spvasm @@ -1,36 +1,54 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 20 +; Bound: 32 ; Schema: 0 OpCapability Shader - %12 = OpExtInstImport "GLSL.std.450" + %14 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %frexp_e8b4d1 "frexp_e8b4d1" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %uint = OpTypeInt 32 0 %v2uint = OpTypeVector %uint 2 %_ptr_Workgroup_v2uint = OpTypePointer Workgroup %v2uint %arg_1 = OpVariable %_ptr_Workgroup_v2uint Workgroup +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void + %7 = OpTypeFunction %void %float = OpTypeFloat 32 %v2float = OpTypeVector %float 2 - %13 = OpConstantNull %v2float + %15 = OpConstantNull %v2float %_ptr_Function_v2float = OpTypePointer Function %v2float -%frexp_e8b4d1 = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v2float Function %13 - %9 = OpExtInst %v2float %12 Frexp %13 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %27 = OpConstantNull %v2uint + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%frexp_e8b4d1 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_v2float Function %15 + %11 = OpExtInst %v2float %14 Frexp %15 %arg_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %18 = OpLabel - %19 = OpFunctionCall %void %frexp_e8b4d1 +%compute_main = OpFunction %void None %7 + %20 = OpLabel + %21 = OpLoad %uint %tint_symbol + %23 = OpIEqual %bool %21 %uint_0 + OpSelectionMerge %25 None + OpBranchConditional %23 %26 %25 + %26 = OpLabel + OpStore %arg_1 %27 + OpBranch %25 + %25 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %31 = OpFunctionCall %void %frexp_e8b4d1 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl index ef2195e965..5f0f807832 100644 --- a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl +++ b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl @@ -4,8 +4,17 @@ void modf_1d59e5() { float4 res = modf(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1); } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = float4(0.0f, 0.0f, 0.0f, 0.0f); + } + GroupMemoryBarrierWithGroupSync(); modf_1d59e5(); return; } diff --git a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm index 6c82548b00..b9996649b1 100644 --- a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm +++ b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm @@ -1,34 +1,52 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 18 +; Bound: 30 ; Schema: 0 OpCapability Shader - %10 = OpExtInstImport "GLSL.std.450" + %13 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %modf_1d59e5 "modf_1d59e5" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %float = OpTypeFloat 32 %v4float = OpTypeVector %float 4 %_ptr_Workgroup_v4float = OpTypePointer Workgroup %v4float %arg_1 = OpVariable %_ptr_Workgroup_v4float Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void - %11 = OpConstantNull %v4float + %8 = OpTypeFunction %void + %14 = OpConstantNull %v4float %_ptr_Function_v4float = OpTypePointer Function %v4float -%modf_1d59e5 = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v4float Function %11 - %9 = OpExtInst %v4float %10 Modf %11 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%modf_1d59e5 = OpFunction %void None %8 + %11 = OpLabel + %res = OpVariable %_ptr_Function_v4float Function %14 + %12 = OpExtInst %v4float %13 Modf %14 %arg_1 + OpStore %res %12 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %16 = OpLabel - %17 = OpFunctionCall %void %modf_1d59e5 +%compute_main = OpFunction %void None %8 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpStore %arg_1 %14 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %modf_1d59e5 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl b/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl index af8e285203..b900c7bfad 100644 --- a/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl +++ b/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl @@ -4,8 +4,17 @@ void modf_a128ab() { float2 res = modf(float2(0.0f, 0.0f), arg_1); } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = float2(0.0f, 0.0f); + } + GroupMemoryBarrierWithGroupSync(); modf_a128ab(); return; } diff --git a/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm b/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm index dcc8f91dfb..ef64e6d8aa 100644 --- a/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm +++ b/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm @@ -1,34 +1,52 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 18 +; Bound: 30 ; Schema: 0 OpCapability Shader - %10 = OpExtInstImport "GLSL.std.450" + %13 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %modf_a128ab "modf_a128ab" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %float = OpTypeFloat 32 %v2float = OpTypeVector %float 2 %_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float %arg_1 = OpVariable %_ptr_Workgroup_v2float Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void - %11 = OpConstantNull %v2float + %8 = OpTypeFunction %void + %14 = OpConstantNull %v2float %_ptr_Function_v2float = OpTypePointer Function %v2float -%modf_a128ab = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v2float Function %11 - %9 = OpExtInst %v2float %10 Modf %11 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%modf_a128ab = OpFunction %void None %8 + %11 = OpLabel + %res = OpVariable %_ptr_Function_v2float Function %14 + %12 = OpExtInst %v2float %13 Modf %14 %arg_1 + OpStore %res %12 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %16 = OpLabel - %17 = OpFunctionCall %void %modf_a128ab +%compute_main = OpFunction %void None %8 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpStore %arg_1 %14 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %modf_a128ab OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl b/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl index dfd17b4493..22f570df06 100644 --- a/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl +++ b/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl @@ -4,8 +4,17 @@ void modf_bb9088() { float3 res = modf(float3(0.0f, 0.0f, 0.0f), arg_1); } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = float3(0.0f, 0.0f, 0.0f); + } + GroupMemoryBarrierWithGroupSync(); modf_bb9088(); return; } diff --git a/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm b/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm index 3d73465848..df3e8e6434 100644 --- a/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm +++ b/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm @@ -1,34 +1,52 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 18 +; Bound: 30 ; Schema: 0 OpCapability Shader - %10 = OpExtInstImport "GLSL.std.450" + %13 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %modf_bb9088 "modf_bb9088" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %float = OpTypeFloat 32 %v3float = OpTypeVector %float 3 %_ptr_Workgroup_v3float = OpTypePointer Workgroup %v3float %arg_1 = OpVariable %_ptr_Workgroup_v3float Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %5 = OpTypeFunction %void - %11 = OpConstantNull %v3float + %8 = OpTypeFunction %void + %14 = OpConstantNull %v3float %_ptr_Function_v3float = OpTypePointer Function %v3float -%modf_bb9088 = OpFunction %void None %5 - %8 = OpLabel - %res = OpVariable %_ptr_Function_v3float Function %11 - %9 = OpExtInst %v3float %10 Modf %11 %arg_1 - OpStore %res %9 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%modf_bb9088 = OpFunction %void None %8 + %11 = OpLabel + %res = OpVariable %_ptr_Function_v3float Function %14 + %12 = OpExtInst %v3float %13 Modf %14 %arg_1 + OpStore %res %12 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %5 - %16 = OpLabel - %17 = OpFunctionCall %void %modf_bb9088 +%compute_main = OpFunction %void None %8 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpStore %arg_1 %14 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %modf_bb9088 OpReturn OpFunctionEnd diff --git a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl index abfd456d83..726d9b51d0 100644 --- a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl +++ b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl @@ -4,8 +4,17 @@ void modf_e38ae6() { float res = modf(1.0f, arg_1); } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void compute_main() { +void compute_main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + arg_1 = 0.0f; + } + GroupMemoryBarrierWithGroupSync(); modf_e38ae6(); return; } diff --git a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm index af3e129d60..3dfafee29b 100644 --- a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm +++ b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm @@ -1,34 +1,52 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 18 +; Bound: 30 ; Schema: 0 OpCapability Shader - %9 = OpExtInstImport "GLSL.std.450" + %12 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %compute_main "compute_main" + OpEntryPoint GLCompute %compute_main "compute_main" %tint_symbol OpExecutionMode %compute_main LocalSize 1 1 1 OpName %arg_1 "arg_1" + OpName %tint_symbol "tint_symbol" OpName %modf_e38ae6 "modf_e38ae6" OpName %res "res" OpName %compute_main "compute_main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %float = OpTypeFloat 32 %_ptr_Workgroup_float = OpTypePointer Workgroup %float %arg_1 = OpVariable %_ptr_Workgroup_float Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %7 = OpTypeFunction %void %float_1 = OpConstant %float 1 %_ptr_Function_float = OpTypePointer Function %float - %14 = OpConstantNull %float -%modf_e38ae6 = OpFunction %void None %4 - %7 = OpLabel - %res = OpVariable %_ptr_Function_float Function %14 - %8 = OpExtInst %float %9 Modf %float_1 %arg_1 - OpStore %res %8 + %17 = OpConstantNull %float + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 +%modf_e38ae6 = OpFunction %void None %7 + %10 = OpLabel + %res = OpVariable %_ptr_Function_float Function %17 + %11 = OpExtInst %float %12 Modf %float_1 %arg_1 + OpStore %res %11 OpReturn OpFunctionEnd -%compute_main = OpFunction %void None %4 - %16 = OpLabel - %17 = OpFunctionCall %void %modf_e38ae6 +%compute_main = OpFunction %void None %7 + %19 = OpLabel + %20 = OpLoad %uint %tint_symbol + %22 = OpIEqual %bool %20 %uint_0 + OpSelectionMerge %24 None + OpBranchConditional %22 %25 %24 + %25 = OpLabel + OpStore %arg_1 %17 + OpBranch %24 + %24 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %29 = OpFunctionCall %void %modf_e38ae6 OpReturn OpFunctionEnd diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl index c694432345..559d47d14b 100644 --- a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl @@ -1,7 +1,16 @@ groupshared int i; +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void main() { +void main(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + i = 0; + } + GroupMemoryBarrierWithGroupSync(); i = 123; const int use = (i + 1); return; diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl index 4c51f84b40..124f375918 100644 --- a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl @@ -1,10 +1,14 @@ #include using namespace metal; -kernel void tint_symbol() { - threadgroup int tint_symbol_1 = 0; - tint_symbol_1 = 123; - int const use = (tint_symbol_1 + 1); +kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) { + threadgroup int tint_symbol_2; + if ((local_invocation_index == 0u)) { + tint_symbol_2 = int(); + } + threadgroup_barrier(mem_flags::mem_threadgroup); + tint_symbol_2 = 123; + int const use = (tint_symbol_2 + 1); return; } diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm index 6cfbdd43fb..fdefba891c 100644 --- a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm @@ -1,25 +1,44 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 12 +; Bound: 25 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %main "main" + OpEntryPoint GLCompute %main "main" %tint_symbol OpExecutionMode %main LocalSize 1 1 1 OpName %i "i" + OpName %tint_symbol "tint_symbol" OpName %main "main" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %i = OpVariable %_ptr_Workgroup_int Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %4 = OpTypeFunction %void + %7 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %17 = OpConstantNull %int + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 %int_123 = OpConstant %int 123 %int_1 = OpConstant %int 1 - %main = OpFunction %void None %4 - %7 = OpLabel + %main = OpFunction %void None %7 + %10 = OpLabel + %11 = OpLoad %uint %tint_symbol + %13 = OpIEqual %bool %11 %uint_0 + OpSelectionMerge %15 None + OpBranchConditional %13 %16 %15 + %16 = OpLabel + OpStore %i %17 + OpBranch %15 + %15 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 OpStore %i %int_123 - %9 = OpLoad %int %i - %11 = OpIAdd %int %9 %int_1 + %22 = OpLoad %int %i + %24 = OpIAdd %int %22 %int_1 OpReturn OpFunctionEnd diff --git a/test/var/workgroup.wgsl.expected.hlsl b/test/var/workgroup.wgsl.expected.hlsl index bf7fdb3b18..0ae4ba79b9 100644 --- a/test/var/workgroup.wgsl.expected.hlsl +++ b/test/var/workgroup.wgsl.expected.hlsl @@ -25,22 +25,50 @@ void outer() { no_uses(); } +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void main1() { +void main1(tint_symbol_1 tint_symbol) { + const uint local_invocation_index = tint_symbol.local_invocation_index; + if ((local_invocation_index == 0u)) { + a = 0; + } + GroupMemoryBarrierWithGroupSync(); a = 42; uses_a(); return; } +struct tint_symbol_3 { + uint local_invocation_index_1 : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void main2() { +void main2(tint_symbol_3 tint_symbol_2) { + const uint local_invocation_index_1 = tint_symbol_2.local_invocation_index_1; + if ((local_invocation_index_1 == 0u)) { + b = 0; + } + GroupMemoryBarrierWithGroupSync(); b = 7; uses_b(); return; } +struct tint_symbol_5 { + uint local_invocation_index_2 : SV_GroupIndex; +}; + [numthreads(1, 1, 1)] -void main3() { +void main3(tint_symbol_5 tint_symbol_4) { + const uint local_invocation_index_2 = tint_symbol_4.local_invocation_index_2; + if ((local_invocation_index_2 == 0u)) { + a = 0; + b = 0; + } + GroupMemoryBarrierWithGroupSync(); outer(); no_uses(); return; diff --git a/test/var/workgroup.wgsl.expected.msl b/test/var/workgroup.wgsl.expected.msl index 88cd0db963..a0c3800dc0 100644 --- a/test/var/workgroup.wgsl.expected.msl +++ b/test/var/workgroup.wgsl.expected.msl @@ -1,47 +1,60 @@ #include using namespace metal; -void uses_a(threadgroup int* const tint_symbol) { - *(tint_symbol) = (*(tint_symbol) + 1); +void uses_a(threadgroup int* const tint_symbol_3) { + *(tint_symbol_3) = (*(tint_symbol_3) + 1); } -void uses_b(threadgroup int* const tint_symbol_1) { - *(tint_symbol_1) = (*(tint_symbol_1) * 2); +void uses_b(threadgroup int* const tint_symbol_4) { + *(tint_symbol_4) = (*(tint_symbol_4) * 2); } -void uses_a_and_b(threadgroup int* const tint_symbol_2, threadgroup int* const tint_symbol_3) { - *(tint_symbol_2) = *(tint_symbol_3); +void uses_a_and_b(threadgroup int* const tint_symbol_5, threadgroup int* const tint_symbol_6) { + *(tint_symbol_5) = *(tint_symbol_6); } void no_uses() { } -void outer(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) { - *(tint_symbol_4) = 0; - uses_a(tint_symbol_4); - uses_a_and_b(tint_symbol_5, tint_symbol_4); - uses_b(tint_symbol_5); +void outer(threadgroup int* const tint_symbol_7, threadgroup int* const tint_symbol_8) { + *(tint_symbol_7) = 0; + uses_a(tint_symbol_7); + uses_a_and_b(tint_symbol_8, tint_symbol_7); + uses_b(tint_symbol_8); no_uses(); } -kernel void main1() { - threadgroup int tint_symbol_6 = 0; - tint_symbol_6 = 42; - uses_a(&(tint_symbol_6)); +kernel void main1(uint local_invocation_index [[thread_index_in_threadgroup]]) { + threadgroup int tint_symbol_9; + if ((local_invocation_index == 0u)) { + tint_symbol_9 = int(); + } + threadgroup_barrier(mem_flags::mem_threadgroup); + tint_symbol_9 = 42; + uses_a(&(tint_symbol_9)); return; } -kernel void main2() { - threadgroup int tint_symbol_7 = 0; - tint_symbol_7 = 7; - uses_b(&(tint_symbol_7)); +kernel void main2(uint local_invocation_index_1 [[thread_index_in_threadgroup]]) { + threadgroup int tint_symbol_10; + if ((local_invocation_index_1 == 0u)) { + tint_symbol_10 = int(); + } + threadgroup_barrier(mem_flags::mem_threadgroup); + tint_symbol_10 = 7; + uses_b(&(tint_symbol_10)); return; } -kernel void main3() { - threadgroup int tint_symbol_8 = 0; - threadgroup int tint_symbol_9 = 0; - outer(&(tint_symbol_8), &(tint_symbol_9)); +kernel void main3(uint local_invocation_index_2 [[thread_index_in_threadgroup]]) { + threadgroup int tint_symbol_11; + threadgroup int tint_symbol_12; + if ((local_invocation_index_2 == 0u)) { + tint_symbol_11 = int(); + tint_symbol_12 = int(); + } + threadgroup_barrier(mem_flags::mem_threadgroup); + outer(&(tint_symbol_11), &(tint_symbol_12)); no_uses(); return; } diff --git a/test/var/workgroup.wgsl.expected.spvasm b/test/var/workgroup.wgsl.expected.spvasm index 66c1f0ccbf..7007d0e506 100644 --- a/test/var/workgroup.wgsl.expected.spvasm +++ b/test/var/workgroup.wgsl.expected.spvasm @@ -1,13 +1,13 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 45 +; Bound: 70 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %main1 "main1" - OpEntryPoint GLCompute %main2 "main2" - OpEntryPoint GLCompute %main3 "main3" + OpEntryPoint GLCompute %main1 "main1" %tint_symbol + OpEntryPoint GLCompute %main2 "main2" %tint_symbol_1 + OpEntryPoint GLCompute %main3 "main3" %tint_symbol_2 OpEntryPoint GLCompute %main4 "main4" OpExecutionMode %main1 LocalSize 1 1 1 OpExecutionMode %main2 LocalSize 1 1 1 @@ -16,6 +16,9 @@ OpName %a "a" OpName %b "b" OpName %c "c" + OpName %tint_symbol "tint_symbol" + OpName %tint_symbol_1 "tint_symbol_1" + OpName %tint_symbol_2 "tint_symbol_2" OpName %uses_a "uses_a" OpName %uses_b "uses_b" OpName %uses_a_and_b "uses_a_and_b" @@ -25,71 +28,112 @@ OpName %main2 "main2" OpName %main3 "main3" OpName %main4 "main4" + OpDecorate %tint_symbol BuiltIn LocalInvocationIndex + OpDecorate %tint_symbol_1 BuiltIn LocalInvocationIndex + OpDecorate %tint_symbol_2 BuiltIn LocalInvocationIndex %int = OpTypeInt 32 1 %_ptr_Workgroup_int = OpTypePointer Workgroup %int %a = OpVariable %_ptr_Workgroup_int Workgroup %b = OpVariable %_ptr_Workgroup_int Workgroup %c = OpVariable %_ptr_Workgroup_int Workgroup + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%tint_symbol = OpVariable %_ptr_Input_uint Input +%tint_symbol_1 = OpVariable %_ptr_Input_uint Input +%tint_symbol_2 = OpVariable %_ptr_Input_uint Input %void = OpTypeVoid - %6 = OpTypeFunction %void + %11 = OpTypeFunction %void %int_1 = OpConstant %int 1 %int_2 = OpConstant %int 2 %int_0 = OpConstant %int 0 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %43 = OpConstantNull %int + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 %int_42 = OpConstant %int 42 %int_7 = OpConstant %int 7 - %uses_a = OpFunction %void None %6 - %9 = OpLabel - %10 = OpLoad %int %a - %12 = OpIAdd %int %10 %int_1 - OpStore %a %12 - OpReturn - OpFunctionEnd - %uses_b = OpFunction %void None %6 + %uses_a = OpFunction %void None %11 %14 = OpLabel - %15 = OpLoad %int %b - %17 = OpIMul %int %15 %int_2 - OpStore %b %17 + %15 = OpLoad %int %a + %17 = OpIAdd %int %15 %int_1 + OpStore %a %17 OpReturn OpFunctionEnd -%uses_a_and_b = OpFunction %void None %6 + %uses_b = OpFunction %void None %11 %19 = OpLabel - %20 = OpLoad %int %a - OpStore %b %20 + %20 = OpLoad %int %b + %22 = OpIMul %int %20 %int_2 + OpStore %b %22 OpReturn OpFunctionEnd - %no_uses = OpFunction %void None %6 - %22 = OpLabel - OpReturn - OpFunctionEnd - %outer = OpFunction %void None %6 +%uses_a_and_b = OpFunction %void None %11 %24 = OpLabel + %25 = OpLoad %int %a + OpStore %b %25 + OpReturn + OpFunctionEnd + %no_uses = OpFunction %void None %11 + %27 = OpLabel + OpReturn + OpFunctionEnd + %outer = OpFunction %void None %11 + %29 = OpLabel OpStore %a %int_0 - %26 = OpFunctionCall %void %uses_a - %27 = OpFunctionCall %void %uses_a_and_b - %28 = OpFunctionCall %void %uses_b - %29 = OpFunctionCall %void %no_uses + %31 = OpFunctionCall %void %uses_a + %32 = OpFunctionCall %void %uses_a_and_b + %33 = OpFunctionCall %void %uses_b + %34 = OpFunctionCall %void %no_uses OpReturn OpFunctionEnd - %main1 = OpFunction %void None %6 - %31 = OpLabel + %main1 = OpFunction %void None %11 + %36 = OpLabel + %37 = OpLoad %uint %tint_symbol + %39 = OpIEqual %bool %37 %uint_0 + OpSelectionMerge %41 None + OpBranchConditional %39 %42 %41 + %42 = OpLabel + OpStore %a %43 + OpBranch %41 + %41 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 OpStore %a %int_42 - %33 = OpFunctionCall %void %uses_a + %48 = OpFunctionCall %void %uses_a OpReturn OpFunctionEnd - %main2 = OpFunction %void None %6 - %35 = OpLabel + %main2 = OpFunction %void None %11 + %50 = OpLabel + %51 = OpLoad %uint %tint_symbol_1 + %52 = OpIEqual %bool %51 %uint_0 + OpSelectionMerge %53 None + OpBranchConditional %52 %54 %53 + %54 = OpLabel + OpStore %b %43 + OpBranch %53 + %53 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 OpStore %b %int_7 - %37 = OpFunctionCall %void %uses_b + %57 = OpFunctionCall %void %uses_b OpReturn OpFunctionEnd - %main3 = OpFunction %void None %6 - %39 = OpLabel - %40 = OpFunctionCall %void %outer - %41 = OpFunctionCall %void %no_uses + %main3 = OpFunction %void None %11 + %59 = OpLabel + %60 = OpLoad %uint %tint_symbol_2 + %61 = OpIEqual %bool %60 %uint_0 + OpSelectionMerge %62 None + OpBranchConditional %61 %63 %62 + %63 = OpLabel + OpStore %a %43 + OpStore %b %43 + OpBranch %62 + %62 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %65 = OpFunctionCall %void %outer + %66 = OpFunctionCall %void %no_uses OpReturn OpFunctionEnd - %main4 = OpFunction %void None %6 - %43 = OpLabel - %44 = OpFunctionCall %void %no_uses + %main4 = OpFunction %void None %11 + %68 = OpLabel + %69 = OpFunctionCall %void %no_uses OpReturn OpFunctionEnd