tint: Add end-to-end tests for expressions using f16 types

This patch add DXC compile flag "-enable-16bit-types" and change profile
to SM6.2 when validating generated HLSL using DXC if f16 extension is
enabled in the WGSL program.
The patch add Tint end-to-end test cases for expressions using f16 type,
including constructor, binary operator, splat, zero-init, and others.
Testcases that use f16 types in uniform or storage buffer are SKIPped,
because such usage is not implemented yet.

Bug: tint:1473, tint:1502
Change-Id: I481ab3d12cbb822f11ef85ba807bca3f9770089b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/96252
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@chromium.org>
Commit-Queue: Zhaoming Jiang <zhaoming.jiang@intel.com>
This commit is contained in:
Zhaoming Jiang
2022-08-03 08:45:25 +00:00
committed by Dawn LUCI CQ
parent 7d04cedce3
commit 7d7976d590
1039 changed files with 14821 additions and 8 deletions

View File

@@ -0,0 +1,7 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
let a = 1.h;
let b = 0.h;
let r : f16 = a % b;
}

View File

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

View File

@@ -0,0 +1,10 @@
SKIP: FAILED
[numthreads(1, 1, 1)]
void f() {
const float16_t r = (float16_t(1.0h) % float16_t(0.0h));
return;
}
FXC validation failure:
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x000002639CD34EB0(3,9-17): error X3000: unrecognized identifier 'float16_t'

View File

@@ -0,0 +1,17 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
float16_t tint_float_modulo(float16_t lhs, float16_t rhs) {
return (lhs - rhs * trunc(lhs / rhs));
}
void f() {
float16_t r = tint_float_modulo(1.0hf, 0.0hf);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

View File

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

View File

@@ -0,0 +1,24 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 9
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%half = OpTypeFloat 16
%half_0x1p_0 = OpConstant %half 0x1p+0
%7 = OpConstantNull %half
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpFRem %half %half_0x1p_0 %7
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,8 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
let a = 1.0h;
let b = 0.0h;
let r : f16 = (a % b);
}

View File

@@ -0,0 +1,7 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
let a = vec3<f16>(1.h, 2.h, 3.h);
let b = vec3<f16>(0.h, 5.h, 0.h);
let r : vec3<f16> = a % b;
}

View File

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

View File

@@ -0,0 +1,14 @@
SKIP: FAILED
[numthreads(1, 1, 1)]
void f() {
const vector<float16_t, 3> a = vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h));
const vector<float16_t, 3> b = vector<float16_t, 3>(float16_t(0.0h), float16_t(5.0h), float16_t(0.0h));
const vector<float16_t, 3> r = (a % b);
return;
}
FXC validation failure:
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000014E5DAA6DB0(3,16-24): error X3000: syntax error: unexpected token 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000014E5DAA6DB0(4,16-24): error X3000: syntax error: unexpected token 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000014E5DAA6DB0(5,16-24): error X3000: syntax error: unexpected token 'float16_t'

View File

@@ -0,0 +1,19 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
f16vec3 tint_float_modulo(f16vec3 lhs, f16vec3 rhs) {
return (lhs - rhs * trunc(lhs / rhs));
}
void f() {
f16vec3 a = f16vec3(1.0hf, 2.0hf, 3.0hf);
f16vec3 b = f16vec3(0.0hf, 5.0hf, 0.0hf);
f16vec3 r = tint_float_modulo(a, b);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

View File

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

View File

@@ -0,0 +1,30 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 15
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
%half_0x1p_0 = OpConstant %half 0x1p+0
%half_0x1p_1 = OpConstant %half 0x1p+1
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%10 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_1 %half_0x1_8p_1
%11 = OpConstantNull %half
%half_0x1_4p_2 = OpConstant %half 0x1.4p+2
%13 = OpConstantComposite %v3half %11 %half_0x1_4p_2 %11
%f = OpFunction %void None %1
%4 = OpLabel
%14 = OpFRem %v3half %10 %13
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,8 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
let a = vec3<f16>(1.0h, 2.0h, 3.0h);
let b = vec3<f16>(0.0h, 5.0h, 0.0h);
let r : vec3<f16> = (a % b);
}

