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