MSL writer: make signed int overflow defined behaviour

Bug: tint:124
Change-Id: Icf545b633d6390ceb7f639e80111390005e311a1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60100
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
Antonio Maiorano
2021-07-29 13:51:47 +00:00
committed by Tint LUCI CQ
parent 1f0200a3ff
commit e5dbe24e94
651 changed files with 3963 additions and 3048 deletions

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int const a = 1;
int const b = 2;
int const r = (a + b);
int const r = as_type<int>((as_type<uint>(a) + as_type<uint>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
int3 const r = (a + b);
int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
int3 const r = (a + b);
int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
int3 const r = (a + b);
int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint3>(b)));
return;
}

View File

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

View File

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

View File

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

View File

@@ -0,0 +1,21 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 10
; 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
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
%f = OpFunction %void None %1
%4 = OpLabel
%9 = OpShiftLeftLogical %int %int_1 %uint_2
OpReturn
OpFunctionEnd

View File

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

View File

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

View File

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

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint const a = 1u;
uint const b = 2u;
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_2 = OpConstant %uint 2
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpShiftLeftLogical %uint %uint_1 %uint_2
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1u;
let b = 2u;
let r : 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 = vec3<u32>(4u, 5u, 6u);
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 uint3 b = uint3(4u, 5u, 6u);
const int3 r = (a << b);
return;
}

View File

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

View File

@@ -0,0 +1,29 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 18
; 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
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_4 = OpConstant %uint 4
%uint_5 = OpConstant %uint 5
%uint_6 = OpConstant %uint 6
%16 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
%f = OpFunction %void None %1
%4 = OpLabel
%17 = OpShiftLeftLogical %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 = vec3<u32>(4u, 5u, 6u);
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>(4u, 5u, 6u);
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(4u, 5u, 6u);
const uint3 r = (a << b);
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(4u, 5u, 6u);
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
%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_4 = OpConstant %uint 4
%uint_5 = OpConstant %uint 5
%uint_6 = OpConstant %uint 6
%14 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
%f = OpFunction %void None %1
%4 = OpLabel
%15 = OpShiftLeftLogical %v3uint %10 %14
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>(4u, 5u, 6u);
let r : vec3<u32> = (a << b);
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int const a = 1;
int const b = 2;
int const r = (a * b);
int const r = as_type<int>((as_type<uint>(a) * as_type<uint>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
int3 const r = (a * b);
int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
int3 const r = (a * b);
int3 const r = as_type<int3>((as_type<uint3>(a) * as_type<uint>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
int3 const r = (a * b);
int3 const r = as_type<int3>((as_type<uint3>(a) * as_type<uint3>(b)));
return;
}

View File

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

View File

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

View File

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

View File

@@ -0,0 +1,21 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 10
; 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
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
%f = OpFunction %void None %1
%4 = OpLabel
%9 = OpShiftRightArithmetic %int %int_1 %uint_2
OpReturn
OpFunctionEnd

View File

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

View File

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

View File

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

View File

@@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint const a = 1u;
uint const b = 2u;
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_2 = OpConstant %uint 2
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpShiftRightLogical %uint %uint_1 %uint_2
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1u;
let b = 2u;
let r : 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 = vec3<u32>(4u, 5u, 6u);
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 uint3 b = uint3(4u, 5u, 6u);
const int3 r = (a >> b);
return;
}

View File

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

View File

@@ -0,0 +1,29 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 18
; 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
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_4 = OpConstant %uint 4
%uint_5 = OpConstant %uint 5
%uint_6 = OpConstant %uint 6
%16 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
%f = OpFunction %void None %1
%4 = OpLabel
%17 = OpShiftRightArithmetic %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 = vec3<u32>(4u, 5u, 6u);
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>(4u, 5u, 6u);
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(4u, 5u, 6u);
const uint3 r = (a >> b);
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(4u, 5u, 6u);
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
%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_4 = OpConstant %uint 4
%uint_5 = OpConstant %uint 5
%uint_6 = OpConstant %uint 6
%14 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
%f = OpFunction %void None %1
%4 = OpLabel
%15 = OpShiftRightLogical %v3uint %10 %14
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>(4u, 5u, 6u);
let r : vec3<u32> = (a >> b);
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int const a = 1;
int const b = 2;
int const r = (a - b);
int const r = as_type<int>((as_type<uint>(a) - as_type<uint>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
int3 const r = (a - b);
int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
int3 const r = (a - b);
int3 const r = as_type<int3>((as_type<uint3>(a) - as_type<uint>(b)));
return;
}

View File

@@ -4,7 +4,7 @@ using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
int3 const r = (a - b);
int3 const r = as_type<int3>((as_type<uint3>(a) - as_type<uint3>(b)));
return;
}

View File

@@ -2,8 +2,8 @@
using namespace metal;
void f() {
int2 v2 = int2((1 + 2));
int3 v3 = int3((1 + 2));
int4 v4 = int4((1 + 2));
int2 v2 = int2(as_type<int>((as_type<uint>(1) + as_type<uint>(2))));
int3 v3 = int3(as_type<int>((as_type<uint>(1) + as_type<uint>(2))));
int4 v4 = int4(as_type<int>((as_type<uint>(1) + as_type<uint>(2))));
}

View File

@@ -2,7 +2,7 @@
using namespace metal;
void f() {
int v = (1 + 2);
int v = as_type<int>((as_type<uint>(1) + as_type<uint>(2)));
int2 v2 = int2(v);
int3 v3 = int3(v);
int4 v4 = int4(v);

View File

@@ -0,0 +1,7 @@
fn i(x : i32) -> i32 {
return -x;
}
fn vi(x : vec4<i32>) -> vec4<i32> {
return -x;
}

View File

@@ -0,0 +1,12 @@
[numthreads(1, 1, 1)]
void unused_entry_point() {
return;
}
int i(int x) {
return -(x);
}
int4 vi(int4 x) {
return -(x);
}

View File

@@ -0,0 +1,20 @@
#include <metal_stdlib>
using namespace metal;
int tint_unary_minus(const int v) {
return select(-v, v, v == -2147483648);
}
int4 tint_unary_minus_1(const int4 v) {
return select(-v, v, v == -2147483648);
}
int i(int x) {
return tint_unary_minus(x);
}
int4 vi(int4 x) {
return tint_unary_minus_1(x);
}

View File

@@ -0,0 +1,36 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
OpExecutionMode %unused_entry_point LocalSize 1 1 1
OpName %unused_entry_point "unused_entry_point"
OpName %i "i"
OpName %x "x"
OpName %vi "vi"
OpName %x_0 "x"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%5 = OpTypeFunction %int %int
%v4int = OpTypeVector %int 4
%11 = OpTypeFunction %v4int %v4int
%unused_entry_point = OpFunction %void None %1
%4 = OpLabel
OpReturn
OpFunctionEnd
%i = OpFunction %int None %5
%x = OpFunctionParameter %int
%9 = OpLabel
%10 = OpSNegate %int %x
OpReturnValue %10
OpFunctionEnd
%vi = OpFunction %v4int None %11
%x_0 = OpFunctionParameter %v4int
%15 = OpLabel
%16 = OpSNegate %v4int %x_0
OpReturnValue %16
OpFunctionEnd

View File

@@ -0,0 +1,7 @@
fn i(x : i32) -> i32 {
return -(x);
}
fn vi(x : vec4<i32>) -> vec4<i32> {
return -(x);
}