From 5e2d8af6adf50b664a6984344b1435b6198fcf3d Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Wed, 30 Jun 2021 15:06:00 +0000 Subject: [PATCH] test: Add missing .expected files Theses were missing from https://dawn-review.googlesource.com/c/tint/+/56140 Change-Id: Iece1c69923024da2ce9473e205c18eb9568872b9 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56548 Kokoro: Kokoro Commit-Queue: Ben Clayton Reviewed-by: Antonio Maiorano --- test/bug/tint/534.wgsl.expected.hlsl | 67 ++++++++ test/bug/tint/534.wgsl.expected.msl | 51 ++++++ test/bug/tint/534.wgsl.expected.spvasm | 224 +++++++++++++++++++++++++ test/bug/tint/534.wgsl.expected.wgsl | 59 +++++++ 4 files changed, 401 insertions(+) create mode 100644 test/bug/tint/534.wgsl.expected.hlsl create mode 100644 test/bug/tint/534.wgsl.expected.msl create mode 100644 test/bug/tint/534.wgsl.expected.spvasm create mode 100644 test/bug/tint/534.wgsl.expected.wgsl diff --git a/test/bug/tint/534.wgsl.expected.hlsl b/test/bug/tint/534.wgsl.expected.hlsl new file mode 100644 index 0000000000..19f5f6bab0 --- /dev/null +++ b/test/bug/tint/534.wgsl.expected.hlsl @@ -0,0 +1,67 @@ +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; + } +} + +Texture2D src : register(t0, space0); +Texture2D tint_symbol : register(t1, space0); +RWByteAddressBuffer output : register(u2, space0); +cbuffer cbuffer_uniforms : register(b3, space0) { + uint4 uniforms[1]; +}; + +uint ConvertToFp16FloatValue(float fp32) { + return 1u; +} + +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); + int2 size = tint_tmp; + int2 dstTexCoord = int2(GlobalInvocationID.xy); + int2 srcTexCoord = dstTexCoord; + const int scalar_offset = (0u) / 4; + if ((uniforms[scalar_offset / 4][scalar_offset % 4] == 1u)) { + srcTexCoord.y = ((size.y - dstTexCoord.y) - 1); + } + float4 srcColor = src.Load(int3(srcTexCoord, 0)); + float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0)); + bool success = true; + uint4 srcColorBits = uint4(0u, 0u, 0u, 0u); + uint4 dstColorBits = uint4(dstColor); + { + uint i = 0u; + while (true) { + const int scalar_offset_1 = (12u) / 4; + if (!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))) { + break; + } + Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i])); + bool tint_tmp_1 = success; + if (tint_tmp_1) { + tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]); + } + success = (tint_tmp_1); + { + i = (i + 1u); + } + } + } + uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x); + if (success) { + output.Store((4u * outputIndex), asuint(uint(1))); + } else { + output.Store((4u * outputIndex), asuint(uint(0))); + } + return; +} diff --git a/test/bug/tint/534.wgsl.expected.msl b/test/bug/tint/534.wgsl.expected.msl new file mode 100644 index 0000000000..429c94b930 --- /dev/null +++ b/test/bug/tint/534.wgsl.expected.msl @@ -0,0 +1,51 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint dstTextureFlipY; + /* 0x0004 */ uint isFloat16; + /* 0x0008 */ uint isRGB10A2Unorm; + /* 0x000c */ uint channelCount; +}; +struct OutputBuf { + /* 0x0000 */ uint result[1]; +}; + +uint ConvertToFp16FloatValue(float fp32) { + return 1u; +} + +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 size = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height()); + int2 dstTexCoord = int2(GlobalInvocationID.xy); + int2 srcTexCoord = dstTexCoord; + if ((uniforms.dstTextureFlipY == 1u)) { + srcTexCoord.y = ((size.y - dstTexCoord.y) - 1); + } + float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0); + float4 dstColor = tint_symbol_3.read(uint2(dstTexCoord), 0); + bool success = true; + uint4 srcColorBits = 0u; + uint4 dstColorBits = uint4(dstColor); + { + uint i = 0u; + while (true) { + if (!((i < uniforms.channelCount))) { + break; + } + srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]); + success = (success && (srcColorBits[i] == dstColorBits[i])); + { + i = (i + 1u); + } + } + } + uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x); + if (success) { + output.result[outputIndex] = uint(1); + } else { + output.result[outputIndex] = uint(0); + } + return; +} + diff --git a/test/bug/tint/534.wgsl.expected.spvasm b/test/bug/tint/534.wgsl.expected.spvasm new file mode 100644 index 0000000000..9f983e6f02 --- /dev/null +++ b/test/bug/tint/534.wgsl.expected.spvasm @@ -0,0 +1,224 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 135 +; Schema: 0 + OpCapability Shader + OpCapability ImageQuery + 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 "isFloat16" + OpMemberName %Uniforms 2 "isRGB10A2Unorm" + OpMemberName %Uniforms 3 "channelCount" + OpName %uniforms "uniforms" + OpName %tint_symbol "tint_symbol" + OpName %ConvertToFp16FloatValue "ConvertToFp16FloatValue" + OpName %fp32 "fp32" + OpName %main "main" + OpName %size "size" + OpName %dstTexCoord "dstTexCoord" + OpName %srcTexCoord "srcTexCoord" + OpName %srcColor "srcColor" + OpName %dstColor "dstColor" + OpName %success "success" + OpName %srcColorBits "srcColorBits" + OpName %dstColorBits "dstColorBits" + OpName %i "i" + OpName %outputIndex "outputIndex" + 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 12 + 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 + %Uniforms = OpTypeStruct %uint %uint %uint %uint +%_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 + %17 = OpTypeFunction %uint %float + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %22 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 + %int_0 = OpConstant %int 0 +%_ptr_Function_v2int = OpTypePointer Function %v2int + %33 = OpConstantNull %v2int + %v2uint = OpTypeVector %uint 2 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %bool = OpTypeBool +%_ptr_Function_int = OpTypePointer Function %int + %int_1 = OpConstant %int 1 + %v4float = OpTypeVector %float 4 +%_ptr_Function_v4float = OpTypePointer Function %v4float + %64 = OpConstantNull %v4float + %true = OpConstantTrue %bool +%_ptr_Function_bool = OpTypePointer Function %bool + %72 = OpConstantNull %bool + %v4uint = OpTypeVector %uint 4 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %76 = OpConstantNull %v4uint +%_ptr_Function_uint = OpTypePointer Function %uint + %82 = OpConstantNull %uint + %uint_3 = OpConstant %uint 3 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Input_uint = OpTypePointer Input %uint +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint +%ConvertToFp16FloatValue = OpFunction %uint None %17 + %fp32 = OpFunctionParameter %float + %20 = OpLabel + OpReturnValue %uint_1 + OpFunctionEnd + %main = OpFunction %void None %22 + %25 = OpLabel + %size = OpVariable %_ptr_Function_v2int Function %33 +%dstTexCoord = OpVariable %_ptr_Function_v2int Function %33 +%srcTexCoord = OpVariable %_ptr_Function_v2int Function %33 + %srcColor = OpVariable %_ptr_Function_v4float Function %64 + %dstColor = OpVariable %_ptr_Function_v4float Function %64 + %success = OpVariable %_ptr_Function_bool Function %72 +%srcColorBits = OpVariable %_ptr_Function_v4uint Function %76 +%dstColorBits = OpVariable %_ptr_Function_v4uint Function %76 + %i = OpVariable %_ptr_Function_uint Function %82 +%outputIndex = OpVariable %_ptr_Function_uint Function %82 + %29 = OpLoad %3 %src + %26 = OpImageQuerySizeLod %v2int %29 %int_0 + OpStore %size %26 + %36 = OpLoad %v3uint %tint_symbol + %37 = OpVectorShuffle %v2uint %36 %36 0 1 + %34 = OpBitcast %v2int %37 + OpStore %dstTexCoord %34 + %39 = OpLoad %v2int %dstTexCoord + OpStore %srcTexCoord %39 + %43 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %44 = OpLoad %uint %43 + %45 = OpIEqual %bool %44 %uint_1 + OpSelectionMerge %47 None + OpBranchConditional %45 %48 %47 + %48 = OpLabel + %50 = OpAccessChain %_ptr_Function_int %srcTexCoord %uint_1 + %51 = OpAccessChain %_ptr_Function_int %size %uint_1 + %52 = OpLoad %int %51 + %53 = OpAccessChain %_ptr_Function_int %dstTexCoord %uint_1 + %54 = OpLoad %int %53 + %55 = OpISub %int %52 %54 + %57 = OpISub %int %55 %int_1 + OpStore %50 %57 + OpBranch %47 + %47 = OpLabel + %60 = OpLoad %3 %src + %61 = OpLoad %v2int %srcTexCoord + %58 = OpImageFetch %v4float %60 %61 Lod %int_0 + OpStore %srcColor %58 + %66 = OpLoad %3 %dst + %67 = OpLoad %v2int %dstTexCoord + %65 = OpImageFetch %v4float %66 %67 Lod %int_0 + OpStore %dstColor %65 + OpStore %success %true + %78 = OpLoad %v4float %dstColor + %77 = OpConvertFToU %v4uint %78 + OpStore %dstColorBits %77 + OpStore %i %uint_0 + OpBranch %83 + %83 = OpLabel + OpLoopMerge %84 %85 None + OpBranch %86 + %86 = OpLabel + %88 = OpLoad %uint %i + %90 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 + %91 = OpLoad %uint %90 + %92 = OpULessThan %bool %88 %91 + %87 = OpLogicalNot %bool %92 + OpSelectionMerge %93 None + OpBranchConditional %87 %94 %93 + %94 = OpLabel + OpBranch %84 + %93 = OpLabel + %95 = OpLoad %uint %i + %96 = OpAccessChain %_ptr_Function_uint %srcColorBits %95 + %98 = OpLoad %uint %i + %100 = OpAccessChain %_ptr_Function_float %srcColor %98 + %101 = OpLoad %float %100 + %97 = OpFunctionCall %uint %ConvertToFp16FloatValue %101 + OpStore %96 %97 + %102 = OpLoad %bool %success + OpSelectionMerge %103 None + OpBranchConditional %102 %104 %103 + %104 = OpLabel + %105 = OpLoad %uint %i + %106 = OpAccessChain %_ptr_Function_uint %srcColorBits %105 + %107 = OpLoad %uint %106 + %108 = OpLoad %uint %i + %109 = OpAccessChain %_ptr_Function_uint %dstColorBits %108 + %110 = OpLoad %uint %109 + %111 = OpIEqual %bool %107 %110 + OpBranch %103 + %103 = OpLabel + %112 = OpPhi %bool %102 %93 %111 %104 + OpStore %success %112 + OpBranch %85 + %85 = OpLabel + %113 = OpLoad %uint %i + %114 = OpIAdd %uint %113 %uint_1 + OpStore %i %114 + OpBranch %83 + %84 = OpLabel + %116 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1 + %117 = OpLoad %uint %116 + %119 = OpAccessChain %_ptr_Function_int %size %uint_0 + %120 = OpLoad %int %119 + %118 = OpBitcast %uint %120 + %121 = OpIMul %uint %117 %118 + %122 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0 + %123 = OpLoad %uint %122 + %124 = OpIAdd %uint %121 %123 + OpStore %outputIndex %124 + %126 = OpLoad %bool %success + OpSelectionMerge %127 None + OpBranchConditional %126 %128 %129 + %128 = OpLabel + %130 = OpLoad %uint %outputIndex + %132 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %130 + OpStore %132 %uint_1 + OpBranch %127 + %129 = OpLabel + %133 = OpLoad %uint %outputIndex + %134 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %133 + OpStore %134 %uint_0 + OpBranch %127 + %127 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/bug/tint/534.wgsl.expected.wgsl b/test/bug/tint/534.wgsl.expected.wgsl new file mode 100644 index 0000000000..91a330b77d --- /dev/null +++ b/test/bug/tint/534.wgsl.expected.wgsl @@ -0,0 +1,59 @@ +[[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; + +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); + { + var i : u32 = 0u; + loop { + if (!((i < uniforms.channelCount))) { + break; + } + srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]); + success = (success && (srcColorBits[i] == dstColorBits[i])); + + continuing { + i = (i + 1u); + } + } + } + var outputIndex : u32 = ((GlobalInvocationID.y * u32(size.x)) + GlobalInvocationID.x); + if (success) { + output.result[outputIndex] = u32(1); + } else { + output.result[outputIndex] = u32(0); + } +}