Tint: Implement bitcast for f16
This CL implement the missing bitcast from/to f16 types, including
resolver validation and HLSL/GLSL backends. MSL and SPIRV backends have
native support. Related unittests and tint tests are also implemented.
Bug: tint:1977, dawn:1510, tint:1502
Change-Id: I04499864df4115eeb130aceb359151e8e1ffbfa9
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/140500
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Zhaoming Jiang <zhaoming.jiang@intel.com>
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl
new file mode 100644
index 0000000..b608ac2
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..29c8009
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const float b = a;
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..29c8009
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const float b = a;
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.glsl
new file mode 100644
index 0000000..926bc4a
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ float a = 2.003662109375f;
+ float b = a;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..79da111
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 2.003662109375f;
+ float const b = as_type<float>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..bd59e7b
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.spvasm
@@ -0,0 +1,19 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 8
+; 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
+ %float = OpTypeFloat 32
+%float_2_00366211 = OpConstant %float 2.00366211
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpCopyObject %float %float_2_00366211
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..fdbb7ea
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl
new file mode 100644
index 0000000..d86e3c9
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..af3719c
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const int b = asint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..af3719c
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const int b = asint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.glsl
new file mode 100644
index 0000000..0351d1a
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ float a = 2.003662109375f;
+ int b = floatBitsToInt(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..0ba9a1f
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 2.003662109375f;
+ int const b = as_type<int>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..3cf024e
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.spvasm
@@ -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
+ %float = OpTypeFloat 32
+%float_2_00366211 = OpConstant %float 2.00366211
+ %int = OpTypeInt 32 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %int %float_2_00366211
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..5889390
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl
new file mode 100644
index 0000000..8f9a7d2
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..98f802d
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const uint b = asuint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..98f802d
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const uint b = asuint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.glsl
new file mode 100644
index 0000000..e1fe92a
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ float a = 2.003662109375f;
+ uint b = floatBitsToUint(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..b2ebfde
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 2.003662109375f;
+ uint const b = as_type<uint>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..6d4e8b0
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.spvasm
@@ -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
+ %float = OpTypeFloat 32
+%float_2_00366211 = OpConstant %float 2.00366211
+ %uint = OpTypeInt 32 0
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %uint %float_2_00366211
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..bfd5887
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl
new file mode 100644
index 0000000..ef1e578
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..d859d97
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.dxc.hlsl
@@ -0,0 +1,13 @@
+vector<float16_t, 2> tint_bitcast_to_f16(float src) {
+ uint v = asuint(src);
+ float t_low = f16tof32(v & 0xffff);
+ float t_high = f16tof32((v >> 16) & 0xffff);
+ return vector<float16_t, 2>(t_low.x, t_high.x);
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const vector<float16_t, 2> b = tint_bitcast_to_f16(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..15a981f
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+vector<float16_t, 2> tint_bitcast_to_f16(float src) {
+ uint v = asuint(src);
+ float t_low = f16tof32(v & 0xffff);
+ float t_high = f16tof32((v >> 16) & 0xffff);
+ return vector<float16_t, 2>(t_low.x, t_high.x);
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 2.003662109375f;
+ const vector<float16_t, 2> b = tint_bitcast_to_f16(a);
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x000002C96174CDD0(1,8-16): error X3000: syntax error: unexpected token 'float16_t'
+
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.glsl
new file mode 100644
index 0000000..bb10cb4
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+f16vec2 tint_bitcast_to_f16(float src) {
+ uint r = floatBitsToUint(src);
+ return unpackFloat2x16(r);
+}
+
+void f() {
+ float a = 2.003662109375f;
+ f16vec2 b = tint_bitcast_to_f16(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.msl
new file mode 100644
index 0000000..dee0136
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 2.003662109375f;
+ half2 const b = as_type<half2>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.spvasm
new file mode 100644
index 0000000..0120613
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; 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
+ %float = OpTypeFloat 32
+%float_2_00366211 = OpConstant %float 2.00366211
+ %half = OpTypeFloat 16
+ %v2half = OpTypeVector %half 2
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %v2half %float_2_00366211
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.wgsl
new file mode 100644
index 0000000..4a8265f
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/f32-vec2f16.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : f32 = f32(2.003662109375f);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl
new file mode 100644
index 0000000..ccfd3e6
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..123510b
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const float b = asfloat(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..123510b
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const float b = asfloat(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.glsl
new file mode 100644
index 0000000..27d7c79
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ int a = 1073757184;
+ float b = intBitsToFloat(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..93e6613
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 1073757184;
+ float const b = as_type<float>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..33772ce
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.spvasm
@@ -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
+ %int = OpTypeInt 32 1
+%int_1073757184 = OpConstant %int 1073757184
+ %float = OpTypeFloat 32
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %float %int_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..429b08e3
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl
new file mode 100644
index 0000000..96a6e7d
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..6cf8029
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const int b = a;
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..6cf8029
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const int b = a;
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.glsl
new file mode 100644
index 0000000..8569758
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ int a = 1073757184;
+ int b = a;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..0ad16d6
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 1073757184;
+ int const b = as_type<int>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..c40927f
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.spvasm
@@ -0,0 +1,19 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 8
+; 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_1073757184 = OpConstant %int 1073757184
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpCopyObject %int %int_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..2181d6c
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl
new file mode 100644
index 0000000..cee6401
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b1c1381
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const uint b = asuint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b1c1381
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const uint b = asuint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.glsl
new file mode 100644
index 0000000..e7bfa52
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ int a = 1073757184;
+ uint b = uint(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..58d38f1
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 1073757184;
+ uint const b = as_type<uint>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..7976301
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.spvasm
@@ -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
+ %int = OpTypeInt 32 1
+%int_1073757184 = OpConstant %int 1073757184
+ %uint = OpTypeInt 32 0
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %uint %int_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..47781bf
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl
new file mode 100644
index 0000000..0637d05
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..1b76cc5
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.dxc.hlsl
@@ -0,0 +1,13 @@
+vector<float16_t, 2> tint_bitcast_to_f16(int src) {
+ uint v = asuint(src);
+ float t_low = f16tof32(v & 0xffff);
+ float t_high = f16tof32((v >> 16) & 0xffff);
+ return vector<float16_t, 2>(t_low.x, t_high.x);
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const vector<float16_t, 2> b = tint_bitcast_to_f16(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7a19f9d
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+vector<float16_t, 2> tint_bitcast_to_f16(int src) {
+ uint v = asuint(src);
+ float t_low = f16tof32(v & 0xffff);
+ float t_high = f16tof32((v >> 16) & 0xffff);
+ return vector<float16_t, 2>(t_low.x, t_high.x);
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 1073757184;
+ const vector<float16_t, 2> b = tint_bitcast_to_f16(a);
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x00000234EF107030(1,8-16): error X3000: syntax error: unexpected token 'float16_t'
+
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.glsl
new file mode 100644
index 0000000..138bdb8
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+f16vec2 tint_bitcast_to_f16(int src) {
+ uint r = uint(src);
+ return unpackFloat2x16(r);
+}
+
+void f() {
+ int a = 1073757184;
+ f16vec2 b = tint_bitcast_to_f16(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.msl
new file mode 100644
index 0000000..193b2f9
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 1073757184;
+ half2 const b = as_type<half2>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.spvasm
new file mode 100644
index 0000000..2d19adb
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; 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
+ %int = OpTypeInt 32 1
+%int_1073757184 = OpConstant %int 1073757184
+ %half = OpTypeFloat 16
+ %v2half = OpTypeVector %half 2
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %v2half %int_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.wgsl
new file mode 100644
index 0000000..f2a8cdf
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/i32-vec2f16.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : i32 = i32(1073757184i);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl
new file mode 100644
index 0000000..6a166e7
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..3054f48
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const float b = asfloat(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..3054f48
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const float b = asfloat(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.glsl
new file mode 100644
index 0000000..957fb56
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ uint a = 1073757184u;
+ float b = uintBitsToFloat(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..139ad95
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 1073757184u;
+ float const b = as_type<float>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..88e7f95
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.spvasm
@@ -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_1073757184 = OpConstant %uint 1073757184
+ %float = OpTypeFloat 32
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %float %uint_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..3340b88
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl
new file mode 100644
index 0000000..3571b0b
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..cf37bd5
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const int b = asint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..cf37bd5
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const int b = asint(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.glsl
new file mode 100644
index 0000000..9846453
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ uint a = 1073757184u;
+ int b = int(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..d91edb8
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 1073757184u;
+ int const b = as_type<int>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..9a1cf21
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.spvasm
@@ -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_1073757184 = OpConstant %uint 1073757184
+ %int = OpTypeInt 32 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %int %uint_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..1de1638
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl
new file mode 100644
index 0000000..b5b8e8f
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..6309306
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const uint b = a;
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..6309306
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const uint b = a;
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.glsl
new file mode 100644
index 0000000..e189e0b
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ uint a = 1073757184u;
+ uint b = a;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..99b84ac
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 1073757184u;
+ uint const b = as_type<uint>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..e705859
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.spvasm
@@ -0,0 +1,19 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 8
+; 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_1073757184 = OpConstant %uint 1073757184
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpCopyObject %uint %uint_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..dda8627
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl
new file mode 100644
index 0000000..f0e0801
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..2e7bcd9
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.dxc.hlsl
@@ -0,0 +1,13 @@
+vector<float16_t, 2> tint_bitcast_to_f16(uint src) {
+ uint v = asuint(src);
+ float t_low = f16tof32(v & 0xffff);
+ float t_high = f16tof32((v >> 16) & 0xffff);
+ return vector<float16_t, 2>(t_low.x, t_high.x);
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const vector<float16_t, 2> b = tint_bitcast_to_f16(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..9099fe9
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+vector<float16_t, 2> tint_bitcast_to_f16(uint src) {
+ uint v = asuint(src);
+ float t_low = f16tof32(v & 0xffff);
+ float t_high = f16tof32((v >> 16) & 0xffff);
+ return vector<float16_t, 2>(t_low.x, t_high.x);
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 1073757184u;
+ const vector<float16_t, 2> b = tint_bitcast_to_f16(a);
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x0000020E6D499E20(1,8-16): error X3000: syntax error: unexpected token 'float16_t'
+
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.glsl
new file mode 100644
index 0000000..e19bd66
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+f16vec2 tint_bitcast_to_f16(uint src) {
+ uint r = uint(src);
+ return unpackFloat2x16(r);
+}
+
+void f() {
+ uint a = 1073757184u;
+ f16vec2 b = tint_bitcast_to_f16(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.msl
new file mode 100644
index 0000000..c8d0f0a
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 1073757184u;
+ half2 const b = as_type<half2>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.spvasm
new file mode 100644
index 0000000..a741773
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; 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
+ %uint = OpTypeInt 32 0
+%uint_1073757184 = OpConstant %uint 1073757184
+ %half = OpTypeFloat 16
+ %v2half = OpTypeVector %half 2
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %v2half %uint_1073757184
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.wgsl
new file mode 100644
index 0000000..3daf0a8
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/u32-vec2f16.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : u32 = u32(1073757184u);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl
new file mode 100644
index 0000000..151c425
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..480ada0
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,11 @@
+float tint_bitcast_from_f16(vector<float16_t, 2> src) {
+ uint2 r = f32tof16(float2(src));
+ return asfloat(uint((r.x & 0xffff) | ((r.y & 0xffff) << 16)));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const float b = tint_bitcast_from_f16(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..71f7b7a
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+float tint_bitcast_from_f16(vector<float16_t, 2> src) {
+ uint2 r = f32tof16(float2(src));
+ return asfloat(uint((r.x & 0xffff) | ((r.y & 0xffff) << 16)));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const float b = tint_bitcast_from_f16(a);
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x000001EDAADDFBA0(1,36-44): error X3000: syntax error: unexpected token 'float16_t'
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x000001EDAADDFBA0(2,29-31): error X3004: undeclared identifier 'src'
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x000001EDAADDFBA0(2,22-32): error X3014: incorrect number of arguments to numeric-type constructor
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.glsl
new file mode 100644
index 0000000..d4dc611
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+float tint_bitcast_from_f16(f16vec2 src) {
+ uint r = packFloat2x16(src);
+ return uintBitsToFloat(r);
+}
+
+void f() {
+ f16vec2 a = f16vec2(1.0hf, 2.0hf);
+ float b = tint_bitcast_from_f16(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.msl
new file mode 100644
index 0000000..34c54d2
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ half2 const a = half2(1.0h, 2.0h);
+ float const b = as_type<float>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..f196339
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.spvasm
@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; 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
+ %v2half = OpTypeVector %half 2
+%half_0x1p_0 = OpConstant %half 0x1p+0
+%half_0x1p_1 = OpConstant %half 0x1p+1
+ %9 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_1
+ %float = OpTypeFloat 32
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %10 = OpBitcast %float %9
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..01e9331
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-f32.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl
new file mode 100644
index 0000000..247d051
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..8ad784d
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,11 @@
+int tint_bitcast_from_f16(vector<float16_t, 2> src) {
+ uint2 r = f32tof16(float2(src));
+ return asint(uint((r.x & 0xffff) | ((r.y & 0xffff) << 16)));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const int b = tint_bitcast_from_f16(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..ef2dd6d
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+int tint_bitcast_from_f16(vector<float16_t, 2> src) {
+ uint2 r = f32tof16(float2(src));
+ return asint(uint((r.x & 0xffff) | ((r.y & 0xffff) << 16)));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const int b = tint_bitcast_from_f16(a);
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x000001DC9BD890B0(1,34-42): error X3000: syntax error: unexpected token 'float16_t'
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x000001DC9BD890B0(2,29-31): error X3004: undeclared identifier 'src'
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x000001DC9BD890B0(2,22-32): error X3014: incorrect number of arguments to numeric-type constructor
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.glsl
new file mode 100644
index 0000000..26dcde7
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+int tint_bitcast_from_f16(f16vec2 src) {
+ uint r = packFloat2x16(src);
+ return int(r);
+}
+
+void f() {
+ f16vec2 a = f16vec2(1.0hf, 2.0hf);
+ int b = tint_bitcast_from_f16(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.msl
new file mode 100644
index 0000000..a0b1f99
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ half2 const a = half2(1.0h, 2.0h);
+ int const b = as_type<int>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..4f4902c
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.spvasm
@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; 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
+ %v2half = OpTypeVector %half 2
+%half_0x1p_0 = OpConstant %half 0x1p+0
+%half_0x1p_1 = OpConstant %half 0x1p+1
+ %9 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_1
+ %int = OpTypeInt 32 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %10 = OpBitcast %int %9
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..364c99a
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-i32.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl
new file mode 100644
index 0000000..eb74f13
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..48e040e
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.dxc.hlsl
@@ -0,0 +1,11 @@
+uint tint_bitcast_from_f16(vector<float16_t, 2> src) {
+ uint2 r = f32tof16(float2(src));
+ return asuint(uint((r.x & 0xffff) | ((r.y & 0xffff) << 16)));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const uint b = tint_bitcast_from_f16(a);
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..2332ca6
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+uint tint_bitcast_from_f16(vector<float16_t, 2> src) {
+ uint2 r = f32tof16(float2(src));
+ return asuint(uint((r.x & 0xffff) | ((r.y & 0xffff) << 16)));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const uint b = tint_bitcast_from_f16(a);
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x0000011D0C566F70(1,35-43): error X3000: syntax error: unexpected token 'float16_t'
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x0000011D0C566F70(2,29-31): error X3004: undeclared identifier 'src'
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x0000011D0C566F70(2,22-32): error X3014: incorrect number of arguments to numeric-type constructor
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.glsl
new file mode 100644
index 0000000..233cc25
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+uint tint_bitcast_from_f16(f16vec2 src) {
+ uint r = packFloat2x16(src);
+ return uint(r);
+}
+
+void f() {
+ f16vec2 a = f16vec2(1.0hf, 2.0hf);
+ uint b = tint_bitcast_from_f16(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.msl
new file mode 100644
index 0000000..36a2afe
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ half2 const a = half2(1.0h, 2.0h);
+ uint const b = as_type<uint>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..bd1959a
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.spvasm
@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; 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
+ %v2half = OpTypeVector %half 2
+%half_0x1p_0 = OpConstant %half 0x1p+0
+%half_0x1p_1 = OpConstant %half 0x1p+1
+ %9 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_1
+ %uint = OpTypeInt 32 0
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %10 = OpBitcast %uint %9
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..04b2414
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-u32.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl
new file mode 100644
index 0000000..48294a6
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.dxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..eef9d47
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.dxc.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const vector<float16_t, 2> b = a;
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.fxc.hlsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..efecfa7
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.fxc.hlsl
@@ -0,0 +1,12 @@
+SKIP: FAILED
+
+[numthreads(1, 1, 1)]
+void f() {
+ const vector<float16_t, 2> a = vector<float16_t, 2>(float16_t(1.0h), float16_t(2.0h));
+ const vector<float16_t, 2> b = a;
+ return;
+}
+FXC validation failure:
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x0000019B1F6BF7B0(3,16-24): error X3000: syntax error: unexpected token 'float16_t'
+D:\Projects\RampUp\dawn\test\tint\expressions\bitcast\Shader@0x0000019B1F6BF7B0(4,16-24): error X3000: syntax error: unexpected token 'float16_t'
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.glsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.glsl
new file mode 100644
index 0000000..b3f50d0
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.glsl
@@ -0,0 +1,13 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void f() {
+ f16vec2 a = f16vec2(1.0hf, 2.0hf);
+ f16vec2 b = a;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.msl b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.msl
new file mode 100644
index 0000000..8a9cbd9
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ half2 const a = half2(1.0h, 2.0h);
+ half2 const b = as_type<half2>(a);
+ return;
+}
+
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.spvasm b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.spvasm
new file mode 100644
index 0000000..29f12c6
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.spvasm
@@ -0,0 +1,26 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 11
+; 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
+ %v2half = OpTypeVector %half 2
+%half_0x1p_0 = OpConstant %half 0x1p+0
+%half_0x1p_1 = OpConstant %half 0x1p+1
+ %9 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %10 = OpCopyObject %v2half %9
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.wgsl b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.wgsl
new file mode 100644
index 0000000..bebc4e7
--- /dev/null
+++ b/test/tint/expressions/bitcast/let/32bit/vec2f16-vec2f16.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+enable f16;
+
+@compute @workgroup_size(1)
+fn f() {
+ let a : vec2<f16> = vec2<f16>(1.0h, 2.0h);
+ let b : vec2<f16> = bitcast<vec2<f16>>(a);
+}
