Fix FXC compile errors on modulo by zero

Just like for divide, FXC fails with the exact same error when
performing a modulo on a value that FXC determines to be zero. We
address it in the same way as we do for divide.

This also fixes a couple of the vk-gl-cts tests for which I manually
generated expectation files for.

Bug: tint:1083
Change-Id: Ia388bf002112afded542adb791d37e88e35a77ff
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/74220
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
This commit is contained in:
Antonio Maiorano
2021-12-22 15:02:09 +00:00
committed by Tint LUCI CQ
parent cc4d97b6e3
commit 9943de6813
159 changed files with 1965 additions and 69 deletions

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1.;
let b = 0.;
let r : f32 = a % b;
}

View File

@@ -0,0 +1,5 @@
[numthreads(1, 1, 1)]
void f() {
const float r = (1.0f % 0.0f);
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float const a = 1.0f;
float const b = 0.0f;
float const r = fmod(a, b);
return;
}

View File

@@ -0,0 +1,20 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 9
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%float_1 = OpConstant %float 1
%float_0 = OpConstant %float 0
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpFRem %float %float_1 %float_0
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1.0;
let b = 0.0;
let r : f32 = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1;
let b = 0;
let r : i32 = a % b;
}

View File

@@ -0,0 +1,5 @@
[numthreads(1, 1, 1)]
void f() {
const int r = (1 % 1);
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int const a = 1;
int const b = 0;
int const r = (a % b);
return;
}

View File

@@ -0,0 +1,20 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 9
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%int_0 = OpConstant %int 0
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpSMod %int %int_1 %int_0
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1;
let b = 0;
let r : i32 = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1u;
let b = 0u;
let r : u32 = a % b;
}

View File

@@ -0,0 +1,5 @@
[numthreads(1, 1, 1)]
void f() {
const uint r = (1u % 1u);
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint const a = 1u;
uint const b = 0u;
uint const r = (a % b);
return;
}

View File

@@ -0,0 +1,20 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 9
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpUMod %uint %uint_1 %uint_0
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1u;
let b = 0u;
let r : u32 = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4;
let b = vec3<i32>(0, 2, 0);
let r : vec3<i32> = a % b;
}

View File

@@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const int a = 4;
const int3 b = int3(0, 2, 0);
const int3 r = (a % int3(1, 2, 1));
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int const a = 4;
int3 const b = int3(0, 2, 0);
int3 const r = (a % b);
return;
}

View File

@@ -0,0 +1,27 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 16
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%int_4 = OpConstant %int 4
%v3int = OpTypeVector %int 3
%int_0 = OpConstant %int 0
%int_2 = OpConstant %int 2
%10 = OpConstantComposite %v3int %int_0 %int_2 %int_0
%_ptr_Function_v3int = OpTypePointer Function %v3int
%14 = OpConstantNull %v3int
%f = OpFunction %void None %1
%4 = OpLabel
%12 = OpVariable %_ptr_Function_v3int Function %14
%15 = OpCompositeConstruct %v3int %int_4 %int_4 %int_4
%11 = OpSMod %v3int %15 %10
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4;
let b = vec3<i32>(0, 2, 0);
let r : vec3<i32> = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4u;
let b = vec3<u32>(0u, 2u, 0u);
let r : vec3<u32> = a % b;
}

View File

@@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const uint a = 4u;
const uint3 b = uint3(0u, 2u, 0u);
const uint3 r = (a % uint3(1u, 2u, 1u));
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint const a = 4u;
uint3 const b = uint3(0u, 2u, 0u);
uint3 const r = (a % b);
return;
}

View File

@@ -0,0 +1,27 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 16
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%v3uint = OpTypeVector %uint 3
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%10 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%14 = OpConstantNull %v3uint
%f = OpFunction %void None %1
%4 = OpLabel
%12 = OpVariable %_ptr_Function_v3uint Function %14
%15 = OpCompositeConstruct %v3uint %uint_4 %uint_4 %uint_4
%11 = OpUMod %v3uint %15 %10
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4u;
let b = vec3<u32>(0u, 2u, 0u);
let r : vec3<u32> = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = 0;
let r : vec3<i32> = a % b;
}

View File

