diff --git a/src/val/hlsl.cc b/src/val/hlsl.cc index c1c3315a05..e7e3aea777 100644 --- a/src/val/hlsl.cc +++ b/src/val/hlsl.cc @@ -20,9 +20,9 @@ #include "src/utils/io/tmpfile.h" #ifdef _WIN32 +#include #include #include -#include #include using Microsoft::WRL::ComPtr; diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 9e81adb5da..bea22ce0bb 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -26,19 +27,23 @@ #include "src/ast/interpolate_decoration.h" #include "src/ast/override_decoration.h" #include "src/ast/variable_decl_statement.h" +#include "src/debug.h" #include "src/sem/array.h" #include "src/sem/atomic_type.h" +#include "src/sem/block_statement.h" #include "src/sem/call.h" #include "src/sem/depth_texture_type.h" #include "src/sem/function.h" #include "src/sem/member_accessor_expression.h" #include "src/sem/multisampled_texture_type.h" #include "src/sem/sampled_texture_type.h" +#include "src/sem/statement.h" #include "src/sem/storage_texture_type.h" #include "src/sem/struct.h" #include "src/sem/variable.h" #include "src/transform/calculate_array_length.h" #include "src/transform/hlsl.h" +#include "src/utils/get_or_create.h" #include "src/utils/scoped_assignment.h" #include "src/writer/append_vector.h" #include "src/writer/float_to_string.h" @@ -120,6 +125,10 @@ bool GeneratorImpl::Generate() { const TypeInfo* last_kind = nullptr; std::streampos last_padding_pos; + if (!FindAndEmitVectorAssignmentInLoopFunctions()) { + return false; + } + for (auto* decl : builder_.AST().GlobalDeclarations()) { if (decl->Is()) { continue; // Ignore aliases. @@ -164,6 +173,112 @@ bool GeneratorImpl::Generate() { return true; } +bool GeneratorImpl::FindAndEmitVectorAssignmentInLoopFunctions() { + auto is_in_loop = [](const sem::Expression* expr) { + auto* block = expr->Stmt()->Block(); + if (!block) { + return false; + } + return block->FindFirstParent() != nullptr; + }; + + auto emit_function_once = [&](const sem::Vector* vec) { + utils::GetOrCreate(vector_assignment_in_loop_funcs_, vec, [&] { + std::ostringstream ss; + EmitType(ss, vec, tint::ast::StorageClass::kInvalid, + ast::Access::kUndefined, ""); + auto func_name = UniqueIdentifier("Set_" + ss.str()); + { + auto out = line(); + out << "void " << func_name << "(inout "; + EmitType(out, vec, ast::StorageClass::kInvalid, ast::Access::kUndefined, + ""); + out << " vec, int idx, "; + EmitType(out, vec->type(), ast::StorageClass::kInvalid, + ast::Access::kUndefined, ""); + out << " val) {"; + } + { + ScopedIndent si(this); + line() << "switch(idx) {"; + { + ScopedIndent si2(this); + for (size_t i = 0; i < vec->size(); ++i) { + auto sidx = std::to_string(i); + line() << "case " + sidx + ": vec[" + sidx + "] = val; break;"; + } + } + line() << "}"; + } + line() << "}"; + return func_name; + }); + }; + + // Find vector assignments via an accessor expression (index) within loops so + // that we can replace them later with calls to setter functions. Also emit + // the setter functions per vector type as we find them. We do this to avoid + // having FCX fail to unroll loops with "error X3511: forced to unroll loop, + // but unrolling failed." See crbug.com/tint/534. + + for (auto* ast_node : program_->ASTNodes().Objects()) { + auto* ast_assign = ast_node->As(); + if (!ast_assign) { + continue; + } + + auto* ast_access_expr = + ast_assign->lhs()->As(); + if (!ast_access_expr) { + continue; + } + + auto* array_expr = builder_.Sem().Get(ast_access_expr->array()); + auto* vec = array_expr->Type()->UnwrapRef()->As(); + + // Skip non-vectors + if (!vec) { + continue; + } + + // Skip if not part of a loop + if (!is_in_loop(array_expr)) { + continue; + } + + // Save this assignment along with the vector type + vector_assignments_in_loops_.emplace(ast_assign, vec); + + // Emit the function if it hasn't already + emit_function_once(vec); + } + + return true; +} + +bool GeneratorImpl::EmitVectorAssignmentInLoopCall( + const ast::AssignmentStatement* stmt, + const sem::Vector* vec) { + auto* ast_access_expr = stmt->lhs()->As(); + + auto out = line(); + out << vector_assignment_in_loop_funcs_.at(vec) << "("; + if (!EmitExpression(out, ast_access_expr->array())) { + return false; + } + out << ", "; + if (!EmitExpression(out, ast_access_expr->idx_expr())) { + return false; + } + out << ", "; + if (!EmitExpression(out, stmt->rhs())) { + return false; + } + out << ");"; + + return true; +} + bool GeneratorImpl::EmitArrayAccessor(std::ostream& out, ast::ArrayAccessorExpression* expr) { if (!EmitExpression(out, expr->array())) { @@ -202,6 +317,11 @@ bool GeneratorImpl::EmitBitcast(std::ostream& out, } bool GeneratorImpl::EmitAssign(ast::AssignmentStatement* stmt) { + auto iter = vector_assignments_in_loops_.find(stmt); + if (iter != vector_assignments_in_loops_.end()) { + return EmitVectorAssignmentInLoopCall(iter->first, iter->second); + } + auto out = line(); if (!EmitExpression(out, stmt->lhs())) { return false; diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index 58b791df96..61d09dae7c 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/src/writer/hlsl/generator_impl.h @@ -344,6 +344,21 @@ class GeneratorImpl : public TextGenerator { /// @returns true if the variable was emitted bool EmitProgramConstVariable(const ast::Variable* var); + /// Finds vector assignments via an accessor expression within loops, storing + /// the assignment/vector node pair in `vector_assignments_in_loops`, and + /// emits function definitions per vector type found. Required to work around + /// an FXC bug, see crbug.com/tint/534. + /// @returns true on success + bool FindAndEmitVectorAssignmentInLoopFunctions(); + /// Emits call to vector assignment function for the input assignment + /// statement and vector type. + /// @param stmt assignment statement that corresponds to a vector assingment + /// via an accessor expression + /// @param vec the vector type being assigned to + /// @returns true on success + bool EmitVectorAssignmentInLoopCall(const ast::AssignmentStatement* stmt, + const sem::Vector* vec); + /// Handles generating a builtin method name /// @param intrinsic the semantic info for the intrinsic /// @returns the name or "" if not valid @@ -373,6 +388,10 @@ class GeneratorImpl : public TextGenerator { std::function emit_continuing_; std::unordered_map structure_builders_; + std::unordered_map + vector_assignments_in_loops_; + std::unordered_map + vector_assignment_in_loop_funcs_; }; } // namespace hlsl diff --git a/test/bug/tint/534.wgsl b/test/bug/tint/534.wgsl new file mode 100644 index 0000000000..56e8e0a08d --- /dev/null +++ b/test/bug/tint/534.wgsl @@ -0,0 +1,48 @@ +[[block]] struct Uniforms { + dstTextureFlipY : u32; + isFloat16 : u32; + isRGB10A2Unorm : u32; + channelCount : u32; +}; +[[block]] struct OutputBuf { + result : [[stride(4)]] array; +}; +[[group(0), binding(0)]] var src : texture_2d; +[[group(0), binding(1)]] var dst : texture_2d; +[[group(0), binding(2)]] var output : OutputBuf; +[[group(0), binding(3)]] var uniforms : Uniforms; +//[[builtin(global_invocation_id)]] var GlobalInvocationID : vec3; +// Fp16 logic +// Infinity and NaN won't happen in this test case. +fn ConvertToFp16FloatValue(fp32 : f32) -> u32 { + return 1u; +} + +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { + var size : vec2 = textureDimensions(src); + var dstTexCoord : vec2 = vec2(GlobalInvocationID.xy); + var srcTexCoord : vec2 = dstTexCoord; + if (uniforms.dstTextureFlipY == 1u) { + srcTexCoord.y = size.y - dstTexCoord.y - 1; + } + + var srcColor : vec4 = textureLoad(src, srcTexCoord, 0); + var dstColor : vec4 = textureLoad(dst, dstTexCoord, 0); + var success : bool = true; + + var srcColorBits : vec4; + var dstColorBits : vec4 = vec4(dstColor); + + for (var i : u32 = 0u; i < uniforms.channelCount; i = i + 1u) { + srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]); + success = success && (srcColorBits[i] == dstColorBits[i]); + } + + var outputIndex : u32 = GlobalInvocationID.y * u32(size.x) + GlobalInvocationID.x; + if (success) { + output.result[outputIndex] = u32(1); + } else { + output.result[outputIndex] = u32(0); + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl new file mode 100644 index 0000000000..8dc751527a --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl @@ -0,0 +1,21 @@ +var v2f : vec2; +var v3i : vec3; +var v4u : vec4; +var v2b : vec2; + +fn foo() { + for (var i : i32 = 0; i < 2; i = i + 1) { + v2f[i] = 1.0; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; + } +} + +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + + for (var i : i32 = 0; i < 2; i = i + 1) { + foo(); + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl new file mode 100644 index 0000000000..5b47c06610 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl @@ -0,0 +1,66 @@ +void Set_float2(inout float2 vec, int idx, float val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_int3(inout int3 vec, int idx, int val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + } +} +void Set_uint4(inout uint4 vec, int idx, uint val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + case 3: vec[3] = val; break; + } +} +void Set_bool2(inout bool2 vec, int idx, bool val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +static float2 v2f = float2(0.0f, 0.0f); +static int3 v3i = int3(0, 0, 0); +static uint4 v4u = uint4(0u, 0u, 0u, 0u); +static bool2 v2b = bool2(false, false); + +void foo() { + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + Set_float2(v2f, i, 1.0f); + Set_int3(v3i, i, 1); + Set_uint4(v4u, i, 1u); + Set_bool2(v2b, i, true); + { + i = (i + 1); + } + } + } +} + +[numthreads(1, 1, 1)] +void main() { + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + foo(); + { + i = (i + 1); + } + } + } + return; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl new file mode 100644 index 0000000000..ab45d31ac8 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl @@ -0,0 +1,41 @@ +#include + +using namespace metal; +void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) { + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + (*(tint_symbol_1))[i] = 1.0f; + (*(tint_symbol_2))[i] = 1; + (*(tint_symbol_3))[i] = 1u; + (*(tint_symbol_4))[i] = true; + { + i = (i + 1); + } + } + } +} + +kernel void tint_symbol() { + thread float2 tint_symbol_5 = 0.0f; + thread int3 tint_symbol_6 = 0; + thread uint4 tint_symbol_7 = 0u; + thread bool2 tint_symbol_8 = false; + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8)); + { + i = (i + 1); + } + } + } + return; +} + diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm new file mode 100644 index 0000000000..212a926008 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm @@ -0,0 +1,116 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 72 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %v2f "v2f" + OpName %v3i "v3i" + OpName %v4u "v4u" + OpName %v2b "v2b" + OpName %foo "foo" + OpName %i "i" + OpName %main "main" + OpName %i_0 "i" + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 +%_ptr_Private_v2float = OpTypePointer Private %v2float + %5 = OpConstantNull %v2float + %v2f = OpVariable %_ptr_Private_v2float Private %5 + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 +%_ptr_Private_v3int = OpTypePointer Private %v3int + %10 = OpConstantNull %v3int + %v3i = OpVariable %_ptr_Private_v3int Private %10 + %uint = OpTypeInt 32 0 + %v4uint = OpTypeVector %uint 4 +%_ptr_Private_v4uint = OpTypePointer Private %v4uint + %15 = OpConstantNull %v4uint + %v4u = OpVariable %_ptr_Private_v4uint Private %15 + %bool = OpTypeBool + %v2bool = OpTypeVector %bool 2 +%_ptr_Private_v2bool = OpTypePointer Private %v2bool + %20 = OpConstantNull %v2bool + %v2b = OpVariable %_ptr_Private_v2bool Private %20 + %void = OpTypeVoid + %21 = OpTypeFunction %void + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %28 = OpConstantNull %int + %int_2 = OpConstant %int 2 +%_ptr_Private_float = OpTypePointer Private %float + %float_1 = OpConstant %float 1 +%_ptr_Private_int = OpTypePointer Private %int + %int_1 = OpConstant %int 1 +%_ptr_Private_uint = OpTypePointer Private %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Private_bool = OpTypePointer Private %bool + %true = OpConstantTrue %bool + %foo = OpFunction %void None %21 + %24 = OpLabel + %i = OpVariable %_ptr_Function_int Function %28 + OpStore %i %int_0 + OpBranch %29 + %29 = OpLabel + OpLoopMerge %30 %31 None + OpBranch %32 + %32 = OpLabel + %34 = OpLoad %int %i + %36 = OpSLessThan %bool %34 %int_2 + %33 = OpLogicalNot %bool %36 + OpSelectionMerge %37 None + OpBranchConditional %33 %38 %37 + %38 = OpLabel + OpBranch %30 + %37 = OpLabel + %39 = OpLoad %int %i + %41 = OpAccessChain %_ptr_Private_float %v2f %39 + OpStore %41 %float_1 + %43 = OpLoad %int %i + %45 = OpAccessChain %_ptr_Private_int %v3i %43 + OpStore %45 %int_1 + %47 = OpLoad %int %i + %49 = OpAccessChain %_ptr_Private_uint %v4u %47 + OpStore %49 %uint_1 + %51 = OpLoad %int %i + %53 = OpAccessChain %_ptr_Private_bool %v2b %51 + OpStore %53 %true + OpBranch %31 + %31 = OpLabel + %55 = OpLoad %int %i + %56 = OpIAdd %int %55 %int_1 + OpStore %i %56 + OpBranch %29 + %30 = OpLabel + OpReturn + OpFunctionEnd + %main = OpFunction %void None %21 + %58 = OpLabel + %i_0 = OpVariable %_ptr_Function_int Function %28 + OpStore %i_0 %int_0 + OpBranch %60 + %60 = OpLabel + OpLoopMerge %61 %62 None + OpBranch %63 + %63 = OpLabel + %65 = OpLoad %int %i_0 + %66 = OpSLessThan %bool %65 %int_2 + %64 = OpLogicalNot %bool %66 + OpSelectionMerge %67 None + OpBranchConditional %64 %68 %67 + %68 = OpLabel + OpBranch %61 + %67 = OpLabel + %69 = OpFunctionCall %void %foo + OpBranch %62 + %62 = OpLabel + %70 = OpLoad %int %i_0 + %71 = OpIAdd %int %70 %int_1 + OpStore %i_0 %71 + OpBranch %60 + %61 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl new file mode 100644 index 0000000000..bc302f118c --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl @@ -0,0 +1,43 @@ +var v2f : vec2; + +var v3i : vec3; + +var v4u : vec4; + +var v2b : vec2; + +fn foo() { + { + var i : i32 = 0; + loop { + if (!((i < 2))) { + break; + } + v2f[i] = 1.0; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; + + continuing { + i = (i + 1); + } + } + } +} + +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + { + var i : i32 = 0; + loop { + if (!((i < 2))) { + break; + } + foo(); + + continuing { + i = (i + 1); + } + } + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl new file mode 100644 index 0000000000..418155d2c5 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl @@ -0,0 +1,20 @@ +var v2f : vec2; +var v3i : vec3; +var v4u : vec4; +var v2b : vec2; + +fn foo() { + var i : i32 = 0; + v2f[i] = 1.0; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; +} + +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + + for (var i : i32 = 0; i < 2; i = i + 1) { + foo(); + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl new file mode 100644 index 0000000000..6e34e5613a --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl @@ -0,0 +1,29 @@ +static float2 v2f = float2(0.0f, 0.0f); +static int3 v3i = int3(0, 0, 0); +static uint4 v4u = uint4(0u, 0u, 0u, 0u); +static bool2 v2b = bool2(false, false); + +void foo() { + int i = 0; + v2f[i] = 1.0f; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; +} + +[numthreads(1, 1, 1)] +void main() { + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + foo(); + { + i = (i + 1); + } + } + } + return; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl new file mode 100644 index 0000000000..2685e75daa --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl @@ -0,0 +1,31 @@ +#include + +using namespace metal; +void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) { + int i = 0; + (*(tint_symbol_1))[i] = 1.0f; + (*(tint_symbol_2))[i] = 1; + (*(tint_symbol_3))[i] = 1u; + (*(tint_symbol_4))[i] = true; +} + +kernel void tint_symbol() { + thread float2 tint_symbol_5 = 0.0f; + thread int3 tint_symbol_6 = 0; + thread uint4 tint_symbol_7 = 0u; + thread bool2 tint_symbol_8 = false; + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8)); + { + i = (i + 1); + } + } + } + return; +} + diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm new file mode 100644 index 0000000000..d0b731b01a --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm @@ -0,0 +1,96 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 61 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %v2f "v2f" + OpName %v3i "v3i" + OpName %v4u "v4u" + OpName %v2b "v2b" + OpName %foo "foo" + OpName %i "i" + OpName %main "main" + OpName %i_0 "i" + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 +%_ptr_Private_v2float = OpTypePointer Private %v2float + %5 = OpConstantNull %v2float + %v2f = OpVariable %_ptr_Private_v2float Private %5 + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 +%_ptr_Private_v3int = OpTypePointer Private %v3int + %10 = OpConstantNull %v3int + %v3i = OpVariable %_ptr_Private_v3int Private %10 + %uint = OpTypeInt 32 0 + %v4uint = OpTypeVector %uint 4 +%_ptr_Private_v4uint = OpTypePointer Private %v4uint + %15 = OpConstantNull %v4uint + %v4u = OpVariable %_ptr_Private_v4uint Private %15 + %bool = OpTypeBool + %v2bool = OpTypeVector %bool 2 +%_ptr_Private_v2bool = OpTypePointer Private %v2bool + %20 = OpConstantNull %v2bool + %v2b = OpVariable %_ptr_Private_v2bool Private %20 + %void = OpTypeVoid + %21 = OpTypeFunction %void + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %28 = OpConstantNull %int +%_ptr_Private_float = OpTypePointer Private %float + %float_1 = OpConstant %float 1 +%_ptr_Private_int = OpTypePointer Private %int + %int_1 = OpConstant %int 1 +%_ptr_Private_uint = OpTypePointer Private %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Private_bool = OpTypePointer Private %bool + %true = OpConstantTrue %bool + %int_2 = OpConstant %int 2 + %foo = OpFunction %void None %21 + %24 = OpLabel + %i = OpVariable %_ptr_Function_int Function %28 + OpStore %i %int_0 + %29 = OpLoad %int %i + %31 = OpAccessChain %_ptr_Private_float %v2f %29 + OpStore %31 %float_1 + %33 = OpLoad %int %i + %35 = OpAccessChain %_ptr_Private_int %v3i %33 + OpStore %35 %int_1 + %37 = OpLoad %int %i + %39 = OpAccessChain %_ptr_Private_uint %v4u %37 + OpStore %39 %uint_1 + %41 = OpLoad %int %i + %43 = OpAccessChain %_ptr_Private_bool %v2b %41 + OpStore %43 %true + OpReturn + OpFunctionEnd + %main = OpFunction %void None %21 + %46 = OpLabel + %i_0 = OpVariable %_ptr_Function_int Function %28 + OpStore %i_0 %int_0 + OpBranch %48 + %48 = OpLabel + OpLoopMerge %49 %50 None + OpBranch %51 + %51 = OpLabel + %53 = OpLoad %int %i_0 + %55 = OpSLessThan %bool %53 %int_2 + %52 = OpLogicalNot %bool %55 + OpSelectionMerge %56 None + OpBranchConditional %52 %57 %56 + %57 = OpLabel + OpBranch %49 + %56 = OpLabel + %58 = OpFunctionCall %void %foo + OpBranch %50 + %50 = OpLabel + %59 = OpLoad %int %i_0 + %60 = OpIAdd %int %59 %int_1 + OpStore %i_0 %60 + OpBranch %48 + %49 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl new file mode 100644 index 0000000000..ae1667c454 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl @@ -0,0 +1,32 @@ +var v2f : vec2; + +var v3i : vec3; + +var v4u : vec4; + +var v2b : vec2; + +fn foo() { + var i : i32 = 0; + v2f[i] = 1.0; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; +} + +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + { + var i : i32 = 0; + loop { + if (!((i < 2))) { + break; + } + foo(); + + continuing { + i = (i + 1); + } + } + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl new file mode 100644 index 0000000000..ef2ec78d6a --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl @@ -0,0 +1,30 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + + for (var i : i32 = 0; i < 2; i = i + 1) { + v2f[i] = 1.0; + v3f[i] = 1.0; + v4f[i] = 1.0; + v2i[i] = 1; + v3i[i] = 1; + v4i[i] = 1; + v2u[i] = 1u; + v3u[i] = 1u; + v4u[i] = 1u; + v2b[i] = true; + v3b[i] = true; + v4b[i] = true; + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl new file mode 100644 index 0000000000..3600e494db --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl @@ -0,0 +1,123 @@ +void Set_float2(inout float2 vec, int idx, float val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_float3(inout float3 vec, int idx, float val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + } +} +void Set_float4(inout float4 vec, int idx, float val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + case 3: vec[3] = val; break; + } +} +void Set_int2(inout int2 vec, int idx, int val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_int3(inout int3 vec, int idx, int val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + } +} +void Set_int4(inout int4 vec, int idx, int val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + case 3: vec[3] = val; break; + } +} +void Set_uint2(inout uint2 vec, int idx, uint val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_uint3(inout uint3 vec, int idx, uint val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + } +} +void Set_uint4(inout uint4 vec, int idx, uint val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + case 3: vec[3] = val; break; + } +} +void Set_bool2(inout bool2 vec, int idx, bool val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_bool3(inout bool3 vec, int idx, bool val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + } +} +void Set_bool4(inout bool4 vec, int idx, bool val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + case 3: vec[3] = val; break; + } +} +[numthreads(1, 1, 1)] +void main() { + float2 v2f = float2(0.0f, 0.0f); + float3 v3f = float3(0.0f, 0.0f, 0.0f); + float4 v4f = float4(0.0f, 0.0f, 0.0f, 0.0f); + int2 v2i = int2(0, 0); + int3 v3i = int3(0, 0, 0); + int4 v4i = int4(0, 0, 0, 0); + uint2 v2u = uint2(0u, 0u); + uint3 v3u = uint3(0u, 0u, 0u); + uint4 v4u = uint4(0u, 0u, 0u, 0u); + bool2 v2b = bool2(false, false); + bool3 v3b = bool3(false, false, false); + bool4 v4b = bool4(false, false, false, false); + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + Set_float2(v2f, i, 1.0f); + Set_float3(v3f, i, 1.0f); + Set_float4(v4f, i, 1.0f); + Set_int2(v2i, i, 1); + Set_int3(v3i, i, 1); + Set_int4(v4i, i, 1); + Set_uint2(v2u, i, 1u); + Set_uint3(v3u, i, 1u); + Set_uint4(v4u, i, 1u); + Set_bool2(v2b, i, true); + Set_bool3(v3b, i, true); + Set_bool4(v4b, i, true); + { + i = (i + 1); + } + } + } + return; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl new file mode 100644 index 0000000000..a460250712 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl @@ -0,0 +1,42 @@ +#include + +using namespace metal; +kernel void tint_symbol() { + float2 v2f = 0.0f; + float3 v3f = 0.0f; + float4 v4f = 0.0f; + int2 v2i = 0; + int3 v3i = 0; + int4 v4i = 0; + uint2 v2u = 0u; + uint3 v3u = 0u; + uint4 v4u = 0u; + bool2 v2b = false; + bool3 v3b = false; + bool4 v4b = false; + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + v2f[i] = 1.0f; + v3f[i] = 1.0f; + v4f[i] = 1.0f; + v2i[i] = 1; + v3i[i] = 1; + v4i[i] = 1; + v2u[i] = 1u; + v3u[i] = 1u; + v4u[i] = 1u; + v2b[i] = true; + v3b[i] = true; + v4b[i] = true; + { + i = (i + 1); + } + } + } + return; +} + diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm new file mode 100644 index 0000000000..5a80cc6b72 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm @@ -0,0 +1,150 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 104 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %main "main" + OpName %v2f "v2f" + OpName %v3f "v3f" + OpName %v4f "v4f" + OpName %v2i "v2i" + OpName %v3i "v3i" + OpName %v4i "v4i" + OpName %v2u "v2u" + OpName %v3u "v3u" + OpName %v4u "v4u" + OpName %v2b "v2b" + OpName %v3b "v3b" + OpName %v4b "v4b" + OpName %i "i" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %9 = OpConstantNull %v2float + %v3float = OpTypeVector %float 3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %13 = OpConstantNull %v3float + %v4float = OpTypeVector %float 4 +%_ptr_Function_v4float = OpTypePointer Function %v4float + %17 = OpConstantNull %v4float + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 +%_ptr_Function_v2int = OpTypePointer Function %v2int + %22 = OpConstantNull %v2int + %v3int = OpTypeVector %int 3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %26 = OpConstantNull %v3int + %v4int = OpTypeVector %int 4 +%_ptr_Function_v4int = OpTypePointer Function %v4int + %30 = OpConstantNull %v4int + %uint = OpTypeInt 32 0 + %v2uint = OpTypeVector %uint 2 +%_ptr_Function_v2uint = OpTypePointer Function %v2uint + %35 = OpConstantNull %v2uint + %v3uint = OpTypeVector %uint 3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %39 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %43 = OpConstantNull %v4uint + %bool = OpTypeBool + %v2bool = OpTypeVector %bool 2 +%_ptr_Function_v2bool = OpTypePointer Function %v2bool + %48 = OpConstantNull %v2bool + %v3bool = OpTypeVector %bool 3 +%_ptr_Function_v3bool = OpTypePointer Function %v3bool + %52 = OpConstantNull %v3bool + %v4bool = OpTypeVector %bool 4 +%_ptr_Function_v4bool = OpTypePointer Function %v4bool + %56 = OpConstantNull %v4bool + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %60 = OpConstantNull %int + %int_2 = OpConstant %int 2 +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %int_1 = OpConstant %int 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Function_bool = OpTypePointer Function %bool + %true = OpConstantTrue %bool + %main = OpFunction %void None %1 + %4 = OpLabel + %v2f = OpVariable %_ptr_Function_v2float Function %9 + %v3f = OpVariable %_ptr_Function_v3float Function %13 + %v4f = OpVariable %_ptr_Function_v4float Function %17 + %v2i = OpVariable %_ptr_Function_v2int Function %22 + %v3i = OpVariable %_ptr_Function_v3int Function %26 + %v4i = OpVariable %_ptr_Function_v4int Function %30 + %v2u = OpVariable %_ptr_Function_v2uint Function %35 + %v3u = OpVariable %_ptr_Function_v3uint Function %39 + %v4u = OpVariable %_ptr_Function_v4uint Function %43 + %v2b = OpVariable %_ptr_Function_v2bool Function %48 + %v3b = OpVariable %_ptr_Function_v3bool Function %52 + %v4b = OpVariable %_ptr_Function_v4bool Function %56 + %i = OpVariable %_ptr_Function_int Function %60 + OpStore %i %int_0 + OpBranch %61 + %61 = OpLabel + OpLoopMerge %62 %63 None + OpBranch %64 + %64 = OpLabel + %66 = OpLoad %int %i + %68 = OpSLessThan %bool %66 %int_2 + %65 = OpLogicalNot %bool %68 + OpSelectionMerge %69 None + OpBranchConditional %65 %70 %69 + %70 = OpLabel + OpBranch %62 + %69 = OpLabel + %71 = OpLoad %int %i + %73 = OpAccessChain %_ptr_Function_float %v2f %71 + OpStore %73 %float_1 + %75 = OpLoad %int %i + %76 = OpAccessChain %_ptr_Function_float %v3f %75 + OpStore %76 %float_1 + %77 = OpLoad %int %i + %78 = OpAccessChain %_ptr_Function_float %v4f %77 + OpStore %78 %float_1 + %79 = OpLoad %int %i + %80 = OpAccessChain %_ptr_Function_int %v2i %79 + OpStore %80 %int_1 + %82 = OpLoad %int %i + %83 = OpAccessChain %_ptr_Function_int %v3i %82 + OpStore %83 %int_1 + %84 = OpLoad %int %i + %85 = OpAccessChain %_ptr_Function_int %v4i %84 + OpStore %85 %int_1 + %86 = OpLoad %int %i + %88 = OpAccessChain %_ptr_Function_uint %v2u %86 + OpStore %88 %uint_1 + %90 = OpLoad %int %i + %91 = OpAccessChain %_ptr_Function_uint %v3u %90 + OpStore %91 %uint_1 + %92 = OpLoad %int %i + %93 = OpAccessChain %_ptr_Function_uint %v4u %92 + OpStore %93 %uint_1 + %94 = OpLoad %int %i + %96 = OpAccessChain %_ptr_Function_bool %v2b %94 + OpStore %96 %true + %98 = OpLoad %int %i + %99 = OpAccessChain %_ptr_Function_bool %v3b %98 + OpStore %99 %true + %100 = OpLoad %int %i + %101 = OpAccessChain %_ptr_Function_bool %v4b %100 + OpStore %101 %true + OpBranch %63 + %63 = OpLabel + %102 = OpLoad %int %i + %103 = OpIAdd %int %102 %int_1 + OpStore %i %103 + OpBranch %61 + %62 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl new file mode 100644 index 0000000000..7d736281bb --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl @@ -0,0 +1,39 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + { + var i : i32 = 0; + loop { + if (!((i < 2))) { + break; + } + v2f[i] = 1.0; + v3f[i] = 1.0; + v4f[i] = 1.0; + v2i[i] = 1; + v3i[i] = 1; + v4i[i] = 1; + v2u[i] = 1u; + v3u[i] = 1u; + v4u[i] = 1u; + v2b[i] = true; + v3b[i] = true; + v4b[i] = true; + + continuing { + i = (i + 1); + } + } + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl new file mode 100644 index 0000000000..d464d4c8d1 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl @@ -0,0 +1,26 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v2f_2 : vec2; + + var v3i : vec3; + var v3i_2 : vec3; + + var v4u : vec4; + var v4u_2 : vec4; + + var v2b : vec2; + var v2b_2 : vec2; + + for (var i : i32 = 0; i < 2; i = i + 1) { + v2f[i] = 1.0; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; + + v2f_2[i] = 1.0; + v3i_2[i] = 1; + v4u_2[i] = 1u; + v2b_2[i] = true; + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl new file mode 100644 index 0000000000..8dc8159939 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl @@ -0,0 +1,58 @@ +void Set_float2(inout float2 vec, int idx, float val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_int3(inout int3 vec, int idx, int val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + } +} +void Set_uint4(inout uint4 vec, int idx, uint val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + case 2: vec[2] = val; break; + case 3: vec[3] = val; break; + } +} +void Set_bool2(inout bool2 vec, int idx, bool val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +[numthreads(1, 1, 1)] +void main() { + float2 v2f = float2(0.0f, 0.0f); + float2 v2f_2 = float2(0.0f, 0.0f); + int3 v3i = int3(0, 0, 0); + int3 v3i_2 = int3(0, 0, 0); + uint4 v4u = uint4(0u, 0u, 0u, 0u); + uint4 v4u_2 = uint4(0u, 0u, 0u, 0u); + bool2 v2b = bool2(false, false); + bool2 v2b_2 = bool2(false, false); + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + Set_float2(v2f, i, 1.0f); + Set_int3(v3i, i, 1); + Set_uint4(v4u, i, 1u); + Set_bool2(v2b, i, true); + Set_float2(v2f_2, i, 1.0f); + Set_int3(v3i_2, i, 1); + Set_uint4(v4u_2, i, 1u); + Set_bool2(v2b_2, i, true); + { + i = (i + 1); + } + } + } + return; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl new file mode 100644 index 0000000000..9f28e5c096 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl @@ -0,0 +1,34 @@ +#include + +using namespace metal; +kernel void tint_symbol() { + float2 v2f = 0.0f; + float2 v2f_2 = 0.0f; + int3 v3i = 0; + int3 v3i_2 = 0; + uint4 v4u = 0u; + uint4 v4u_2 = 0u; + bool2 v2b = false; + bool2 v2b_2 = false; + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + v2f[i] = 1.0f; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; + v2f_2[i] = 1.0f; + v3i_2[i] = 1; + v4u_2[i] = 1u; + v2b_2[i] = true; + { + i = (i + 1); + } + } + } + return; +} + diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm new file mode 100644 index 0000000000..ee5714c43e --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm @@ -0,0 +1,106 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 68 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %main "main" + OpName %v2f "v2f" + OpName %v2f_2 "v2f_2" + OpName %v3i "v3i" + OpName %v3i_2 "v3i_2" + OpName %v4u "v4u" + OpName %v4u_2 "v4u_2" + OpName %v2b "v2b" + OpName %v2b_2 "v2b_2" + OpName %i "i" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %9 = OpConstantNull %v2float + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %15 = OpConstantNull %v3int + %uint = OpTypeInt 32 0 + %v4uint = OpTypeVector %uint 4 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %21 = OpConstantNull %v4uint + %bool = OpTypeBool + %v2bool = OpTypeVector %bool 2 +%_ptr_Function_v2bool = OpTypePointer Function %v2bool + %27 = OpConstantNull %v2bool + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %32 = OpConstantNull %int + %int_2 = OpConstant %int 2 +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %int_1 = OpConstant %int 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Function_bool = OpTypePointer Function %bool + %true = OpConstantTrue %bool + %main = OpFunction %void None %1 + %4 = OpLabel + %v2f = OpVariable %_ptr_Function_v2float Function %9 + %v2f_2 = OpVariable %_ptr_Function_v2float Function %9 + %v3i = OpVariable %_ptr_Function_v3int Function %15 + %v3i_2 = OpVariable %_ptr_Function_v3int Function %15 + %v4u = OpVariable %_ptr_Function_v4uint Function %21 + %v4u_2 = OpVariable %_ptr_Function_v4uint Function %21 + %v2b = OpVariable %_ptr_Function_v2bool Function %27 + %v2b_2 = OpVariable %_ptr_Function_v2bool Function %27 + %i = OpVariable %_ptr_Function_int Function %32 + OpStore %i %int_0 + OpBranch %33 + %33 = OpLabel + OpLoopMerge %34 %35 None + OpBranch %36 + %36 = OpLabel + %38 = OpLoad %int %i + %40 = OpSLessThan %bool %38 %int_2 + %37 = OpLogicalNot %bool %40 + OpSelectionMerge %41 None + OpBranchConditional %37 %42 %41 + %42 = OpLabel + OpBranch %34 + %41 = OpLabel + %43 = OpLoad %int %i + %45 = OpAccessChain %_ptr_Function_float %v2f %43 + OpStore %45 %float_1 + %47 = OpLoad %int %i + %48 = OpAccessChain %_ptr_Function_int %v3i %47 + OpStore %48 %int_1 + %50 = OpLoad %int %i + %52 = OpAccessChain %_ptr_Function_uint %v4u %50 + OpStore %52 %uint_1 + %54 = OpLoad %int %i + %56 = OpAccessChain %_ptr_Function_bool %v2b %54 + OpStore %56 %true + %58 = OpLoad %int %i + %59 = OpAccessChain %_ptr_Function_float %v2f_2 %58 + OpStore %59 %float_1 + %60 = OpLoad %int %i + %61 = OpAccessChain %_ptr_Function_int %v3i_2 %60 + OpStore %61 %int_1 + %62 = OpLoad %int %i + %63 = OpAccessChain %_ptr_Function_uint %v4u_2 %62 + OpStore %63 %uint_1 + %64 = OpLoad %int %i + %65 = OpAccessChain %_ptr_Function_bool %v2b_2 %64 + OpStore %65 %true + OpBranch %35 + %35 = OpLabel + %66 = OpLoad %int %i + %67 = OpIAdd %int %66 %int_1 + OpStore %i %67 + OpBranch %33 + %34 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl new file mode 100644 index 0000000000..02293235ee --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl @@ -0,0 +1,31 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v2f_2 : vec2; + var v3i : vec3; + var v3i_2 : vec3; + var v4u : vec4; + var v4u_2 : vec4; + var v2b : vec2; + var v2b_2 : vec2; + { + var i : i32 = 0; + loop { + if (!((i < 2))) { + break; + } + v2f[i] = 1.0; + v3i[i] = 1; + v4u[i] = 1u; + v2b[i] = true; + v2f_2[i] = 1.0; + v3i_2[i] = 1; + v4u_2[i] = 1u; + v2b_2[i] = true; + + continuing { + i = (i + 1); + } + } + } +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl new file mode 100644 index 0000000000..d07e08d46b --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl @@ -0,0 +1,32 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + + for (var i : i32 = 0; i < 2; i = i + 1) { + v2f[i] = 1.0; + v2i[i] = 1; + v2u[i] = 1u; + v2b[i] = true; + } + + var i : i32 = 0; + v3f[i] = 1.0; + v4f[i] = 1.0; + v3i[i] = 1; + v4i[i] = 1; + v3u[i] = 1u; + v4u[i] = 1u; + v3b[i] = true; + v4b[i] = true; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl new file mode 100644 index 0000000000..953595e88c --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl @@ -0,0 +1,64 @@ +void Set_float2(inout float2 vec, int idx, float val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_int2(inout int2 vec, int idx, int val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_uint2(inout uint2 vec, int idx, uint val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +void Set_bool2(inout bool2 vec, int idx, bool val) { + switch(idx) { + case 0: vec[0] = val; break; + case 1: vec[1] = val; break; + } +} +[numthreads(1, 1, 1)] +void main() { + float2 v2f = float2(0.0f, 0.0f); + float3 v3f = float3(0.0f, 0.0f, 0.0f); + float4 v4f = float4(0.0f, 0.0f, 0.0f, 0.0f); + int2 v2i = int2(0, 0); + int3 v3i = int3(0, 0, 0); + int4 v4i = int4(0, 0, 0, 0); + uint2 v2u = uint2(0u, 0u); + uint3 v3u = uint3(0u, 0u, 0u); + uint4 v4u = uint4(0u, 0u, 0u, 0u); + bool2 v2b = bool2(false, false); + bool3 v3b = bool3(false, false, false); + bool4 v4b = bool4(false, false, false, false); + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + Set_float2(v2f, i, 1.0f); + Set_int2(v2i, i, 1); + Set_uint2(v2u, i, 1u); + Set_bool2(v2b, i, true); + { + i = (i + 1); + } + } + } + int i = 0; + v3f[i] = 1.0f; + v4f[i] = 1.0f; + v3i[i] = 1; + v4i[i] = 1; + v3u[i] = 1u; + v4u[i] = 1u; + v3b[i] = true; + v4b[i] = true; + return; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl new file mode 100644 index 0000000000..fb5319916d --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl @@ -0,0 +1,43 @@ +#include + +using namespace metal; +kernel void tint_symbol() { + float2 v2f = 0.0f; + float3 v3f = 0.0f; + float4 v4f = 0.0f; + int2 v2i = 0; + int3 v3i = 0; + int4 v4i = 0; + uint2 v2u = 0u; + uint3 v3u = 0u; + uint4 v4u = 0u; + bool2 v2b = false; + bool3 v3b = false; + bool4 v4b = false; + { + int i = 0; + while (true) { + if (!((i < 2))) { + break; + } + v2f[i] = 1.0f; + v2i[i] = 1; + v2u[i] = 1u; + v2b[i] = true; + { + i = (i + 1); + } + } + } + int i = 0; + v3f[i] = 1.0f; + v4f[i] = 1.0f; + v3i[i] = 1; + v4i[i] = 1; + v3u[i] = 1u; + v4u[i] = 1u; + v3b[i] = true; + v4b[i] = true; + return; +} + diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm new file mode 100644 index 0000000000..d2938bbe2c --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm @@ -0,0 +1,153 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 105 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %main "main" + OpName %v2f "v2f" + OpName %v3f "v3f" + OpName %v4f "v4f" + OpName %v2i "v2i" + OpName %v3i "v3i" + OpName %v4i "v4i" + OpName %v2u "v2u" + OpName %v3u "v3u" + OpName %v4u "v4u" + OpName %v2b "v2b" + OpName %v3b "v3b" + OpName %v4b "v4b" + OpName %i "i" + OpName %i_0 "i" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %9 = OpConstantNull %v2float + %v3float = OpTypeVector %float 3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %13 = OpConstantNull %v3float + %v4float = OpTypeVector %float 4 +%_ptr_Function_v4float = OpTypePointer Function %v4float + %17 = OpConstantNull %v4float + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 +%_ptr_Function_v2int = OpTypePointer Function %v2int + %22 = OpConstantNull %v2int + %v3int = OpTypeVector %int 3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %26 = OpConstantNull %v3int + %v4int = OpTypeVector %int 4 +%_ptr_Function_v4int = OpTypePointer Function %v4int + %30 = OpConstantNull %v4int + %uint = OpTypeInt 32 0 + %v2uint = OpTypeVector %uint 2 +%_ptr_Function_v2uint = OpTypePointer Function %v2uint + %35 = OpConstantNull %v2uint + %v3uint = OpTypeVector %uint 3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %39 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %43 = OpConstantNull %v4uint + %bool = OpTypeBool + %v2bool = OpTypeVector %bool 2 +%_ptr_Function_v2bool = OpTypePointer Function %v2bool + %48 = OpConstantNull %v2bool + %v3bool = OpTypeVector %bool 3 +%_ptr_Function_v3bool = OpTypePointer Function %v3bool + %52 = OpConstantNull %v3bool + %v4bool = OpTypeVector %bool 4 +%_ptr_Function_v4bool = OpTypePointer Function %v4bool + %56 = OpConstantNull %v4bool + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %60 = OpConstantNull %int + %int_2 = OpConstant %int 2 +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %int_1 = OpConstant %int 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Function_bool = OpTypePointer Function %bool + %true = OpConstantTrue %bool + %main = OpFunction %void None %1 + %4 = OpLabel + %v2f = OpVariable %_ptr_Function_v2float Function %9 + %v3f = OpVariable %_ptr_Function_v3float Function %13 + %v4f = OpVariable %_ptr_Function_v4float Function %17 + %v2i = OpVariable %_ptr_Function_v2int Function %22 + %v3i = OpVariable %_ptr_Function_v3int Function %26 + %v4i = OpVariable %_ptr_Function_v4int Function %30 + %v2u = OpVariable %_ptr_Function_v2uint Function %35 + %v3u = OpVariable %_ptr_Function_v3uint Function %39 + %v4u = OpVariable %_ptr_Function_v4uint Function %43 + %v2b = OpVariable %_ptr_Function_v2bool Function %48 + %v3b = OpVariable %_ptr_Function_v3bool Function %52 + %v4b = OpVariable %_ptr_Function_v4bool Function %56 + %i = OpVariable %_ptr_Function_int Function %60 + %i_0 = OpVariable %_ptr_Function_int Function %60 + OpStore %i %int_0 + OpBranch %61 + %61 = OpLabel + OpLoopMerge %62 %63 None + OpBranch %64 + %64 = OpLabel + %66 = OpLoad %int %i + %68 = OpSLessThan %bool %66 %int_2 + %65 = OpLogicalNot %bool %68 + OpSelectionMerge %69 None + OpBranchConditional %65 %70 %69 + %70 = OpLabel + OpBranch %62 + %69 = OpLabel + %71 = OpLoad %int %i + %73 = OpAccessChain %_ptr_Function_float %v2f %71 + OpStore %73 %float_1 + %75 = OpLoad %int %i + %76 = OpAccessChain %_ptr_Function_int %v2i %75 + OpStore %76 %int_1 + %78 = OpLoad %int %i + %80 = OpAccessChain %_ptr_Function_uint %v2u %78 + OpStore %80 %uint_1 + %82 = OpLoad %int %i + %84 = OpAccessChain %_ptr_Function_bool %v2b %82 + OpStore %84 %true + OpBranch %63 + %63 = OpLabel + %86 = OpLoad %int %i + %87 = OpIAdd %int %86 %int_1 + OpStore %i %87 + OpBranch %61 + %62 = OpLabel + OpStore %i_0 %int_0 + %89 = OpLoad %int %i_0 + %90 = OpAccessChain %_ptr_Function_float %v3f %89 + OpStore %90 %float_1 + %91 = OpLoad %int %i_0 + %92 = OpAccessChain %_ptr_Function_float %v4f %91 + OpStore %92 %float_1 + %93 = OpLoad %int %i_0 + %94 = OpAccessChain %_ptr_Function_int %v3i %93 + OpStore %94 %int_1 + %95 = OpLoad %int %i_0 + %96 = OpAccessChain %_ptr_Function_int %v4i %95 + OpStore %96 %int_1 + %97 = OpLoad %int %i_0 + %98 = OpAccessChain %_ptr_Function_uint %v3u %97 + OpStore %98 %uint_1 + %99 = OpLoad %int %i_0 + %100 = OpAccessChain %_ptr_Function_uint %v4u %99 + OpStore %100 %uint_1 + %101 = OpLoad %int %i_0 + %102 = OpAccessChain %_ptr_Function_bool %v3b %101 + OpStore %102 %true + %103 = OpLoad %int %i_0 + %104 = OpAccessChain %_ptr_Function_bool %v4b %103 + OpStore %104 %true + OpReturn + OpFunctionEnd diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl new file mode 100644 index 0000000000..8c541461fb --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl @@ -0,0 +1,40 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + { + var i : i32 = 0; + loop { + if (!((i < 2))) { + break; + } + v2f[i] = 1.0; + v2i[i] = 1; + v2u[i] = 1u; + v2b[i] = true; + + continuing { + i = (i + 1); + } + } + } + var i : i32 = 0; + v3f[i] = 1.0; + v4f[i] = 1.0; + v3i[i] = 1; + v4i[i] = 1; + v3u[i] = 1u; + v4u[i] = 1u; + v3b[i] = true; + v4b[i] = true; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl new file mode 100644 index 0000000000..8103edfa4f --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl @@ -0,0 +1,29 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + + var i : i32 = 0; + v2f[i] = 1.0; + v3f[i] = 1.0; + v4f[i] = 1.0; + v2i[i] = 1; + v3i[i] = 1; + v4i[i] = 1; + v2u[i] = 1u; + v3u[i] = 1u; + v4u[i] = 1u; + v2b[i] = true; + v3b[i] = true; + v4b[i] = true; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl new file mode 100644 index 0000000000..f836c66e07 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl @@ -0,0 +1,29 @@ +[numthreads(1, 1, 1)] +void main() { + float2 v2f = float2(0.0f, 0.0f); + float3 v3f = float3(0.0f, 0.0f, 0.0f); + float4 v4f = float4(0.0f, 0.0f, 0.0f, 0.0f); + int2 v2i = int2(0, 0); + int3 v3i = int3(0, 0, 0); + int4 v4i = int4(0, 0, 0, 0); + uint2 v2u = uint2(0u, 0u); + uint3 v3u = uint3(0u, 0u, 0u); + uint4 v4u = uint4(0u, 0u, 0u, 0u); + bool2 v2b = bool2(false, false); + bool3 v3b = bool3(false, false, false); + bool4 v4b = bool4(false, false, false, false); + int i = 0; + v2f[i] = 1.0f; + v3f[i] = 1.0f; + v4f[i] = 1.0f; + v2i[i] = 1; + v3i[i] = 1; + v4i[i] = 1; + v2u[i] = 1u; + v3u[i] = 1u; + v4u[i] = 1u; + v2b[i] = true; + v3b[i] = true; + v4b[i] = true; + return; +} diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.msl new file mode 100644 index 0000000000..cb0748eddd --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.msl @@ -0,0 +1,32 @@ +#include + +using namespace metal; +kernel void tint_symbol() { + float2 v2f = 0.0f; + float3 v3f = 0.0f; + float4 v4f = 0.0f; + int2 v2i = 0; + int3 v3i = 0; + int4 v4i = 0; + uint2 v2u = 0u; + uint3 v3u = 0u; + uint4 v4u = 0u; + bool2 v2b = false; + bool3 v3b = false; + bool4 v4b = false; + int i = 0; + v2f[i] = 1.0f; + v3f[i] = 1.0f; + v4f[i] = 1.0f; + v2i[i] = 1; + v3i[i] = 1; + v4i[i] = 1; + v2u[i] = 1u; + v3u[i] = 1u; + v4u[i] = 1u; + v2b[i] = true; + v3b[i] = true; + v4b[i] = true; + return; +} + diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm new file mode 100644 index 0000000000..fc900eeeab --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm @@ -0,0 +1,129 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 92 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %main "main" + OpName %v2f "v2f" + OpName %v3f "v3f" + OpName %v4f "v4f" + OpName %v2i "v2i" + OpName %v3i "v3i" + OpName %v4i "v4i" + OpName %v2u "v2u" + OpName %v3u "v3u" + OpName %v4u "v4u" + OpName %v2b "v2b" + OpName %v3b "v3b" + OpName %v4b "v4b" + OpName %i "i" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v2float = OpTypeVector %float 2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %9 = OpConstantNull %v2float + %v3float = OpTypeVector %float 3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %13 = OpConstantNull %v3float + %v4float = OpTypeVector %float 4 +%_ptr_Function_v4float = OpTypePointer Function %v4float + %17 = OpConstantNull %v4float + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 +%_ptr_Function_v2int = OpTypePointer Function %v2int + %22 = OpConstantNull %v2int + %v3int = OpTypeVector %int 3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %26 = OpConstantNull %v3int + %v4int = OpTypeVector %int 4 +%_ptr_Function_v4int = OpTypePointer Function %v4int + %30 = OpConstantNull %v4int + %uint = OpTypeInt 32 0 + %v2uint = OpTypeVector %uint 2 +%_ptr_Function_v2uint = OpTypePointer Function %v2uint + %35 = OpConstantNull %v2uint + %v3uint = OpTypeVector %uint 3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %39 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %43 = OpConstantNull %v4uint + %bool = OpTypeBool + %v2bool = OpTypeVector %bool 2 +%_ptr_Function_v2bool = OpTypePointer Function %v2bool + %48 = OpConstantNull %v2bool + %v3bool = OpTypeVector %bool 3 +%_ptr_Function_v3bool = OpTypePointer Function %v3bool + %52 = OpConstantNull %v3bool + %v4bool = OpTypeVector %bool 4 +%_ptr_Function_v4bool = OpTypePointer Function %v4bool + %56 = OpConstantNull %v4bool + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %60 = OpConstantNull %int +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %int_1 = OpConstant %int 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Function_bool = OpTypePointer Function %bool + %true = OpConstantTrue %bool + %main = OpFunction %void None %1 + %4 = OpLabel + %v2f = OpVariable %_ptr_Function_v2float Function %9 + %v3f = OpVariable %_ptr_Function_v3float Function %13 + %v4f = OpVariable %_ptr_Function_v4float Function %17 + %v2i = OpVariable %_ptr_Function_v2int Function %22 + %v3i = OpVariable %_ptr_Function_v3int Function %26 + %v4i = OpVariable %_ptr_Function_v4int Function %30 + %v2u = OpVariable %_ptr_Function_v2uint Function %35 + %v3u = OpVariable %_ptr_Function_v3uint Function %39 + %v4u = OpVariable %_ptr_Function_v4uint Function %43 + %v2b = OpVariable %_ptr_Function_v2bool Function %48 + %v3b = OpVariable %_ptr_Function_v3bool Function %52 + %v4b = OpVariable %_ptr_Function_v4bool Function %56 + %i = OpVariable %_ptr_Function_int Function %60 + OpStore %i %int_0 + %61 = OpLoad %int %i + %63 = OpAccessChain %_ptr_Function_float %v2f %61 + OpStore %63 %float_1 + %65 = OpLoad %int %i + %66 = OpAccessChain %_ptr_Function_float %v3f %65 + OpStore %66 %float_1 + %67 = OpLoad %int %i + %68 = OpAccessChain %_ptr_Function_float %v4f %67 + OpStore %68 %float_1 + %69 = OpLoad %int %i + %70 = OpAccessChain %_ptr_Function_int %v2i %69 + OpStore %70 %int_1 + %72 = OpLoad %int %i + %73 = OpAccessChain %_ptr_Function_int %v3i %72 + OpStore %73 %int_1 + %74 = OpLoad %int %i + %75 = OpAccessChain %_ptr_Function_int %v4i %74 + OpStore %75 %int_1 + %76 = OpLoad %int %i + %78 = OpAccessChain %_ptr_Function_uint %v2u %76 + OpStore %78 %uint_1 + %80 = OpLoad %int %i + %81 = OpAccessChain %_ptr_Function_uint %v3u %80 + OpStore %81 %uint_1 + %82 = OpLoad %int %i + %83 = OpAccessChain %_ptr_Function_uint %v4u %82 + OpStore %83 %uint_1 + %84 = OpLoad %int %i + %86 = OpAccessChain %_ptr_Function_bool %v2b %84 + OpStore %86 %true + %88 = OpLoad %int %i + %89 = OpAccessChain %_ptr_Function_bool %v3b %88 + OpStore %89 %true + %90 = OpLoad %int %i + %91 = OpAccessChain %_ptr_Function_bool %v4b %90 + OpStore %91 %true + OpReturn + OpFunctionEnd diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl new file mode 100644 index 0000000000..2d267744b5 --- /dev/null +++ b/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl @@ -0,0 +1,28 @@ +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main() { + var v2f : vec2; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + var i : i32 = 0; + v2f[i] = 1.0; + v3f[i] = 1.0; + v4f[i] = 1.0; + v2i[i] = 1; + v3i[i] = 1; + v4i[i] = 1; + v2u[i] = 1u; + v3u[i] = 1u; + v4u[i] = 1u; + v2b[i] = true; + v3b[i] = true; + v4b[i] = true; +}