View File

@@ -0,0 +1,7 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = 1.h;
var b = 0.h;
let r : f16 = a % (b + b);
}

View File

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

View File

@@ -0,0 +1,13 @@
SKIP: FAILED
[numthreads(1, 1, 1)]
void f() {
float16_t a = float16_t(1.0h);
float16_t b = float16_t(0.0h);
const float16_t r = (a % (b + b));
return;
}
FXC validation failure:
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x00000290AE802EA0(3,3-11): error X3000: unrecognized identifier 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x00000290AE802EA0(3,13): error X3000: unrecognized identifier 'a'

View File

@@ -0,0 +1,19 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
float16_t tint_float_modulo(float16_t lhs, float16_t rhs) {
return (lhs - rhs * trunc(lhs / rhs));
}
void f() {
float16_t a = 1.0hf;
float16_t b = 0.0hf;
float16_t r = tint_float_modulo(a, (b + b));
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

View File

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

View File

@@ -0,0 +1,35 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 16
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%half = OpTypeFloat 16
%half_0x1p_0 = OpConstant %half 0x1p+0
%_ptr_Function_half = OpTypePointer Function %half
%9 = OpConstantNull %half
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_half Function %9
%b = OpVariable %_ptr_Function_half Function %9
OpStore %a %half_0x1p_0
OpStore %b %9
%11 = OpLoad %half %a
%12 = OpLoad %half %b
%13 = OpLoad %half %b
%14 = OpFAdd %half %12 %13
%15 = OpFRem %half %11 %14
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,8 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = 1.0h;
var b = 0.0h;
let r : f16 = (a % (b + b));
}

View File

@@ -0,0 +1,7 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = vec3<f16>(1.h, 2.h, 3.h);
var b = vec3<f16>(0.h, 5.h, 0.h);
let r : vec3<f16> = a % (b + b);
}

View File

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

View File

@@ -0,0 +1,14 @@
SKIP: FAILED
[numthreads(1, 1, 1)]
void f() {
vector<float16_t, 3> a = vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h));
vector<float16_t, 3> b = vector<float16_t, 3>(float16_t(0.0h), float16_t(5.0h), float16_t(0.0h));
const vector<float16_t, 3> r = (a % (b + b));
return;
}
FXC validation failure:
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x000001C6CF8C3C20(3,10-18): error X3000: syntax error: unexpected token 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x000001C6CF8C3C20(4,10-18): error X3000: syntax error: unexpected token 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x000001C6CF8C3C20(5,16-24): error X3000: syntax error: unexpected token 'float16_t'

View File

@@ -0,0 +1,19 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
f16vec3 tint_float_modulo(f16vec3 lhs, f16vec3 rhs) {
return (lhs - rhs * trunc(lhs / rhs));
}
void f() {
f16vec3 a = f16vec3(1.0hf, 2.0hf, 3.0hf);
f16vec3 b = f16vec3(0.0hf, 5.0hf, 0.0hf);
f16vec3 r = tint_float_modulo(a, (b + b));
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

View File

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

View File

@@ -0,0 +1,42 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 23
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
%half_0x1p_0 = OpConstant %half 0x1p+0
%half_0x1p_1 = OpConstant %half 0x1p+1
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%10 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_1 %half_0x1_8p_1
%_ptr_Function_v3half = OpTypePointer Function %v3half
%13 = OpConstantNull %v3half
%14 = OpConstantNull %half
%half_0x1_4p_2 = OpConstant %half 0x1.4p+2
%16 = OpConstantComposite %v3half %14 %half_0x1_4p_2 %14
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_v3half Function %13
%b = OpVariable %_ptr_Function_v3half Function %13
OpStore %a %10
OpStore %b %16
%18 = OpLoad %v3half %a
%19 = OpLoad %v3half %b
%20 = OpLoad %v3half %b
%21 = OpFAdd %v3half %19 %20
%22 = OpFRem %v3half %18 %21
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,8 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = vec3<f16>(1.0h, 2.0h, 3.0h);
var b = vec3<f16>(0.0h, 5.0h, 0.0h);
let r : vec3<f16> = (a % (b + b));
}