@@ -0,0 +1,6 @@
[numthreads(1, 1, 1)]
void f() {
const int3 a = int3(1, 2, 3);
const int3 r = (a % 1);
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 0;
int3 const r = (a % b);
return;
}

View File

@@ -0,0 +1,28 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v3int = OpTypeVector %int 3
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%int_0 = OpConstant %int 0
%_ptr_Function_v3int = OpTypePointer Function %v3int
%15 = OpConstantNull %v3int
%f = OpFunction %void None %1
%4 = OpLabel
%13 = OpVariable %_ptr_Function_v3int Function %15
%16 = OpCompositeConstruct %v3int %int_0 %int_0 %int_0
%12 = OpSMod %v3int %10 %16
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = 0;
let r : vec3<i32> = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = 0u;
let r : vec3<u32> = a % b;
}

View File

@@ -0,0 +1,6 @@
[numthreads(1, 1, 1)]
void f() {
const uint3 a = uint3(1u, 2u, 3u);
const uint3 r = (a % 1u);
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint3 const a = uint3(1u, 2u, 3u);
uint const b = 0u;
uint3 const r = (a % b);
return;
}

View File

@@ -0,0 +1,28 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%uint_0 = OpConstant %uint 0
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%15 = OpConstantNull %v3uint
%f = OpFunction %void None %1
%4 = OpLabel
%13 = OpVariable %_ptr_Function_v3uint Function %15
%16 = OpCompositeConstruct %v3uint %uint_0 %uint_0 %uint_0
%12 = OpUMod %v3uint %10 %16
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = 0u;
let r : vec3<u32> = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<f32>(1., 2., 3.);
let b = vec3<f32>(0., 5., 0.);
let r : vec3<f32> = a % b;
}

View File

@@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const float3 a = float3(1.0f, 2.0f, 3.0f);
const float3 b = float3(0.0f, 5.0f, 0.0f);
const float3 r = (a % b);
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float3 const a = float3(1.0f, 2.0f, 3.0f);
float3 const b = float3(0.0f, 5.0f, 0.0f);
float3 const r = fmod(a, b);
return;
}

View File

@@ -0,0 +1,26 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 15
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%float_0 = OpConstant %float 0
%float_5 = OpConstant %float 5
%13 = OpConstantComposite %v3float %float_0 %float_5 %float_0
%f = OpFunction %void None %1
%4 = OpLabel
%14 = OpFRem %v3float %10 %13
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<f32>(1.0, 2.0, 3.0);
let b = vec3<f32>(0.0, 5.0, 0.0);
let r : vec3<f32> = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = vec3<i32>(0, 5, 0);
let r : vec3<i32> = a % b;
}

View File

@@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const int3 a = int3(1, 2, 3);
const int3 b = int3(0, 5, 0);
const int3 r = (a % int3(1, 5, 1));
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(0, 5, 0);
int3 const r = (a % b);
return;
}

View File

@@ -0,0 +1,26 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 15
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v3int = OpTypeVector %int 3
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%int_0 = OpConstant %int 0
%int_5 = OpConstant %int 5
%13 = OpConstantComposite %v3int %int_0 %int_5 %int_0
%f = OpFunction %void None %1
%4 = OpLabel
%14 = OpSMod %v3int %10 %13
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = vec3<i32>(0, 5, 0);
let r : vec3<i32> = (a % b);
}

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = vec3<u32>(0u, 5u, 0u);
let r : vec3<u32> = a % b;
}

View File

@@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const uint3 a = uint3(1u, 2u, 3u);
const uint3 b = uint3(0u, 5u, 0u);
const uint3 r = (a % uint3(1u, 5u, 1u));
return;
}

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint3 const a = uint3(1u, 2u, 3u);
uint3 const b = uint3(0u, 5u, 0u);
uint3 const r = (a % b);
return;
}

View File

@@ -0,0 +1,26 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 15
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%uint_0 = OpConstant %uint 0
%uint_5 = OpConstant %uint 5
%13 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0
%f = OpFunction %void None %1
%4 = OpLabel
%14 = OpUMod %v3uint %10 %13
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = vec3<u32>(0u, 5u, 0u);
let r : vec3<u32> = (a % b);
}