Additional tests for constant array matrix f16
Change-Id: I04b36c5956a7036e3ce84ca16eb4db5c0683aac6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/244174
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/test/tint/const/array/array_matrix_f16.wgsl b/test/tint/const/array/array_matrix_f16.wgsl
new file mode 100644
index 0000000..552d123
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl
@@ -0,0 +1,14 @@
+// flags: --hlsl-shader-model 62
+enable f16;
+
+@group(0) @binding(0) var<storage, read_write> s: array<u32>;
+
+@compute @workgroup_size(1u)
+fn main(){
+ const kArray = array(
+ mat3x2(vec2(0.0h,1.0h),vec2(2.0h,3.0h),vec2(2.0h,3.0h)),
+ mat3x2(vec2(0.0h,1.0h),vec2(2.0h,3.0h),vec2(2.0h,3.0h)),
+ );
+ var q = 0u;
+ s[0] = u32(kArray[q][0][0]);
+}
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.dxc.hlsl b/test/tint/const/array/array_matrix_f16.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..49a6ceb
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.dxc.hlsl
@@ -0,0 +1,12 @@
+RWByteAddressBuffer s : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_1 = 0u;
+ s.GetDimensions(tint_symbol_1);
+ uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ uint q = 0u;
+ matrix<float16_t, 3, 2> tint_symbol_3[2] = {matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h))), matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)))};
+ s.Store((4u * min(0u, (tint_symbol_2 - 1u))), asuint(uint(tint_symbol_3[min(q, 1u)][0][0])));
+ return;
+}
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.fxc.hlsl b/test/tint/const/array/array_matrix_f16.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..d79d545
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.fxc.hlsl
@@ -0,0 +1,14 @@
+SKIP: INVALID
+
+RWByteAddressBuffer s : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_1 = 0u;
+ s.GetDimensions(tint_symbol_1);
+ uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ uint q = 0u;
+ matrix<float16_t, 3, 2> tint_symbol_3[2] = {matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h))), matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)))};
+ s.Store((4u * min(0u, (tint_symbol_2 - 1u))), asuint(uint(tint_symbol_3[min(q, 1u)][0][0])));
+ return;
+}
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.glsl b/test/tint/const/array/array_matrix_f16.wgsl.expected.glsl
new file mode 100644
index 0000000..27821ea
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.glsl
@@ -0,0 +1,17 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float: require
+
+layout(binding = 0, std430)
+buffer s_block_1_ssbo {
+ uint inner[];
+} v;
+uint tint_f16_to_u32(float16_t value) {
+ return mix(4294967295u, mix(0u, uint(value), (value >= 0.0hf)), (value <= 65504.0hf));
+}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ uint q = 0u;
+ uint v_1 = (uint(v.inner.length()) - 1u);
+ uint v_2 = min(uint(0), v_1);
+ v.inner[v_2] = tint_f16_to_u32(f16mat3x2[2](f16mat3x2(f16vec2(0.0hf, 1.0hf), f16vec2(2.0hf, 3.0hf), f16vec2(2.0hf, 3.0hf)), f16mat3x2(f16vec2(0.0hf, 1.0hf), f16vec2(2.0hf, 3.0hf), f16vec2(2.0hf, 3.0hf)))[min(q, 1u)][0u].x);
+}
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.ir.dxc.hlsl b/test/tint/const/array/array_matrix_f16.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..2dd5805
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,17 @@
+
+RWByteAddressBuffer s : register(u0);
+uint tint_f16_to_u32(float16_t value) {
+ return (((value <= float16_t(65504.0h))) ? ((((value >= float16_t(0.0h))) ? (uint(value)) : (0u))) : (4294967295u));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint q = 0u;
+ uint v = 0u;
+ s.GetDimensions(v);
+ uint v_1 = ((v / 4u) - 1u);
+ uint v_2 = (min(uint(int(0)), v_1) * 4u);
+ matrix<float16_t, 3, 2> v_3[2] = {matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h))), matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)))};
+ s.Store((0u + v_2), tint_f16_to_u32(v_3[min(q, 1u)][0u].x));
+}
+
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.ir.fxc.hlsl b/test/tint/const/array/array_matrix_f16.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7082472
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: INVALID
+
+RWByteAddressBuffer s : register(u0);
+uint tint_f16_to_u32(float16_t value) {
+ return (((value <= float16_t(65504.0h))) ? ((((value >= float16_t(0.0h))) ? (uint(value)) : (0u))) : (4294967295u));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint q = 0u;
+ uint v = 0u;
+ s.GetDimensions(v);
+ uint v_1 = ((v / 4u) - 1u);
+ uint v_2 = (min(uint(int(0)), v_1) * 4u);
+ matrix<float16_t, 3, 2> v_3[2] = {matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h))), matrix<float16_t, 3, 2>(vector<float16_t, 2>(float16_t(0.0h), float16_t(1.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 2>(float16_t(2.0h), float16_t(3.0h)))};
+ s.Store((0u + v_2), tint_f16_to_u32(v_3[min(q, 1u)][0u].x));
+}
+
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.msl b/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
new file mode 100644
index 0000000..a84f77f
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
@@ -0,0 +1,31 @@
+#include <metal_stdlib>
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_module_vars_struct {
+ device tint_array<uint, 1>* s;
+ const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
+};
+
+uint tint_f16_to_u32(half value) {
+ return select(4294967295u, select(0u, uint(value), (value >= 0.0h)), (value <= 65504.0h));
+}
+
+kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
+ uint q = 0u;
+ uint const v_1 = (((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u) - 1u);
+ device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
+ (*v_2) = tint_f16_to_u32(tint_array<half3x2, 2>{half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h)), half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h))}[min(q, 1u)][0u].x);
+}
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.spvasm b/test/tint/const/array/array_matrix_f16.wgsl.expected.spvasm
new file mode 100644
index 0000000..8ec1c7b
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.spvasm
@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 60
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ %37 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpMemberName %s_block_tint_explicit_layout 0 "inner"
+ OpName %s_block_tint_explicit_layout "s_block_tint_explicit_layout"
+ OpName %main "main"
+ OpName %q "q"
+ OpName %tint_f16_to_u32 "tint_f16_to_u32"
+ OpName %value "value"
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpMemberDecorate %s_block_tint_explicit_layout 0 Offset 0
+ OpDecorate %s_block_tint_explicit_layout Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %1 Coherent
+ %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+%s_block_tint_explicit_layout = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer_s_block_tint_explicit_layout = OpTypePointer StorageBuffer %s_block_tint_explicit_layout
+ %1 = OpVariable %_ptr_StorageBuffer_s_block_tint_explicit_layout StorageBuffer
+ %half = OpTypeFloat 16
+ %v2half = OpTypeVector %half 2
+ %mat3v2half = OpTypeMatrix %v2half 3
+ %uint_2 = OpConstant %uint 2
+%_arr_mat3v2half_uint_2 = OpTypeArray %mat3v2half %uint_2
+%_ptr_Private__arr_mat3v2half_uint_2 = OpTypePointer Private %_arr_mat3v2half_uint_2
+%half_0x0p_0 = OpConstant %half 0x0p+0
+%half_0x1p_0 = OpConstant %half 0x1p+0
+ %15 = OpConstantComposite %v2half %half_0x0p_0 %half_0x1p_0
+%half_0x1p_1 = OpConstant %half 0x1p+1
+%half_0x1_8p_1 = OpConstant %half 0x1.8p+1
+ %18 = OpConstantComposite %v2half %half_0x1p_1 %half_0x1_8p_1
+ %14 = OpConstantComposite %mat3v2half %15 %18 %18
+ %13 = OpConstantComposite %_arr_mat3v2half_uint_2 %14 %14
+ %6 = OpVariable %_ptr_Private__arr_mat3v2half_uint_2 Private %13
+ %void = OpTypeVoid
+ %23 = OpTypeFunction %void
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__runtimearr_uint = OpTypePointer StorageBuffer %_runtimearr_uint
+ %uint_1 = OpConstant %uint 1
+ %int = OpTypeInt 32 1
+ %int_0 = OpConstant %int 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_Private_v2half = OpTypePointer Private %v2half
+%_ptr_Private_half = OpTypePointer Private %half
+ %50 = OpTypeFunction %uint %half
+ %bool = OpTypeBool
+%half_0x1_ffcp_15 = OpConstant %half 0x1.ffcp+15
+%uint_4294967295 = OpConstant %uint 4294967295
+ %main = OpFunction %void None %23
+ %24 = OpLabel
+ %q = OpVariable %_ptr_Function_uint Function
+ OpStore %q %uint_0
+ %28 = OpAccessChain %_ptr_StorageBuffer__runtimearr_uint %1 %uint_0
+ %30 = OpArrayLength %uint %1 0
+ %31 = OpISub %uint %30 %uint_1
+ %33 = OpBitcast %uint %int_0
+ %36 = OpExtInst %uint %37 UMin %33 %31
+ %38 = OpAccessChain %_ptr_StorageBuffer_uint %1 %uint_0 %36
+ %40 = OpLoad %uint %q None
+ %41 = OpExtInst %uint %37 UMin %40 %uint_1
+ %42 = OpAccessChain %_ptr_Private_v2half %6 %41 %uint_0
+ %44 = OpAccessChain %_ptr_Private_half %42 %uint_0
+ %46 = OpLoad %half %44 None
+ %47 = OpFunctionCall %uint %tint_f16_to_u32 %46
+ OpStore %38 %47 None
+ OpReturn
+ OpFunctionEnd
+%tint_f16_to_u32 = OpFunction %uint None %50
+ %value = OpFunctionParameter %half
+ %51 = OpLabel
+ %52 = OpConvertFToU %uint %value
+ %53 = OpFOrdGreaterThanEqual %bool %value %half_0x0p_0
+ %55 = OpSelect %uint %53 %52 %uint_0
+ %56 = OpFOrdLessThanEqual %bool %value %half_0x1_ffcp_15
+ %58 = OpSelect %uint %56 %55 %uint_4294967295
+ OpReturnValue %58
+ OpFunctionEnd
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.wgsl b/test/tint/const/array/array_matrix_f16.wgsl.expected.wgsl
new file mode 100644
index 0000000..846076d
--- /dev/null
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+enable f16;
+
+@group(0) @binding(0) var<storage, read_write> s : array<u32>;
+
+@compute @workgroup_size(1u)
+fn main() {
+ const kArray = array(mat3x2(vec2(0.0h, 1.0h), vec2(2.0h, 3.0h), vec2(2.0h, 3.0h)), mat3x2(vec2(0.0h, 1.0h), vec2(2.0h, 3.0h), vec2(2.0h, 3.0h)));
+ var q = 0u;
+ s[0] = u32(kArray[q][0][0]);
+}