View File

@@ -0,0 +1,7 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = 1.h;
var b = 0.h;
let r : f16 = a % b;
}

View File

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

View File

@@ -0,0 +1,13 @@
SKIP: FAILED
[numthreads(1, 1, 1)]
void f() {
float16_t a = float16_t(1.0h);
float16_t b = float16_t(0.0h);
const float16_t r = (a % b);
return;
}
FXC validation failure:
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000019B7B10C910(3,3-11): error X3000: unrecognized identifier 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000019B7B10C910(3,13): error X3000: unrecognized identifier 'a'

View File

@@ -0,0 +1,19 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
float16_t tint_float_modulo(float16_t lhs, float16_t rhs) {
return (lhs - rhs * trunc(lhs / rhs));
}
void f() {
float16_t a = 1.0hf;
float16_t b = 0.0hf;
float16_t r = tint_float_modulo(a, b);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

View File

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

View File

@@ -0,0 +1,33 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 14
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%half = OpTypeFloat 16
%half_0x1p_0 = OpConstant %half 0x1p+0
%_ptr_Function_half = OpTypePointer Function %half
%9 = OpConstantNull %half
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_half Function %9
%b = OpVariable %_ptr_Function_half Function %9
OpStore %a %half_0x1p_0
OpStore %b %9
%11 = OpLoad %half %a
%12 = OpLoad %half %b
%13 = OpFRem %half %11 %12
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,8 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = 1.0h;
var b = 0.0h;
let r : f16 = (a % b);
}

View File

@@ -0,0 +1,7 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = vec3<f16>(1.h, 2.h, 3.h);
var b = vec3<f16>(0.h, 5.h, 0.h);
let r : vec3<f16> = a % b;
}

View File

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

View File

@@ -0,0 +1,14 @@
SKIP: FAILED
[numthreads(1, 1, 1)]
void f() {
vector<float16_t, 3> a = vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h));
vector<float16_t, 3> b = vector<float16_t, 3>(float16_t(0.0h), float16_t(5.0h), float16_t(0.0h));
const vector<float16_t, 3> r = (a % b);
return;
}
FXC validation failure:
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000024A088362F0(3,10-18): error X3000: syntax error: unexpected token 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000024A088362F0(4,10-18): error X3000: syntax error: unexpected token 'float16_t'
D:\Projects\RampUp\dawn\test\tint\expressions\Shader@0x0000024A088362F0(5,16-24): error X3000: syntax error: unexpected token 'float16_t'

View File

@@ -0,0 +1,19 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
f16vec3 tint_float_modulo(f16vec3 lhs, f16vec3 rhs) {
return (lhs - rhs * trunc(lhs / rhs));
}
void f() {
f16vec3 a = f16vec3(1.0hf, 2.0hf, 3.0hf);
f16vec3 b = f16vec3(0.0hf, 5.0hf, 0.0hf);
f16vec3 r = tint_float_modulo(a, b);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

View File

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

View File

@@ -0,0 +1,40 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 21
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
%half_0x1p_0 = OpConstant %half 0x1p+0
%half_0x1p_1 = OpConstant %half 0x1p+1
%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
%10 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_1 %half_0x1_8p_1
%_ptr_Function_v3half = OpTypePointer Function %v3half
%13 = OpConstantNull %v3half
%14 = OpConstantNull %half
%half_0x1_4p_2 = OpConstant %half 0x1.4p+2
%16 = OpConstantComposite %v3half %14 %half_0x1_4p_2 %14
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_v3half Function %13
%b = OpVariable %_ptr_Function_v3half Function %13
OpStore %a %10
OpStore %b %16
%18 = OpLoad %v3half %a
%19 = OpLoad %v3half %b
%20 = OpFRem %v3half %18 %19
OpReturn
OpFunctionEnd

View File

@@ -0,0 +1,8 @@
enable f16;
@compute @workgroup_size(1)
fn f() {
var a = vec3<f16>(1.0h, 2.0h, 3.0h);
var b = vec3<f16>(0.0h, 5.0h, 0.0h);
let r : vec3<f16> = (a % b);
}