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 <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
Ben Clayton 2021-07-15 22:36:44 +00:00 committed by Tint LUCI CQ
parent 0bff3fb3b7
commit 1da4073870
5 changed files with 523 additions and 0 deletions

153
test/bug/tint/977.spvasm Normal file
View File

@ -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

View File

@ -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;
}

View File

@ -0,0 +1,71 @@
SKIP: FAILED
#include <metal_stdlib>
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<int>(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))) {
~~~~ ^ ~~~~

View File

@ -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

View File

@ -0,0 +1,80 @@
type RTArr = [[stride(4)]] array<f32>;
type RTArr_1 = [[stride(4)]] array<f32>;
[[block]]
struct ResultMatrix {
numbers : RTArr_1;
};
type RTArr_2 = [[stride(4)]] array<f32>;
[[block]]
struct FirstMatrix {
numbers : RTArr_1;
};
[[block]]
struct SecondMatrix {
numbers : RTArr_1;
};
[[block]]
struct Uniforms {
NAN : f32;
sizeA : i32;
sizeB : i32;
};
var<private> gl_GlobalInvocationID : vec3<u32>;
[[group(0), binding(2)]] var<storage, read_write> resultMatrix : ResultMatrix;
[[group(0), binding(0)]] var<storage, read> firstMatrix : FirstMatrix;
[[group(0), binding(1)]] var<storage, read> secondMatrix : SecondMatrix;
[[group(0), binding(3)]] var<uniform> x_46 : Uniforms;
fn binaryOperation_f1_f1_(a : ptr<function, f32>, b : ptr<function, f32>) -> 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<i32>(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<u32>) {
gl_GlobalInvocationID = gl_GlobalInvocationID_param;
main_1();
}