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]);
+}