From 0f916164aeac6fef7fbad745ad6b2e96488bb9a4 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Sat, 19 Jun 2021 08:18:50 +0000 Subject: [PATCH] writer/hlsl: Add missing parenthesis around UBO ternary op Bug: tint:913 Change-Id: I2edbab363cb03e6ce64b5c6bddf184bf92438521 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55340 Auto-Submit: Ben Clayton Reviewed-by: Sarah Mashayekhi Kokoro: Kokoro Commit-Queue: Ben Clayton --- src/writer/hlsl/generator_impl.cc | 4 +- test/bug/tint/749.spvasm.expected.hlsl | 2 +- test/bug/tint/913.wgsl | 70 ++++++ test/bug/tint/913.wgsl.expected.hlsl | 102 +++++++++ test/bug/tint/913.wgsl.expected.msl | 48 ++++ test/bug/tint/913.wgsl.expected.spvasm | 294 +++++++++++++++++++++++++ test/bug/tint/913.wgsl.expected.wgsl | 55 +++++ 7 files changed, 572 insertions(+), 3 deletions(-) create mode 100644 test/bug/tint/913.wgsl create mode 100644 test/bug/tint/913.wgsl.expected.hlsl create mode 100644 test/bug/tint/913.wgsl.expected.msl create mode 100644 test/bug/tint/913.wgsl.expected.spvasm create mode 100644 test/bug/tint/913.wgsl.expected.wgsl diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 810985118c..b9cd81c2e4 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -564,8 +564,8 @@ bool GeneratorImpl::EmitUniformBufferAccess( make_indent(ss << ";" << std::endl); pre << ss.str(); - out << "(" << scalar_offset << " & 2) ? " << ubo_load - << ".zw : " << ubo_load << ".xy"; + out << "((" << scalar_offset << " & 2) ? " << ubo_load + << ".zw : " << ubo_load << ".xy)"; return true; }; // vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle diff --git a/test/bug/tint/749.spvasm.expected.hlsl b/test/bug/tint/749.spvasm.expected.hlsl index a8f610dec5..8d3b29803e 100644 --- a/test/bug/tint/749.spvasm.expected.hlsl +++ b/test/bug/tint/749.spvasm.expected.hlsl @@ -868,7 +868,7 @@ void main_1() { uv = x_762; const int scalar_offset = (0u) / 4; uint4 ubo_load = x_188[scalar_offset / 4]; - const float2 x_191 = asfloat((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy); + const float2 x_191 = asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)); const QuicksortObject x_763 = obj; const tint_array_wrapper tint_symbol_51 = {{0, 0, 0, 0, 0, 0, 0, 0, 0, 0}}; const QuicksortObject tint_symbol_52 = {tint_symbol_51}; diff --git a/test/bug/tint/913.wgsl b/test/bug/tint/913.wgsl new file mode 100644 index 0000000000..467a9443a0 --- /dev/null +++ b/test/bug/tint/913.wgsl @@ -0,0 +1,70 @@ +[[block]] struct Uniforms { + dstTextureFlipY : u32; + channelCount : u32; + srcCopyOrigin : vec2; + dstCopyOrigin : vec2; + copySize : vec2; +}; +[[block]] struct OutputBuf { + result : 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; +fn aboutEqual(value : f32, expect : f32) -> bool { + // The value diff should be smaller than the hard coded tolerance. + return abs(value - expect) < 0.001; +} +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { + let srcSize : vec2 = textureDimensions(src); + let dstSize : vec2 = textureDimensions(dst); + let dstTexCoord : vec2 = vec2(GlobalInvocationID.xy); + let nonCoveredColor : vec4 = + vec4(0.0, 1.0, 0.0, 1.0); // should be green + + var success : bool = true; + if (dstTexCoord.x < uniforms.dstCopyOrigin.x || + dstTexCoord.y < uniforms.dstCopyOrigin.y || + dstTexCoord.x >= uniforms.dstCopyOrigin.x + uniforms.copySize.x || + dstTexCoord.y >= uniforms.dstCopyOrigin.y + uniforms.copySize.y) { + success = success && + all(textureLoad(dst, vec2(dstTexCoord), 0) == nonCoveredColor); + } else { + // Calculate source texture coord. + var srcTexCoord : vec2 = dstTexCoord - uniforms.dstCopyOrigin + + uniforms.srcCopyOrigin; + // Note that |flipY| equals flip src texture firstly and then do copy from src + // subrect to dst subrect. This helps on blink part to handle some input texture + // which is flipped and need to unpack flip during the copy. + // We need to calculate the expect y coord based on this rule. + if (uniforms.dstTextureFlipY == 1u) { + srcTexCoord.y = u32(srcSize.y) - srcTexCoord.y - 1u; + } + + let srcColor : vec4 = textureLoad(src, vec2(srcTexCoord), 0); + let dstColor : vec4 = textureLoad(dst, vec2(dstTexCoord), 0); + + // Not use loop and variable index format to workaround + // crbug.com/tint/638. + if (uniforms.channelCount == 2u) { // All have rg components. + success = success && + aboutEqual(dstColor.r, srcColor.r) && + aboutEqual(dstColor.g, srcColor.g); + } else { + success = success && + aboutEqual(dstColor.r, srcColor.r) && + aboutEqual(dstColor.g, srcColor.g) && + aboutEqual(dstColor.b, srcColor.b) && + aboutEqual(dstColor.a, srcColor.a); + } + } + let outputIndex : u32 = GlobalInvocationID.y * u32(dstSize.x) + + GlobalInvocationID.x; + if (success) { + output.result[outputIndex] = 1u; + } else { + output.result[outputIndex] = 0u; + } +} \ No newline at end of file diff --git a/test/bug/tint/913.wgsl.expected.hlsl b/test/bug/tint/913.wgsl.expected.hlsl new file mode 100644 index 0000000000..ada977db3f --- /dev/null +++ b/test/bug/tint/913.wgsl.expected.hlsl @@ -0,0 +1,102 @@ +Texture2D src : register(t0, space0); +Texture2D tint_symbol : register(t1, space0); +RWByteAddressBuffer output : register(u2, space0); +cbuffer cbuffer_uniforms : register(b3, space0) { + uint4 uniforms[2]; +}; + +bool aboutEqual(float value, float expect) { + return (abs((value - expect)) < 0.001f); +} + +struct tint_symbol_2 { + uint3 GlobalInvocationID : SV_DispatchThreadID; +}; + +[numthreads(1, 1, 1)] +void main(tint_symbol_2 tint_symbol_1) { + const uint3 GlobalInvocationID = tint_symbol_1.GlobalInvocationID; + int2 tint_tmp; + src.GetDimensions(tint_tmp.x, tint_tmp.y); + const int2 srcSize = tint_tmp; + int2 tint_tmp_1; + tint_symbol.GetDimensions(tint_tmp_1.x, tint_tmp_1.y); + const int2 dstSize = tint_tmp_1; + const uint2 dstTexCoord = uint2(GlobalInvocationID.xy); + const float4 nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f); + bool success = true; + const int scalar_offset = (16u) / 4; + bool tint_tmp_2 = (dstTexCoord.x < uniforms[scalar_offset / 4][scalar_offset % 4]); + if (!tint_tmp_2) { +const int scalar_offset_1 = (20u) / 4; + tint_tmp_2 = (dstTexCoord.y < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]); + } + bool tint_tmp_3 = (tint_tmp_2); + if (!tint_tmp_3) { +const int scalar_offset_2 = (16u) / 4; + const int scalar_offset_3 = (24u) / 4; + tint_tmp_3 = (dstTexCoord.x >= (uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4] + uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4])); + } + bool tint_tmp_4 = (tint_tmp_3); + if (!tint_tmp_4) { +const int scalar_offset_4 = (20u) / 4; + const int scalar_offset_5 = (28u) / 4; + tint_tmp_4 = (dstTexCoord.y >= (uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4] + uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4])); + } +if ((tint_tmp_4)) { + bool tint_tmp_5 = success; + if (tint_tmp_5) { + tint_tmp_5 = all((tint_symbol.Load(int3(dstTexCoord, 0), 0) == nonCoveredColor)); + } +success = (tint_tmp_5); + } else { + const int scalar_offset_6 = (16u) / 4; + uint4 ubo_load = uniforms[scalar_offset_6 / 4]; + const int scalar_offset_7 = (8u) / 4; + uint4 ubo_load_1 = uniforms[scalar_offset_7 / 4]; + uint2 srcTexCoord = ((dstTexCoord - ((scalar_offset_6 & 2) ? ubo_load.zw : ubo_load.xy)) + ((scalar_offset_7 & 2) ? ubo_load_1.zw : ubo_load_1.xy)); + const int scalar_offset_8 = (0u) / 4; + if ((uniforms[scalar_offset_8 / 4][scalar_offset_8 % 4] == 1u)) { + srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u); + } + const float4 srcColor = src.Load(int3(srcTexCoord, 0), 0); + const float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0), 0); + const int scalar_offset_9 = (4u) / 4; + if ((uniforms[scalar_offset_9 / 4][scalar_offset_9 % 4] == 2u)) { + bool tint_tmp_6 = success; + if (tint_tmp_6) { + tint_tmp_6 = aboutEqual(dstColor.r, srcColor.r); + } + bool tint_tmp_7 = (tint_tmp_6); + if (tint_tmp_7) { + tint_tmp_7 = aboutEqual(dstColor.g, srcColor.g); + } +success = (tint_tmp_7); + } else { + bool tint_tmp_8 = success; + if (tint_tmp_8) { + tint_tmp_8 = aboutEqual(dstColor.r, srcColor.r); + } + bool tint_tmp_9 = (tint_tmp_8); + if (tint_tmp_9) { + tint_tmp_9 = aboutEqual(dstColor.g, srcColor.g); + } + bool tint_tmp_10 = (tint_tmp_9); + if (tint_tmp_10) { + tint_tmp_10 = aboutEqual(dstColor.b, srcColor.b); + } + bool tint_tmp_11 = (tint_tmp_10); + if (tint_tmp_11) { + tint_tmp_11 = aboutEqual(dstColor.a, srcColor.a); + } +success = (tint_tmp_11); + } + } + const uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x); + if (success) { + output.Store((4u * outputIndex), asuint(1u)); + } else { + output.Store((4u * outputIndex), asuint(0u)); + } + return; +} diff --git a/test/bug/tint/913.wgsl.expected.msl b/test/bug/tint/913.wgsl.expected.msl new file mode 100644 index 0000000000..e30f694978 --- /dev/null +++ b/test/bug/tint/913.wgsl.expected.msl @@ -0,0 +1,48 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint dstTextureFlipY; + /* 0x0004 */ uint channelCount; + /* 0x0008 */ packed_uint2 srcCopyOrigin; + /* 0x0010 */ packed_uint2 dstCopyOrigin; + /* 0x0018 */ packed_uint2 copySize; +}; +struct OutputBuf { + /* 0x0000 */ uint result[1]; +}; + +bool aboutEqual(float value, float expect) { + return ( fabs((value - expect)) < 0.001f); +} + +kernel void tint_symbol(texture2d tint_symbol_2 [[texture(0)]], texture2d tint_symbol_3 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) { + int2 const srcSize = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height()); + int2 const dstSize = int2(tint_symbol_3.get_width(), tint_symbol_3.get_height()); + uint2 const dstTexCoord = uint2(GlobalInvocationID.xy); + float4 const nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f); + bool success = true; + if (((((dstTexCoord.x < uniforms.dstCopyOrigin.x) || (dstTexCoord.y < uniforms.dstCopyOrigin.y)) || (dstTexCoord.x >= (uniforms.dstCopyOrigin.x + uniforms.copySize.x))) || (dstTexCoord.y >= (uniforms.dstCopyOrigin.y + uniforms.copySize.y)))) { + success = (success && all((tint_symbol_3.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor))); + } else { + uint2 srcTexCoord = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin); + if ((uniforms.dstTextureFlipY == 1u)) { + srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u); + } + float4 const srcColor = tint_symbol_2.read(uint2(int2(srcTexCoord)), 0); + float4 const dstColor = tint_symbol_3.read(uint2(int2(dstTexCoord)), 0); + if ((uniforms.channelCount == 2u)) { + success = ((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g)); + } else { + success = ((((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g)) && aboutEqual(dstColor.b, srcColor.b)) && aboutEqual(dstColor.a, srcColor.a)); + } + } + uint const outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x); + if (success) { + output.result[outputIndex] = 1u; + } else { + output.result[outputIndex] = 0u; + } + return; +} + diff --git a/test/bug/tint/913.wgsl.expected.spvasm b/test/bug/tint/913.wgsl.expected.spvasm new file mode 100644 index 0000000000..0a66bbb823 --- /dev/null +++ b/test/bug/tint/913.wgsl.expected.spvasm @@ -0,0 +1,294 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 191 +; Schema: 0 + OpCapability Shader + OpCapability ImageQuery + %25 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %tint_symbol + OpExecutionMode %main LocalSize 1 1 1 + OpName %src "src" + OpName %dst "dst" + OpName %OutputBuf "OutputBuf" + OpMemberName %OutputBuf 0 "result" + OpName %output "output" + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "dstTextureFlipY" + OpMemberName %Uniforms 1 "channelCount" + OpMemberName %Uniforms 2 "srcCopyOrigin" + OpMemberName %Uniforms 3 "dstCopyOrigin" + OpMemberName %Uniforms 4 "copySize" + OpName %uniforms "uniforms" + OpName %tint_symbol "tint_symbol" + OpName %aboutEqual "aboutEqual" + OpName %value "value" + OpName %expect "expect" + OpName %main "main" + OpName %success "success" + OpName %srcTexCoord "srcTexCoord" + OpDecorate %src DescriptorSet 0 + OpDecorate %src Binding 0 + OpDecorate %dst DescriptorSet 0 + OpDecorate %dst Binding 1 + OpDecorate %OutputBuf Block + OpMemberDecorate %OutputBuf 0 Offset 0 + OpDecorate %_runtimearr_uint ArrayStride 4 + OpDecorate %output DescriptorSet 0 + OpDecorate %output Binding 2 + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpMemberDecorate %Uniforms 2 Offset 8 + OpMemberDecorate %Uniforms 3 Offset 16 + OpMemberDecorate %Uniforms 4 Offset 24 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 0 + OpDecorate %uniforms Binding 3 + OpDecorate %tint_symbol BuiltIn GlobalInvocationId + %float = OpTypeFloat 32 + %3 = OpTypeImage %float 2D 0 0 0 1 Unknown +%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3 + %src = OpVariable %_ptr_UniformConstant_3 UniformConstant + %dst = OpVariable %_ptr_UniformConstant_3 UniformConstant + %uint = OpTypeInt 32 0 +%_runtimearr_uint = OpTypeRuntimeArray %uint + %OutputBuf = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer_OutputBuf = OpTypePointer StorageBuffer %OutputBuf + %output = OpVariable %_ptr_StorageBuffer_OutputBuf StorageBuffer + %v2uint = OpTypeVector %uint 2 + %Uniforms = OpTypeStruct %uint %uint %v2uint %v2uint %v2uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%tint_symbol = OpVariable %_ptr_Input_v3uint Input + %bool = OpTypeBool + %18 = OpTypeFunction %bool %float %float +%float_0_00100000005 = OpConstant %float 0.00100000005 + %void = OpTypeVoid + %29 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 + %int_0 = OpConstant %int 0 + %v4float = OpTypeVector %float 4 + %float_0 = OpConstant %float 0 + %float_1 = OpConstant %float 1 + %46 = OpConstantComposite %v4float %float_0 %float_1 %float_0 %float_1 + %true = OpConstantTrue %bool +%_ptr_Function_bool = OpTypePointer Function %bool + %50 = OpConstantNull %bool + %uint_3 = OpConstant %uint 3 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_1 = OpConstant %uint 1 + %uint_4 = OpConstant %uint 4 + %v4bool = OpTypeVector %bool 4 +%_ptr_Uniform_v2uint = OpTypePointer Uniform %v2uint + %uint_2 = OpConstant %uint 2 +%_ptr_Function_v2uint = OpTypePointer Function %v2uint + %110 = OpConstantNull %v2uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Input_uint = OpTypePointer Input %uint +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %aboutEqual = OpFunction %bool None %18 + %value = OpFunctionParameter %float + %expect = OpFunctionParameter %float + %23 = OpLabel + %26 = OpFSub %float %value %expect + %24 = OpExtInst %float %25 FAbs %26 + %28 = OpFOrdLessThan %bool %24 %float_0_00100000005 + OpReturnValue %28 + OpFunctionEnd + %main = OpFunction %void None %29 + %32 = OpLabel + %success = OpVariable %_ptr_Function_bool Function %50 +%srcTexCoord = OpVariable %_ptr_Function_v2uint Function %110 + %36 = OpLoad %3 %src + %33 = OpImageQuerySizeLod %v2int %36 %int_0 + %39 = OpLoad %3 %dst + %38 = OpImageQuerySizeLod %v2int %39 %int_0 + %41 = OpLoad %v3uint %tint_symbol + %42 = OpVectorShuffle %v2uint %41 %41 0 1 + OpStore %success %true + %51 = OpCompositeExtract %uint %42 0 + %55 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0 + %56 = OpLoad %uint %55 + %57 = OpULessThan %bool %51 %56 + OpSelectionMerge %58 None + OpBranchConditional %57 %58 %59 + %59 = OpLabel + %60 = OpCompositeExtract %uint %42 1 + %62 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1 + %63 = OpLoad %uint %62 + %64 = OpULessThan %bool %60 %63 + OpBranch %58 + %58 = OpLabel + %65 = OpPhi %bool %57 %32 %64 %59 + OpSelectionMerge %66 None + OpBranchConditional %65 %66 %67 + %67 = OpLabel + %68 = OpCompositeExtract %uint %42 0 + %69 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0 + %70 = OpLoad %uint %69 + %72 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_0 + %73 = OpLoad %uint %72 + %74 = OpIAdd %uint %70 %73 + %75 = OpUGreaterThanEqual %bool %68 %74 + OpBranch %66 + %66 = OpLabel + %76 = OpPhi %bool %65 %58 %75 %67 + OpSelectionMerge %77 None + OpBranchConditional %76 %77 %78 + %78 = OpLabel + %79 = OpCompositeExtract %uint %42 1 + %80 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1 + %81 = OpLoad %uint %80 + %82 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_1 + %83 = OpLoad %uint %82 + %84 = OpIAdd %uint %81 %83 + %85 = OpUGreaterThanEqual %bool %79 %84 + OpBranch %77 + %77 = OpLabel + %86 = OpPhi %bool %76 %66 %85 %78 + OpSelectionMerge %87 None + OpBranchConditional %86 %88 %89 + %88 = OpLabel + %90 = OpLoad %bool %success + OpSelectionMerge %91 None + OpBranchConditional %90 %92 %91 + %92 = OpLabel + %95 = OpLoad %3 %dst + %96 = OpBitcast %v2int %42 + %94 = OpImageFetch %v4float %95 %96 Lod %int_0 + %97 = OpFOrdEqual %v4bool %94 %46 + %93 = OpAll %bool %97 + OpBranch %91 + %91 = OpLabel + %99 = OpPhi %bool %90 %88 %93 %92 + OpStore %success %99 + OpBranch %87 + %89 = OpLabel + %101 = OpAccessChain %_ptr_Uniform_v2uint %uniforms %uint_3 + %102 = OpLoad %v2uint %101 + %103 = OpISub %v2uint %42 %102 + %105 = OpAccessChain %_ptr_Uniform_v2uint %uniforms %uint_2 + %106 = OpLoad %v2uint %105 + %107 = OpIAdd %v2uint %103 %106 + OpStore %srcTexCoord %107 + %111 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %112 = OpLoad %uint %111 + %113 = OpIEqual %bool %112 %uint_1 + OpSelectionMerge %114 None + OpBranchConditional %113 %115 %114 + %115 = OpLabel + %117 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1 + %119 = OpCompositeExtract %int %33 1 + %118 = OpBitcast %uint %119 + %120 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1 + %121 = OpLoad %uint %120 + %122 = OpISub %uint %118 %121 + %123 = OpISub %uint %122 %uint_1 + OpStore %117 %123 + OpBranch %114 + %114 = OpLabel + %125 = OpLoad %3 %src + %127 = OpLoad %v2uint %srcTexCoord + %126 = OpBitcast %v2int %127 + %124 = OpImageFetch %v4float %125 %126 Lod %int_0 + %129 = OpLoad %3 %dst + %130 = OpBitcast %v2int %42 + %128 = OpImageFetch %v4float %129 %130 Lod %int_0 + %131 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %132 = OpLoad %uint %131 + %133 = OpIEqual %bool %132 %uint_2 + OpSelectionMerge %134 None + OpBranchConditional %133 %135 %136 + %135 = OpLabel + %137 = OpLoad %bool %success + OpSelectionMerge %138 None + OpBranchConditional %137 %139 %138 + %139 = OpLabel + %141 = OpCompositeExtract %float %128 0 + %142 = OpCompositeExtract %float %124 0 + %140 = OpFunctionCall %bool %aboutEqual %141 %142 + OpBranch %138 + %138 = OpLabel + %143 = OpPhi %bool %137 %135 %140 %139 + OpSelectionMerge %144 None + OpBranchConditional %143 %145 %144 + %145 = OpLabel + %147 = OpCompositeExtract %float %128 1 + %148 = OpCompositeExtract %float %124 1 + %146 = OpFunctionCall %bool %aboutEqual %147 %148 + OpBranch %144 + %144 = OpLabel + %149 = OpPhi %bool %143 %138 %146 %145 + OpStore %success %149 + OpBranch %134 + %136 = OpLabel + %150 = OpLoad %bool %success + OpSelectionMerge %151 None + OpBranchConditional %150 %152 %151 + %152 = OpLabel + %154 = OpCompositeExtract %float %128 0 + %155 = OpCompositeExtract %float %124 0 + %153 = OpFunctionCall %bool %aboutEqual %154 %155 + OpBranch %151 + %151 = OpLabel + %156 = OpPhi %bool %150 %136 %153 %152 + OpSelectionMerge %157 None + OpBranchConditional %156 %158 %157 + %158 = OpLabel + %160 = OpCompositeExtract %float %128 1 + %161 = OpCompositeExtract %float %124 1 + %159 = OpFunctionCall %bool %aboutEqual %160 %161 + OpBranch %157 + %157 = OpLabel + %162 = OpPhi %bool %156 %151 %159 %158 + OpSelectionMerge %163 None + OpBranchConditional %162 %164 %163 + %164 = OpLabel + %166 = OpCompositeExtract %float %128 2 + %167 = OpCompositeExtract %float %124 2 + %165 = OpFunctionCall %bool %aboutEqual %166 %167 + OpBranch %163 + %163 = OpLabel + %168 = OpPhi %bool %162 %157 %165 %164 + OpSelectionMerge %169 None + OpBranchConditional %168 %170 %169 + %170 = OpLabel + %172 = OpCompositeExtract %float %128 3 + %173 = OpCompositeExtract %float %124 3 + %171 = OpFunctionCall %bool %aboutEqual %172 %173 + OpBranch %169 + %169 = OpLabel + %174 = OpPhi %bool %168 %163 %171 %170 + OpStore %success %174 + OpBranch %134 + %134 = OpLabel + OpBranch %87 + %87 = OpLabel + %176 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1 + %177 = OpLoad %uint %176 + %179 = OpCompositeExtract %int %38 0 + %178 = OpBitcast %uint %179 + %180 = OpIMul %uint %177 %178 + %181 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0 + %182 = OpLoad %uint %181 + %183 = OpIAdd %uint %180 %182 + %184 = OpLoad %bool %success + OpSelectionMerge %185 None + OpBranchConditional %184 %186 %187 + %186 = OpLabel + %189 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %183 + OpStore %189 %uint_1 + OpBranch %185 + %187 = OpLabel + %190 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %183 + OpStore %190 %uint_0 + OpBranch %185 + %185 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/bug/tint/913.wgsl.expected.wgsl b/test/bug/tint/913.wgsl.expected.wgsl new file mode 100644 index 0000000000..b38c00103d --- /dev/null +++ b/test/bug/tint/913.wgsl.expected.wgsl @@ -0,0 +1,55 @@ +[[block]] +struct Uniforms { + dstTextureFlipY : u32; + channelCount : u32; + srcCopyOrigin : vec2; + dstCopyOrigin : vec2; + copySize : vec2; +}; + +[[block]] +struct OutputBuf { + result : 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; + +fn aboutEqual(value : f32, expect : f32) -> bool { + return (abs((value - expect)) < 0.001); +} + +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { + let srcSize : vec2 = textureDimensions(src); + let dstSize : vec2 = textureDimensions(dst); + let dstTexCoord : vec2 = vec2(GlobalInvocationID.xy); + let nonCoveredColor : vec4 = vec4(0.0, 1.0, 0.0, 1.0); + var success : bool = true; + if (((((dstTexCoord.x < uniforms.dstCopyOrigin.x) || (dstTexCoord.y < uniforms.dstCopyOrigin.y)) || (dstTexCoord.x >= (uniforms.dstCopyOrigin.x + uniforms.copySize.x))) || (dstTexCoord.y >= (uniforms.dstCopyOrigin.y + uniforms.copySize.y)))) { + success = (success && all((textureLoad(dst, vec2(dstTexCoord), 0) == nonCoveredColor))); + } else { + var srcTexCoord : vec2 = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin); + if ((uniforms.dstTextureFlipY == 1u)) { + srcTexCoord.y = ((u32(srcSize.y) - srcTexCoord.y) - 1u); + } + let srcColor : vec4 = textureLoad(src, vec2(srcTexCoord), 0); + let dstColor : vec4 = textureLoad(dst, vec2(dstTexCoord), 0); + if ((uniforms.channelCount == 2u)) { + success = ((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g)); + } else { + success = ((((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g)) && aboutEqual(dstColor.b, srcColor.b)) && aboutEqual(dstColor.a, srcColor.a)); + } + } + let outputIndex : u32 = ((GlobalInvocationID.y * u32(dstSize.x)) + GlobalInvocationID.x); + if (success) { + output.result[outputIndex] = 1u; + } else { + output.result[outputIndex] = 0u; + } +}