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 <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
diff --git a/test/bug/tint/534.wgsl.expected.hlsl b/test/bug/tint/534.wgsl.expected.hlsl
new file mode 100644
index 0000000..19f5f6b
--- /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<float4> src : register(t0, space0);
+Texture2D<float4> 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 0000000..429c94b
--- /dev/null
+++ b/test/bug/tint/534.wgsl.expected.msl
@@ -0,0 +1,51 @@
+#include <metal_stdlib>
+
+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<float, access::sample> tint_symbol_2 [[texture(0)]], texture2d<float, access::sample> 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 0000000..9f983e6
--- /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 0000000..91a330b
--- /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<u32>;
+};
+
+[[group(0), binding(0)]] var src : texture_2d<f32>;
+
+[[group(0), binding(1)]] var dst : texture_2d<f32>;
+
+[[group(0), binding(2)]] var<storage, read_write> output : OutputBuf;
+
+[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
+
+fn ConvertToFp16FloatValue(fp32 : f32) -> u32 {
+ return 1u;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ var size : vec2<i32> = textureDimensions(src);
+ var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
+ var srcTexCoord : vec2<i32> = dstTexCoord;
+ if ((uniforms.dstTextureFlipY == 1u)) {
+ srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
+ }
+ var srcColor : vec4<f32> = textureLoad(src, srcTexCoord, 0);
+ var dstColor : vec4<f32> = textureLoad(dst, dstTexCoord, 0);
+ var success : bool = true;
+ var srcColorBits : vec4<u32>;
+ var dstColorBits : vec4<u32> = vec4<u32>(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);
+ }
+}
