tint: End-to-end tests for f16 built-in
This patch add Tint end-to-end tests for built-ins using f16 types.
Bug: tint:1473, tint:1502
Change-Id: I09db6e0b7e90541fb246e11d475e27be96a6d07e
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/97340
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Zhaoming Jiang <zhaoming.jiang@intel.com>
diff --git a/test/tint/builtins/gen/literal/log2/38b478.wgsl b/test/tint/builtins/gen/literal/log2/38b478.wgsl
new file mode 100644
index 0000000..493fbb2
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/38b478.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn log2(vec<3, f16>) -> vec<3, f16>
+fn log2_38b478() {
+ var res: vec3<f16> = log2(vec3<f16>(f16()));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_38b478();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_38b478();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_38b478();
+}
diff --git a/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..029b567
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void log2_38b478() {
+ vector<float16_t, 3> res = log2((float16_t(0.0h)).xxx);
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_38b478();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_38b478();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_38b478();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..76ffd38
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.fxc.hlsl
@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+void log2_38b478() {
+ vector<float16_t, 3> res = log2((float16_t(0.0h)).xxx);
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_38b478();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_38b478();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_38b478();
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\builtins\Shader@0x000001D9884A1160(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+
diff --git a/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.glsl b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.glsl
new file mode 100644
index 0000000..e0f07c6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.glsl
@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_38b478() {
+ f16vec3 res = log2(f16vec3(0.0hf));
+}
+
+vec4 vertex_main() {
+ log2_38b478();
+ return vec4(0.0f);
+}
+
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void log2_38b478() {
+ f16vec3 res = log2(f16vec3(0.0hf));
+}
+
+void fragment_main() {
+ log2_38b478();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_38b478() {
+ f16vec3 res = log2(f16vec3(0.0hf));
+}
+
+void compute_main() {
+ log2_38b478();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.msl b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.msl
new file mode 100644
index 0000000..a36b81f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void log2_38b478() {
+ half3 res = log2(half3(0.0h));
+}
+
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+ log2_38b478();
+ return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ log2_38b478();
+ return;
+}
+
+kernel void compute_main() {
+ log2_38b478();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.spvasm
new file mode 100644
index 0000000..de97fed
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.spvasm
@@ -0,0 +1,72 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ %16 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %log2_38b478 "log2_38b478"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %half = OpTypeFloat 16
+ %v3half = OpTypeVector %half 3
+ %17 = OpConstantNull %v3half
+%_ptr_Function_v3half = OpTypePointer Function %v3half
+ %20 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+%log2_38b478 = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v3half Function %17
+ %13 = OpExtInst %v3half %16 Log2 %17
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+ %22 = OpLabel
+ %23 = OpFunctionCall %void %log2_38b478
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %25 = OpLabel
+ %26 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %26
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %29 = OpLabel
+ %30 = OpFunctionCall %void %log2_38b478
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %log2_38b478
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.wgsl
new file mode 100644
index 0000000..9dee10e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/38b478.wgsl.expected.wgsl
@@ -0,0 +1,21 @@
+enable f16;
+
+fn log2_38b478() {
+ var res : vec3<f16> = log2(vec3<f16>(f16()));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_38b478();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_38b478();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_38b478();
+}
diff --git a/test/tint/builtins/gen/literal/log2/776088.wgsl b/test/tint/builtins/gen/literal/log2/776088.wgsl
new file mode 100644
index 0000000..3843754
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/776088.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn log2(vec<4, f16>) -> vec<4, f16>
+fn log2_776088() {
+ var res: vec4<f16> = log2(vec4<f16>(f16()));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_776088();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_776088();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_776088();
+}
diff --git a/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..2a3af8f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void log2_776088() {
+ vector<float16_t, 4> res = log2((float16_t(0.0h)).xxxx);
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_776088();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_776088();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_776088();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..cc463e0
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.fxc.hlsl
@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+void log2_776088() {
+ vector<float16_t, 4> res = log2((float16_t(0.0h)).xxxx);
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_776088();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_776088();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_776088();
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\builtins\Shader@0x00000254F7CB1160(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+
diff --git a/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.glsl b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.glsl
new file mode 100644
index 0000000..ebf4aa1
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.glsl
@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_776088() {
+ f16vec4 res = log2(f16vec4(0.0hf));
+}
+
+vec4 vertex_main() {
+ log2_776088();
+ return vec4(0.0f);
+}
+
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void log2_776088() {
+ f16vec4 res = log2(f16vec4(0.0hf));
+}
+
+void fragment_main() {
+ log2_776088();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_776088() {
+ f16vec4 res = log2(f16vec4(0.0hf));
+}
+
+void compute_main() {
+ log2_776088();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.msl b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.msl
new file mode 100644
index 0000000..eae4b4d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void log2_776088() {
+ half4 res = log2(half4(0.0h));
+}
+
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+ log2_776088();
+ return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ log2_776088();
+ return;
+}
+
+kernel void compute_main() {
+ log2_776088();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.spvasm
new file mode 100644
index 0000000..47d547b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.spvasm
@@ -0,0 +1,72 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ %16 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %log2_776088 "log2_776088"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %half = OpTypeFloat 16
+ %v4half = OpTypeVector %half 4
+ %17 = OpConstantNull %v4half
+%_ptr_Function_v4half = OpTypePointer Function %v4half
+ %20 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+%log2_776088 = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v4half Function %17
+ %13 = OpExtInst %v4half %16 Log2 %17
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+ %22 = OpLabel
+ %23 = OpFunctionCall %void %log2_776088
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %25 = OpLabel
+ %26 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %26
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %29 = OpLabel
+ %30 = OpFunctionCall %void %log2_776088
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %log2_776088
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.wgsl
new file mode 100644
index 0000000..87971b4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/776088.wgsl.expected.wgsl
@@ -0,0 +1,21 @@
+enable f16;
+
+fn log2_776088() {
+ var res : vec4<f16> = log2(vec4<f16>(f16()));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_776088();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_776088();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_776088();
+}
diff --git a/test/tint/builtins/gen/literal/log2/8c10b3.wgsl b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl
new file mode 100644
index 0000000..669e8ca
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn log2(f16) -> f16
+fn log2_8c10b3() {
+ var res: f16 = log2(f16());
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_8c10b3();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_8c10b3();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_8c10b3();
+}
diff --git a/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..9edb6af
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void log2_8c10b3() {
+ float16_t res = log2(float16_t(0.0h));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_8c10b3();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_8c10b3();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_8c10b3();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..a1f50c0
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.fxc.hlsl
@@ -0,0 +1,36 @@
+SKIP: FAILED
+
+void log2_8c10b3() {
+ float16_t res = log2(float16_t(0.0h));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_8c10b3();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_8c10b3();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_8c10b3();
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\builtins\Shader@0x0000021DFAAFB990(2,3-11): error X3000: unrecognized identifier 'float16_t'
+D:\Projects\RampUp\dawn\test\tint\builtins\Shader@0x0000021DFAAFB990(2,13-15): error X3000: unrecognized identifier 'res'
+
diff --git a/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.glsl b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.glsl
new file mode 100644
index 0000000..7b9deb9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.glsl
@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_8c10b3() {
+ float16_t res = log2(0.0hf);
+}
+
+vec4 vertex_main() {
+ log2_8c10b3();
+ return vec4(0.0f);
+}
+
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void log2_8c10b3() {
+ float16_t res = log2(0.0hf);
+}
+
+void fragment_main() {
+ log2_8c10b3();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_8c10b3() {
+ float16_t res = log2(0.0hf);
+}
+
+void compute_main() {
+ log2_8c10b3();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.msl b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.msl
new file mode 100644
index 0000000..0ea021a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void log2_8c10b3() {
+ half res = log2(0.0h);
+}
+
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+ log2_8c10b3();
+ return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ log2_8c10b3();
+ return;
+}
+
+kernel void compute_main() {
+ log2_8c10b3();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.spvasm
new file mode 100644
index 0000000..f269f68
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.spvasm
@@ -0,0 +1,71 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ %15 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %log2_8c10b3 "log2_8c10b3"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %half = OpTypeFloat 16
+ %16 = OpConstantNull %half
+%_ptr_Function_half = OpTypePointer Function %half
+ %19 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+%log2_8c10b3 = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_half Function %16
+ %13 = OpExtInst %half %15 Log2 %16
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %19
+ %21 = OpLabel
+ %22 = OpFunctionCall %void %log2_8c10b3
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %24 = OpLabel
+ %25 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %25
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %28 = OpLabel
+ %29 = OpFunctionCall %void %log2_8c10b3
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %log2_8c10b3
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.wgsl
new file mode 100644
index 0000000..6de59ec
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/8c10b3.wgsl.expected.wgsl
@@ -0,0 +1,21 @@
+enable f16;
+
+fn log2_8c10b3() {
+ var res : f16 = log2(f16());
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_8c10b3();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_8c10b3();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_8c10b3();
+}
diff --git a/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl
new file mode 100644
index 0000000..9f7c93a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn log2(vec<2, f16>) -> vec<2, f16>
+fn log2_fb9f0b() {
+ var res: vec2<f16> = log2(vec2<f16>(f16()));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_fb9f0b();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_fb9f0b();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_fb9f0b();
+}
diff --git a/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..1614719
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void log2_fb9f0b() {
+ vector<float16_t, 2> res = log2((float16_t(0.0h)).xx);
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_fb9f0b();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_fb9f0b();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_fb9f0b();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..13779ea
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.fxc.hlsl
@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+void log2_fb9f0b() {
+ vector<float16_t, 2> res = log2((float16_t(0.0h)).xx);
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ log2_fb9f0b();
+ return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ log2_fb9f0b();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ log2_fb9f0b();
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\builtins\Shader@0x000001D63F9088D0(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+
diff --git a/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.glsl b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.glsl
new file mode 100644
index 0000000..73abe1d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.glsl
@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_fb9f0b() {
+ f16vec2 res = log2(f16vec2(0.0hf));
+}
+
+vec4 vertex_main() {
+ log2_fb9f0b();
+ return vec4(0.0f);
+}
+
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void log2_fb9f0b() {
+ f16vec2 res = log2(f16vec2(0.0hf));
+}
+
+void fragment_main() {
+ log2_fb9f0b();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void log2_fb9f0b() {
+ f16vec2 res = log2(f16vec2(0.0hf));
+}
+
+void compute_main() {
+ log2_fb9f0b();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.msl b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.msl
new file mode 100644
index 0000000..0532f4d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void log2_fb9f0b() {
+ half2 res = log2(half2(0.0h));
+}
+
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+ log2_fb9f0b();
+ return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ log2_fb9f0b();
+ return;
+}
+
+kernel void compute_main() {
+ log2_fb9f0b();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.spvasm
new file mode 100644
index 0000000..2dfd8bd
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.spvasm
@@ -0,0 +1,72 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+ OpCapability Shader
+ OpCapability Float16
+ OpCapability UniformAndStorageBuffer16BitAccess
+ OpCapability StorageBuffer16BitAccess
+ OpCapability StorageInputOutput16
+ %16 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %log2_fb9f0b "log2_fb9f0b"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %half = OpTypeFloat 16
+ %v2half = OpTypeVector %half 2
+ %17 = OpConstantNull %v2half
+%_ptr_Function_v2half = OpTypePointer Function %v2half
+ %20 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+%log2_fb9f0b = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v2half Function %17
+ %13 = OpExtInst %v2half %16 Log2 %17
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+ %22 = OpLabel
+ %23 = OpFunctionCall %void %log2_fb9f0b
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %25 = OpLabel
+ %26 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %26
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %29 = OpLabel
+ %30 = OpFunctionCall %void %log2_fb9f0b
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %log2_fb9f0b
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.wgsl
new file mode 100644
index 0000000..4fdced8
--- /dev/null
+++ b/test/tint/builtins/gen/literal/log2/fb9f0b.wgsl.expected.wgsl
@@ -0,0 +1,21 @@
+enable f16;
+
+fn log2_fb9f0b() {
+ var res : vec2<f16> = log2(vec2<f16>(f16()));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ log2_fb9f0b();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ log2_fb9f0b();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ log2_fb9f0b();
+}
