From 1da4073870c3b3da233cf0985372fa4f7815bc3c Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Thu, 15 Jul 2021 22:36:44 +0000 Subject: [PATCH] test: Add case for tint:977 Bug: tint:977 Change-Id: I50778c6e1778717c0ad9b02b2ea25b13c4a3da97 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58065 Auto-Submit: Ben Clayton Reviewed-by: Antonio Maiorano Commit-Queue: Ben Clayton Kokoro: Kokoro --- test/bug/tint/977.spvasm | 153 +++++++++++++++++++++ test/bug/tint/977.spvasm.expected.hlsl | 55 ++++++++ test/bug/tint/977.spvasm.expected.msl | 71 ++++++++++ test/bug/tint/977.spvasm.expected.spvasm | 164 +++++++++++++++++++++++ test/bug/tint/977.spvasm.expected.wgsl | 80 +++++++++++ 5 files changed, 523 insertions(+) create mode 100644 test/bug/tint/977.spvasm create mode 100644 test/bug/tint/977.spvasm.expected.hlsl create mode 100644 test/bug/tint/977.spvasm.expected.msl create mode 100644 test/bug/tint/977.spvasm.expected.spvasm create mode 100644 test/bug/tint/977.spvasm.expected.wgsl diff --git a/test/bug/tint/977.spvasm b/test/bug/tint/977.spvasm new file mode 100644 index 0000000000..dfd5ab6cae --- /dev/null +++ b/test/bug/tint/977.spvasm @@ -0,0 +1,153 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 10 +; Bound: 82 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %main "main" + OpName %binaryOperation_f1_f1_ "binaryOperation(f1;f1;" + OpName %a "a" + OpName %b "b" + OpName %index "index" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %a_0 "a" + OpName %ResultMatrix "ResultMatrix" + OpMemberName %ResultMatrix 0 "numbers" + OpName %resultMatrix "resultMatrix" + OpName %param "param" + OpName %param_0 "param" + OpName %FirstMatrix "FirstMatrix" + OpMemberName %FirstMatrix 0 "numbers" + OpName %firstMatrix "firstMatrix" + OpName %SecondMatrix "SecondMatrix" + OpMemberName %SecondMatrix 0 "numbers" + OpName %secondMatrix "secondMatrix" + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "NAN" + OpMemberName %Uniforms 1 "sizeA" + OpMemberName %Uniforms 2 "sizeB" + OpName %_ "" + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %_runtimearr_float ArrayStride 4 + OpMemberDecorate %ResultMatrix 0 Offset 0 + OpDecorate %ResultMatrix BufferBlock + OpDecorate %resultMatrix DescriptorSet 0 + OpDecorate %resultMatrix Binding 2 + OpDecorate %_runtimearr_float_0 ArrayStride 4 + OpMemberDecorate %FirstMatrix 0 NonWritable + OpMemberDecorate %FirstMatrix 0 Offset 0 + OpDecorate %FirstMatrix BufferBlock + OpDecorate %firstMatrix DescriptorSet 0 + OpDecorate %firstMatrix Binding 0 + OpDecorate %_runtimearr_float_1 ArrayStride 4 + OpMemberDecorate %SecondMatrix 0 NonWritable + OpMemberDecorate %SecondMatrix 0 Offset 0 + OpDecorate %SecondMatrix BufferBlock + OpDecorate %secondMatrix DescriptorSet 0 + OpDecorate %secondMatrix Binding 1 + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpMemberDecorate %Uniforms 2 Offset 8 + OpDecorate %Uniforms Block + OpDecorate %_ DescriptorSet 0 + OpDecorate %_ Binding 3 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 + %_ptr_Function_float = OpTypePointer Function %float + %8 = OpTypeFunction %float %_ptr_Function_float %_ptr_Function_float + %float_0 = OpConstant %float 0 + %bool = OpTypeBool + %float_1 = OpConstant %float 1 + %float_2 = OpConstant %float 2 + %int = OpTypeInt 32 1 + %_ptr_Function_int = OpTypePointer Function %int + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %_ptr_Input_v3uint = OpTypePointer Input %v3uint + %gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %uint_0 = OpConstant %uint 0 + %_ptr_Input_uint = OpTypePointer Input %uint + %int_n10 = OpConstant %int -10 + %_runtimearr_float = OpTypeRuntimeArray %float + %ResultMatrix = OpTypeStruct %_runtimearr_float +%_ptr_Uniform_ResultMatrix = OpTypePointer Uniform %ResultMatrix + %resultMatrix = OpVariable %_ptr_Uniform_ResultMatrix Uniform + %int_0 = OpConstant %int 0 + %float_n4 = OpConstant %float -4 + %float_n3 = OpConstant %float -3 + %_ptr_Uniform_float = OpTypePointer Uniform %float + %_runtimearr_float_0 = OpTypeRuntimeArray %float + %FirstMatrix = OpTypeStruct %_runtimearr_float_0 + %_ptr_Uniform_FirstMatrix = OpTypePointer Uniform %FirstMatrix + %firstMatrix = OpVariable %_ptr_Uniform_FirstMatrix Uniform + %_runtimearr_float_1 = OpTypeRuntimeArray %float + %SecondMatrix = OpTypeStruct %_runtimearr_float_1 +%_ptr_Uniform_SecondMatrix = OpTypePointer Uniform %SecondMatrix + %secondMatrix = OpVariable %_ptr_Uniform_SecondMatrix Uniform + %Uniforms = OpTypeStruct %float %int %int + %_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %_ = OpVariable %_ptr_Uniform_Uniforms Uniform + %main = OpFunction %void None %3 + %5 = OpLabel + %index = OpVariable %_ptr_Function_int Function + %a_0 = OpVariable %_ptr_Function_int Function + %param = OpVariable %_ptr_Function_float Function + %param_0 = OpVariable %_ptr_Function_float Function + %53 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %54 = OpLoad %uint %53 + %55 = OpBitcast %int %54 + OpStore %index %55 + OpStore %a_0 %int_n10 + %63 = OpLoad %int %index + OpStore %param %float_n4 + OpStore %param_0 %float_n3 + %68 = OpFunctionCall %float %binaryOperation_f1_f1_ %param %param_0 + %70 = OpAccessChain %_ptr_Uniform_float %resultMatrix %int_0 %63 + OpStore %70 %68 + OpReturn + OpFunctionEnd + %binaryOperation_f1_f1_ = OpFunction %float None %8 + %a = OpFunctionParameter %_ptr_Function_float + %b = OpFunctionParameter %_ptr_Function_float + %12 = OpLabel + %26 = OpVariable %_ptr_Function_float Function + %13 = OpLoad %float %b + %16 = OpFOrdEqual %bool %13 %float_0 + OpSelectionMerge %18 None + OpBranchConditional %16 %17 %18 + %17 = OpLabel + OpReturnValue %float_1 + %18 = OpLabel + %21 = OpLoad %float %b + %23 = OpFMod %float %21 %float_2 + %24 = OpExtInst %float %1 Round %23 + %25 = OpFUnordNotEqual %bool %24 %float_1 + OpSelectionMerge %28 None + OpBranchConditional %25 %27 %33 + %27 = OpLabel + %29 = OpLoad %float %a + %30 = OpExtInst %float %1 FAbs %29 + %31 = OpLoad %float %b + %32 = OpExtInst %float %1 Pow %30 %31 + OpStore %26 %32 + OpBranch %28 + %33 = OpLabel + %34 = OpLoad %float %a + %35 = OpExtInst %float %1 FSign %34 + %36 = OpLoad %float %a + %37 = OpExtInst %float %1 FAbs %36 + %38 = OpLoad %float %b + %39 = OpExtInst %float %1 Pow %37 %38 + %40 = OpFMul %float %35 %39 + OpStore %26 %40 + OpBranch %28 + %28 = OpLabel + %41 = OpLoad %float %26 + OpReturnValue %41 + OpFunctionEnd diff --git a/test/bug/tint/977.spvasm.expected.hlsl b/test/bug/tint/977.spvasm.expected.hlsl new file mode 100644 index 0000000000..c37348c381 --- /dev/null +++ b/test/bug/tint/977.spvasm.expected.hlsl @@ -0,0 +1,55 @@ +static uint3 gl_GlobalInvocationID = uint3(0u, 0u, 0u); +RWByteAddressBuffer resultMatrix : register(u2, space0); +ByteAddressBuffer firstMatrix : register(t0, space0); +ByteAddressBuffer secondMatrix : register(t1, space0); +cbuffer cbuffer_x_46 : register(b3, space0) { + uint4 x_46[1]; +}; + +float binaryOperation_f1_f1_(inout float a, inout float b) { + float x_26 = 0.0f; + const float x_13 = b; + if ((x_13 == 0.0f)) { + return 1.0f; + } + const float x_21 = b; + if (!((round((x_21 % 2.0f)) == 1.0f))) { + const float x_29 = a; + const float x_31 = b; + x_26 = pow(abs(x_29), x_31); + } else { + const float x_34 = a; + const float x_36 = a; + const float x_38 = b; + x_26 = (sign(x_34) * pow(abs(x_36), x_38)); + } + return x_26; +} + +void main_1() { + int index = 0; + int a_1 = 0; + float param = 0.0f; + float param_1 = 0.0f; + const uint x_54 = gl_GlobalInvocationID.x; + index = asint(x_54); + a_1 = -10; + const int x_63 = index; + param = -4.0f; + param_1 = -3.0f; + const float x_68 = binaryOperation_f1_f1_(param, param_1); + resultMatrix.Store((4u * uint(x_63)), asuint(x_68)); + return; +} + +struct tint_symbol_1 { + uint3 gl_GlobalInvocationID_param : SV_DispatchThreadID; +}; + +[numthreads(1, 1, 1)] +void main(tint_symbol_1 tint_symbol) { + const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param; + gl_GlobalInvocationID = gl_GlobalInvocationID_param; + main_1(); + return; +} diff --git a/test/bug/tint/977.spvasm.expected.msl b/test/bug/tint/977.spvasm.expected.msl new file mode 100644 index 0000000000..9662b9384d --- /dev/null +++ b/test/bug/tint/977.spvasm.expected.msl @@ -0,0 +1,71 @@ +SKIP: FAILED + +#include + +using namespace metal; +struct ResultMatrix { + /* 0x0000 */ float numbers[1]; +}; +struct FirstMatrix { + /* 0x0000 */ float numbers[1]; +}; +struct SecondMatrix { + /* 0x0000 */ float numbers[1]; +}; +struct Uniforms { + /* 0x0000 */ float tint_symbol; + /* 0x0004 */ int sizeA; + /* 0x0008 */ int sizeB; +}; + +float binaryOperation_f1_f1_(thread float* const a, thread float* const b) { + float x_26 = 0.0f; + float const x_13 = *(b); + if ((x_13 == 0.0f)) { + return 1.0f; + } + float const x_21 = *(b); + if (!((rint((x_21 % 2.0f)) == 1.0f))) { + float const x_29 = *(a); + float const x_31 = *(b); + x_26 = pow(fabs(x_29), x_31); + } else { + float const x_34 = *(a); + float const x_36 = *(a); + float const x_38 = *(b); + x_26 = (sign(x_34) * pow(fabs(x_36), x_38)); + } + float const x_41 = x_26; + return x_41; +} + +void main_1(device ResultMatrix& resultMatrix, thread uint3* const tint_symbol_3) { + int index = 0; + int a_1 = 0; + float param = 0.0f; + float param_1 = 0.0f; + uint const x_54 = (*(tint_symbol_3)).x; + index = as_type(x_54); + a_1 = -10; + int const x_63 = index; + param = -4.0f; + param_1 = -3.0f; + float const x_68 = binaryOperation_f1_f1_(&(param), &(param_1)); + resultMatrix.numbers[x_63] = x_68; + return; +} + +kernel void tint_symbol_1(uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], device ResultMatrix& resultMatrix [[buffer(2)]]) { + thread uint3 tint_symbol_4 = 0u; + tint_symbol_4 = gl_GlobalInvocationID_param; + main_1(resultMatrix, &(tint_symbol_4)); + return; +} + +Compilation failed: + +program_source:26:21: error: invalid operands to binary expression ('const float' and 'float') + if (!((rint((x_21 % 2.0f)) == 1.0f))) { + ~~~~ ^ ~~~~ + + diff --git a/test/bug/tint/977.spvasm.expected.spvasm b/test/bug/tint/977.spvasm.expected.spvasm new file mode 100644 index 0000000000..76dd37d7a7 --- /dev/null +++ b/test/bug/tint/977.spvasm.expected.spvasm @@ -0,0 +1,164 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 95 +; Schema: 0 + OpCapability Shader + %43 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %tint_symbol + OpExecutionMode %main LocalSize 1 1 1 + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %ResultMatrix "ResultMatrix" + OpMemberName %ResultMatrix 0 "numbers" + OpName %resultMatrix "resultMatrix" + OpName %FirstMatrix "FirstMatrix" + OpMemberName %FirstMatrix 0 "numbers" + OpName %firstMatrix "firstMatrix" + OpName %SecondMatrix "SecondMatrix" + OpMemberName %SecondMatrix 0 "numbers" + OpName %secondMatrix "secondMatrix" + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "NAN" + OpMemberName %Uniforms 1 "sizeA" + OpMemberName %Uniforms 2 "sizeB" + OpName %x_46 "x_46" + OpName %tint_symbol "tint_symbol" + OpName %binaryOperation_f1_f1_ "binaryOperation_f1_f1_" + OpName %a "a" + OpName %b "b" + OpName %x_26 "x_26" + OpName %main_1 "main_1" + OpName %index "index" + OpName %a_1 "a_1" + OpName %param "param" + OpName %param_1 "param_1" + OpName %main "main" + OpDecorate %ResultMatrix Block + OpMemberDecorate %ResultMatrix 0 Offset 0 + OpDecorate %_runtimearr_float ArrayStride 4 + OpDecorate %resultMatrix DescriptorSet 0 + OpDecorate %resultMatrix Binding 2 + OpDecorate %FirstMatrix Block + OpMemberDecorate %FirstMatrix 0 Offset 0 + OpDecorate %firstMatrix NonWritable + OpDecorate %firstMatrix DescriptorSet 0 + OpDecorate %firstMatrix Binding 0 + OpDecorate %SecondMatrix Block + OpMemberDecorate %SecondMatrix 0 Offset 0 + OpDecorate %secondMatrix NonWritable + OpDecorate %secondMatrix DescriptorSet 0 + OpDecorate %secondMatrix Binding 1 + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpMemberDecorate %Uniforms 2 Offset 8 + OpDecorate %x_46 NonWritable + OpDecorate %x_46 DescriptorSet 0 + OpDecorate %x_46 Binding 3 + OpDecorate %tint_symbol BuiltIn GlobalInvocationId + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint + %5 = OpConstantNull %v3uint +%gl_GlobalInvocationID = OpVariable %_ptr_Private_v3uint Private %5 + %float = OpTypeFloat 32 +%_runtimearr_float = OpTypeRuntimeArray %float +%ResultMatrix = OpTypeStruct %_runtimearr_float +%_ptr_StorageBuffer_ResultMatrix = OpTypePointer StorageBuffer %ResultMatrix +%resultMatrix = OpVariable %_ptr_StorageBuffer_ResultMatrix StorageBuffer +%FirstMatrix = OpTypeStruct %_runtimearr_float +%_ptr_StorageBuffer_FirstMatrix = OpTypePointer StorageBuffer %FirstMatrix +%firstMatrix = OpVariable %_ptr_StorageBuffer_FirstMatrix StorageBuffer +%SecondMatrix = OpTypeStruct %_runtimearr_float +%_ptr_StorageBuffer_SecondMatrix = OpTypePointer StorageBuffer %SecondMatrix +%secondMatrix = OpVariable %_ptr_StorageBuffer_SecondMatrix StorageBuffer + %int = OpTypeInt 32 1 + %Uniforms = OpTypeStruct %float %int %int +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %x_46 = OpVariable %_ptr_Uniform_Uniforms Uniform +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%tint_symbol = OpVariable %_ptr_Input_v3uint Input +%_ptr_Function_float = OpTypePointer Function %float + %23 = OpTypeFunction %float %_ptr_Function_float %_ptr_Function_float + %30 = OpConstantNull %float + %float_0 = OpConstant %float 0 + %bool = OpTypeBool + %float_1 = OpConstant %float 1 + %float_2 = OpConstant %float 2 + %void = OpTypeVoid + %67 = OpTypeFunction %void +%_ptr_Function_int = OpTypePointer Function %int + %73 = OpConstantNull %int + %uint_0 = OpConstant %uint 0 +%_ptr_Private_uint = OpTypePointer Private %uint + %int_n10 = OpConstant %int -10 + %float_n4 = OpConstant %float -4 + %float_n3 = OpConstant %float -3 +%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float +%binaryOperation_f1_f1_ = OpFunction %float None %23 + %a = OpFunctionParameter %_ptr_Function_float + %b = OpFunctionParameter %_ptr_Function_float + %28 = OpLabel + %x_26 = OpVariable %_ptr_Function_float Function %30 + %32 = OpLoad %float %b + %34 = OpFOrdEqual %bool %32 %float_0 + OpSelectionMerge %36 None + OpBranchConditional %34 %37 %36 + %37 = OpLabel + OpReturnValue %float_1 + %36 = OpLabel + %40 = OpLoad %float %b + %45 = OpFMod %float %40 %float_2 + %42 = OpExtInst %float %43 RoundEven %45 + %46 = OpFOrdEqual %bool %42 %float_1 + %41 = OpLogicalNot %bool %46 + OpSelectionMerge %47 None + OpBranchConditional %41 %48 %49 + %48 = OpLabel + %51 = OpLoad %float %a + %53 = OpLoad %float %b + %55 = OpExtInst %float %43 FAbs %51 + %54 = OpExtInst %float %43 Pow %55 %53 + OpStore %x_26 %54 + OpBranch %47 + %49 = OpLabel + %57 = OpLoad %float %a + %59 = OpLoad %float %a + %61 = OpLoad %float %b + %62 = OpExtInst %float %43 FSign %57 + %64 = OpExtInst %float %43 FAbs %59 + %63 = OpExtInst %float %43 Pow %64 %61 + %65 = OpFMul %float %62 %63 + OpStore %x_26 %65 + OpBranch %47 + %47 = OpLabel + %66 = OpLoad %float %x_26 + OpReturnValue %66 + OpFunctionEnd + %main_1 = OpFunction %void None %67 + %70 = OpLabel + %index = OpVariable %_ptr_Function_int Function %73 + %a_1 = OpVariable %_ptr_Function_int Function %73 + %param = OpVariable %_ptr_Function_float Function %30 + %param_1 = OpVariable %_ptr_Function_float Function %30 + %79 = OpAccessChain %_ptr_Private_uint %gl_GlobalInvocationID %uint_0 + %80 = OpLoad %uint %79 + %81 = OpBitcast %int %80 + OpStore %index %81 + OpStore %a_1 %int_n10 + %83 = OpLoad %int %index + OpStore %param %float_n4 + OpStore %param_1 %float_n3 + %86 = OpFunctionCall %float %binaryOperation_f1_f1_ %param %param_1 + %90 = OpAccessChain %_ptr_StorageBuffer_float %resultMatrix %uint_0 %83 + OpStore %90 %86 + OpReturn + OpFunctionEnd + %main = OpFunction %void None %67 + %92 = OpLabel + %93 = OpLoad %v3uint %tint_symbol + OpStore %gl_GlobalInvocationID %93 + %94 = OpFunctionCall %void %main_1 + OpReturn + OpFunctionEnd diff --git a/test/bug/tint/977.spvasm.expected.wgsl b/test/bug/tint/977.spvasm.expected.wgsl new file mode 100644 index 0000000000..b9f67d48c6 --- /dev/null +++ b/test/bug/tint/977.spvasm.expected.wgsl @@ -0,0 +1,80 @@ +type RTArr = [[stride(4)]] array; + +type RTArr_1 = [[stride(4)]] array; + +[[block]] +struct ResultMatrix { + numbers : RTArr_1; +}; + +type RTArr_2 = [[stride(4)]] array; + +[[block]] +struct FirstMatrix { + numbers : RTArr_1; +}; + +[[block]] +struct SecondMatrix { + numbers : RTArr_1; +}; + +[[block]] +struct Uniforms { + NAN : f32; + sizeA : i32; + sizeB : i32; +}; + +var gl_GlobalInvocationID : vec3; + +[[group(0), binding(2)]] var resultMatrix : ResultMatrix; + +[[group(0), binding(0)]] var firstMatrix : FirstMatrix; + +[[group(0), binding(1)]] var secondMatrix : SecondMatrix; + +[[group(0), binding(3)]] var x_46 : Uniforms; + +fn binaryOperation_f1_f1_(a : ptr, b : ptr) -> f32 { + var x_26 : f32; + let x_13 : f32 = *(b); + if ((x_13 == 0.0)) { + return 1.0; + } + let x_21 : f32 = *(b); + if (!((round((x_21 % 2.0)) == 1.0))) { + let x_29 : f32 = *(a); + let x_31 : f32 = *(b); + x_26 = pow(abs(x_29), x_31); + } else { + let x_34 : f32 = *(a); + let x_36 : f32 = *(a); + let x_38 : f32 = *(b); + x_26 = (sign(x_34) * pow(abs(x_36), x_38)); + } + let x_41 : f32 = x_26; + return x_41; +} + +fn main_1() { + var index : i32; + var a_1 : i32; + var param : f32; + var param_1 : f32; + let x_54 : u32 = gl_GlobalInvocationID.x; + index = bitcast(x_54); + a_1 = -10; + let x_63 : i32 = index; + param = -4.0; + param_1 = -3.0; + let x_68 : f32 = binaryOperation_f1_f1_(&(param), &(param_1)); + resultMatrix.numbers[x_63] = x_68; + return; +} + +[[stage(compute), workgroup_size(1, 1, 1)]] +fn main([[builtin(global_invocation_id)]] gl_GlobalInvocationID_param : vec3) { + gl_GlobalInvocationID = gl_GlobalInvocationID_param; + main_1(); +}