writer/hlsl: Support bitcast of vectors
Add end to end tests
Fixed: tint:1026
Change-Id: I10813cbe6dc4f1bccddf9a8a29e3a249a364c051
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59663
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index a85a555..a924c86 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -268,6 +268,10 @@
bool GeneratorImpl::EmitBitcast(std::ostream& out,
ast::BitcastExpression* expr) {
auto* type = TypeOf(expr);
+ if (auto* vec = type->UnwrapRef()->As<sem::Vector>()) {
+ type = vec->type();
+ }
+
if (!type->is_integer_scalar() && !type->is_float_scalar()) {
diagnostics_.add_error(diag::System::Writer,
"Unable to do bitcast to type " + type->type_name());
diff --git a/test/expressions/bitcast/scalar/f32-f32.wgsl b/test/expressions/bitcast/scalar/f32-f32.wgsl
new file mode 100644
index 0000000..a91bec0
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-f32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : f32 = 1.;
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2e5fee6
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float b = asfloat(1.0f);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.msl b/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..ea55957
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 1.0f;
+ float const b = as_type<float>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..9153635
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %float 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpCopyObject %float %float_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..1d57da4
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : f32 = 1.0;
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/f32-i32.wgsl b/test/expressions/bitcast/scalar/f32-i32.wgsl
new file mode 100644
index 0000000..c54ab72
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-i32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : f32 = 1.;
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..e7f1a21
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int b = asint(1.0f);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.msl b/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..a2bdd13
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 1.0f;
+ int const b = as_type<int>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..bea1419
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %float 1
+ %int = OpTypeInt 32 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %int %float_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..331cca2
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : f32 = 1.0;
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/f32-u32.wgsl b/test/expressions/bitcast/scalar/f32-u32.wgsl
new file mode 100644
index 0000000..0d79999
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-u32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : f32 = 1.;
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..289a1d5
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint b = asuint(1.0f);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.msl b/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..f0e4685
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 1.0f;
+ uint const b = as_type<uint>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..cc94dcd
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %float 1
+ %uint = OpTypeInt 32 0
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %uint %float_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..e52e25e
--- /dev/null
+++ b/test/expressions/bitcast/scalar/f32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : f32 = 1.0;
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/i32-f32.wgsl b/test/expressions/bitcast/scalar/i32-f32.wgsl
new file mode 100644
index 0000000..f398e96
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-f32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : i32 = 1;
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2fcce4f
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float b = asfloat(1);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.msl b/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..4368427
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 1;
+ float const b = as_type<float>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..1804815
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %int 1
+ %float = OpTypeFloat 32
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %float %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..9471a56
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : i32 = 1;
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/i32-i32.wgsl b/test/expressions/bitcast/scalar/i32-i32.wgsl
new file mode 100644
index 0000000..eaacdd1
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-i32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : i32 = 1;
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..4fc4846
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int b = asint(1);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.msl b/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..782176a
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 1;
+ int const b = as_type<int>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..35c7822
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %int 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpCopyObject %int %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..ed26420
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : i32 = 1;
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/i32-u32.wgsl b/test/expressions/bitcast/scalar/i32-u32.wgsl
new file mode 100644
index 0000000..f11e65d
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-u32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : i32 = 1;
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..d85da9a
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint b = asuint(1);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.msl b/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..dd82cdd
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 1;
+ uint const b = as_type<uint>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..c91c503
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %int 1
+ %uint = OpTypeInt 32 0
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %uint %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6c7b9f9
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : i32 = 1;
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/u32-f32.wgsl b/test/expressions/bitcast/scalar/u32-f32.wgsl
new file mode 100644
index 0000000..82c62b7
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-f32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : u32 = 1u;
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..1a00f3b
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float b = asfloat(1u);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.msl b/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..f019daf
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 1u;
+ float const b = as_type<float>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..daf433a
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %uint 1
+ %float = OpTypeFloat 32
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %float %uint_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..9efbcfd
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : u32 = 1u;
+ let b : f32 = bitcast<f32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/u32-i32.wgsl b/test/expressions/bitcast/scalar/u32-i32.wgsl
new file mode 100644
index 0000000..91b4230
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-i32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : u32 = 1u;
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2650dff
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int b = asint(1u);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.msl b/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..c6bd3b9
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 1u;
+ int const b = as_type<int>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..69997c4
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %uint 1
+ %int = OpTypeInt 32 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpBitcast %int %uint_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..9ac119d
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : u32 = 1u;
+ let b : i32 = bitcast<i32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/u32-u32.wgsl b/test/expressions/bitcast/scalar/u32-u32.wgsl
new file mode 100644
index 0000000..ceddfa2
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-u32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : u32 = 1u;
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..62d26be
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.hlsl
@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint b = asuint(1u);
+ return;
+}
diff --git a/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.msl b/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..01f3462
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 1u;
+ uint const b = as_type<uint>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..0a402a5
--- /dev/null
+++ b/test/expressions/bitcast/scalar/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_1 = OpConstant %uint 1
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %7 = OpCopyObject %uint %uint_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..4fde963
--- /dev/null
+++ b/test/expressions/bitcast/scalar/u32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : u32 = 1u;
+ let b : u32 = bitcast<u32>(a);
+}
diff --git a/test/expressions/bitcast/vector/f32-f32.wgsl b/test/expressions/bitcast/vector/f32-f32.wgsl
new file mode 100644
index 0000000..2e863d3
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-f32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<f32> = vec3<f32>(1., 2., 3.);
+ let b : vec3<f32> = bitcast<vec3<f32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/f32-f32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..e6ce403
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float3 a = float3(1.0f, 2.0f, 3.0f);
+ const float3 b = asfloat(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/f32-f32.wgsl.expected.msl b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..66dbb55
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float3 const a = float3(1.0f, 2.0f, 3.0f);
+ float3 const b = as_type<float3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/f32-f32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..2956553
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; 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
+ %v3float = OpTypeVector %float 3
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpCopyObject %v3float %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/f32-f32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..5ed2541
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<f32> = vec3<f32>(1.0, 2.0, 3.0);
+ let b : vec3<f32> = bitcast<vec3<f32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/f32-i32.wgsl b/test/expressions/bitcast/vector/f32-i32.wgsl
new file mode 100644
index 0000000..493dddd
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-i32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<f32> = vec3<f32>(1., 2., 3.);
+ let b : vec3<i32> = bitcast<vec3<i32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/f32-i32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..3ebeb51
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float3 a = float3(1.0f, 2.0f, 3.0f);
+ const int3 b = asint(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/f32-i32.wgsl.expected.msl b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..2d64209
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float3 const a = float3(1.0f, 2.0f, 3.0f);
+ int3 const b = as_type<int3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/f32-i32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..9061d7a
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; 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
+ %v3float = OpTypeVector %float 3
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpBitcast %v3int %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/f32-i32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..e4363c8
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<f32> = vec3<f32>(1.0, 2.0, 3.0);
+ let b : vec3<i32> = bitcast<vec3<i32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/f32-u32.wgsl b/test/expressions/bitcast/vector/f32-u32.wgsl
new file mode 100644
index 0000000..7daf410
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-u32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<f32> = vec3<f32>(1., 2., 3.);
+ let b : vec3<u32> = bitcast<vec3<u32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/f32-u32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..9490820
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float3 a = float3(1.0f, 2.0f, 3.0f);
+ const uint3 b = asuint(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/f32-u32.wgsl.expected.msl b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..b8a22f2
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float3 const a = float3(1.0f, 2.0f, 3.0f);
+ uint3 const b = as_type<uint3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/f32-u32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..85ebe09
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; 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
+ %v3float = OpTypeVector %float 3
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpBitcast %v3uint %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/f32-u32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6cd21fa
--- /dev/null
+++ b/test/expressions/bitcast/vector/f32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<f32> = vec3<f32>(1.0, 2.0, 3.0);
+ let b : vec3<u32> = bitcast<vec3<u32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/i32-f32.wgsl b/test/expressions/bitcast/vector/i32-f32.wgsl
new file mode 100644
index 0000000..e3bc280
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-f32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<i32> = vec3<i32>(1, 2, 3);
+ let b : vec3<f32> = bitcast<vec3<f32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/i32-f32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..73bb118
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int3 a = int3(1, 2, 3);
+ const float3 b = asfloat(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/i32-f32.wgsl.expected.msl b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..14d44ea
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int3 const a = int3(1, 2, 3);
+ float3 const b = as_type<float3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/i32-f32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..06884c2
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; 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
+ %v3int = OpTypeVector %int 3
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpBitcast %v3float %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/i32-f32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..a7a19ea
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<i32> = vec3<i32>(1, 2, 3);
+ let b : vec3<f32> = bitcast<vec3<f32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/i32-i32.wgsl b/test/expressions/bitcast/vector/i32-i32.wgsl
new file mode 100644
index 0000000..3149395
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-i32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<i32> = vec3<i32>(1, 2, 3);
+ let b : vec3<i32> = bitcast<vec3<i32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/i32-i32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2304f93
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int3 a = int3(1, 2, 3);
+ const int3 b = asint(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/i32-i32.wgsl.expected.msl b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..3c58a22
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int3 const a = int3(1, 2, 3);
+ int3 const b = as_type<int3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/i32-i32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..a56c4b6
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; 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
+ %v3int = OpTypeVector %int 3
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpCopyObject %v3int %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/i32-i32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..94f3acd
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<i32> = vec3<i32>(1, 2, 3);
+ let b : vec3<i32> = bitcast<vec3<i32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/i32-u32.wgsl b/test/expressions/bitcast/vector/i32-u32.wgsl
new file mode 100644
index 0000000..e3cf1fa
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-u32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<i32> = vec3<i32>(1, 2, 3);
+ let b : vec3<u32> = bitcast<vec3<u32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/i32-u32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..80e6646
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int3 a = int3(1, 2, 3);
+ const uint3 b = asuint(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/i32-u32.wgsl.expected.msl b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..34f7964
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int3 const a = int3(1, 2, 3);
+ uint3 const b = as_type<uint3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/i32-u32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..c4c7544
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; 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
+ %v3int = OpTypeVector %int 3
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpBitcast %v3uint %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/i32-u32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..64dd507
--- /dev/null
+++ b/test/expressions/bitcast/vector/i32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<i32> = vec3<i32>(1, 2, 3);
+ let b : vec3<u32> = bitcast<vec3<u32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/u32-f32.wgsl b/test/expressions/bitcast/vector/u32-f32.wgsl
new file mode 100644
index 0000000..4323361
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-f32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<u32> = vec3<u32>(1u, 2u, 3u);
+ let b : vec3<f32> = bitcast<vec3<f32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/u32-f32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..72084e9
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint3 a = uint3(1u, 2u, 3u);
+ const float3 b = asfloat(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/u32-f32.wgsl.expected.msl b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.msl
new file mode 100644
index 0000000..2376290
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint3 const a = uint3(1u, 2u, 3u);
+ float3 const b = as_type<float3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/u32-f32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..8545908
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; 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
+ %v3uint = OpTypeVector %uint 3
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+ %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpBitcast %v3float %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/u32-f32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..e74cdd3
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-f32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<u32> = vec3<u32>(1u, 2u, 3u);
+ let b : vec3<f32> = bitcast<vec3<f32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/u32-i32.wgsl b/test/expressions/bitcast/vector/u32-i32.wgsl
new file mode 100644
index 0000000..6efefc3
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-i32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<u32> = vec3<u32>(1u, 2u, 3u);
+ let b : vec3<i32> = bitcast<vec3<i32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/u32-i32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..6622323
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint3 a = uint3(1u, 2u, 3u);
+ const int3 b = asint(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/u32-i32.wgsl.expected.msl b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.msl
new file mode 100644
index 0000000..51c8f92
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint3 const a = uint3(1u, 2u, 3u);
+ int3 const b = as_type<int3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/u32-i32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..d59bdc2
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; 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
+ %v3uint = OpTypeVector %uint 3
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+ %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpBitcast %v3int %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/u32-i32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..db1bd62
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-i32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<u32> = vec3<u32>(1u, 2u, 3u);
+ let b : vec3<i32> = bitcast<vec3<i32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/u32-u32.wgsl b/test/expressions/bitcast/vector/u32-u32.wgsl
new file mode 100644
index 0000000..99d2211
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-u32.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<u32> = vec3<u32>(1u, 2u, 3u);
+ let b : vec3<u32> = bitcast<vec3<u32>>(a);
+}
diff --git a/test/expressions/bitcast/vector/u32-u32.wgsl.expected.hlsl b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..71f2935
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint3 a = uint3(1u, 2u, 3u);
+ const uint3 b = asuint(a);
+ return;
+}
diff --git a/test/expressions/bitcast/vector/u32-u32.wgsl.expected.msl b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.msl
new file mode 100644
index 0000000..43d1c6b
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint3 const a = uint3(1u, 2u, 3u);
+ uint3 const b = as_type<uint3>(a);
+ return;
+}
+
diff --git a/test/expressions/bitcast/vector/u32-u32.wgsl.expected.spvasm b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..08458e5
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; 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
+ %v3uint = OpTypeVector %uint 3
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+ %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %11 = OpCopyObject %v3uint %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/bitcast/vector/u32-u32.wgsl.expected.wgsl b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6bad22d
--- /dev/null
+++ b/test/expressions/bitcast/vector/u32-u32.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a : vec3<u32> = vec3<u32>(1u, 2u, 3u);
+ let b : vec3<u32> = bitcast<vec3<u32>>(a);
+}
