Add end2end tests for recent SPIR-V reader updates
- reduce output location value
- texture_depth_multisampled_2d
- matrix stride handling
- ldexp for unsigned second argument
Change-Id: Ib0782838dfa93fc637118cbf1faea67cf2406dd5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60460
Auto-Submit: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm
new file mode 100644
index 0000000..99a3dea
--- /dev/null
+++ b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm
@@ -0,0 +1,99 @@
+; Test: ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos SPIR-V Tools Assembler; 0
+; Bound: 100
+; Schema: 0
+OpCapability Shader
+OpCapability Sampled1D
+OpCapability Image1D
+OpCapability StorageImageExtendedFormats
+OpCapability ImageQuery
+OpMemoryModel Logical Simple
+OpEntryPoint Fragment %1 "main"
+OpExecutionMode %1 OriginUpperLeft
+OpName %f1 "f1"
+OpName %vf12 "vf12"
+OpName %vf123 "vf123"
+OpName %vf1234 "vf1234"
+OpName %u1 "u1"
+OpName %vu12 "vu12"
+OpName %vu123 "vu123"
+OpName %vu1234 "vu1234"
+OpName %i1 "i1"
+OpName %vi12 "vi12"
+OpName %vi123 "vi123"
+OpName %vi1234 "vi1234"
+OpName %offsets2d "offsets2d"
+OpDecorate %20 DescriptorSet 2
+OpDecorate %20 Binding 1
+%void = OpTypeVoid
+%16 = OpTypeFunction %void
+%float = OpTypeFloat 32
+%uint = OpTypeInt 32 0
+%int = OpTypeInt 32 1
+%int_0 = OpConstant %int 0
+%int_1 = OpConstant %int 1
+%int_2 = OpConstant %int 2
+%int_3 = OpConstant %int 3
+%int_4 = OpConstant %int 4
+%uint_0 = OpConstant %uint 0
+%uint_1 = OpConstant %uint 1
+%uint_2 = OpConstant %uint 2
+%uint_3 = OpConstant %uint 3
+%uint_4 = OpConstant %uint 4
+%uint_100 = OpConstant %uint 100
+%v2int = OpTypeVector %int 2
+%v3int = OpTypeVector %int 3
+%v4int = OpTypeVector %int 4
+%v2uint = OpTypeVector %uint 2
+%v3uint = OpTypeVector %uint 3
+%v4uint = OpTypeVector %uint 4
+%v2float = OpTypeVector %float 2
+%v3float = OpTypeVector %float 3
+%v4float = OpTypeVector %float 4
+%41 = OpConstantNull %float
+%float_0 = OpConstant %float 0
+%float_1 = OpConstant %float 1
+%float_2 = OpConstant %float 2
+%float_3 = OpConstant %float 3
+%float_4 = OpConstant %float 4
+%float_7 = OpConstant %float 7
+%48 = OpConstantNull %v2float
+%49 = OpConstantNull %v3float
+%50 = OpConstantNull %v4float
+%51 = OpConstantComposite %v2int %int_1 %int_2
+%52 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+%53 = OpConstantComposite %v4int %int_1 %int_2 %int_3 %int_4
+%54 = OpConstantComposite %v2uint %uint_1 %uint_2
+%55 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+%56 = OpConstantComposite %v4uint %uint_1 %uint_2 %uint_3 %uint_4
+%57 = OpConstantComposite %v2float %float_1 %float_2
+%58 = OpConstantComposite %v2float %float_2 %float_1
+%59 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+%60 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
+%float_0_200000003 = OpConstant %float 0.200000003
+%62 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_62 = OpTypePointer UniformConstant %62
+%20 = OpVariable %_ptr_UniformConstant_62 UniformConstant
+%1 = OpFunction %void None %16
+%64 = OpLabel
+%f1 = OpCopyObject %float %float_1
+%vf12 = OpCopyObject %v2float %57
+%vf123 = OpCopyObject %v3float %59
+%vf1234 = OpCopyObject %v4float %60
+%i1 = OpCopyObject %int %int_1
+%vi12 = OpCopyObject %v2int %51
+%vi123 = OpCopyObject %v3int %52
+%vi1234 = OpCopyObject %v4int %53
+%u1 = OpCopyObject %uint %uint_1
+%vu12 = OpCopyObject %v2uint %54
+%vu123 = OpCopyObject %v3uint %55
+%vu1234 = OpCopyObject %v4uint %56
+%65 = OpCompositeConstruct %v2int %int_3 %int_4
+%offsets2d = OpCopyObject %v2int %65
+%66 = OpLoad %62 %20
+%99 = OpImageFetch %v4float %66 %vi12 Sample %i1
+OpReturn
+OpFunctionEnd
+
diff --git a/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.hlsl b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.hlsl
new file mode 100644
index 0000000..94c1196
--- /dev/null
+++ b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.hlsl
@@ -0,0 +1,24 @@
+Texture2DMS<float4> x_20 : register(t1, space2);
+
+void main_1() {
+ const float f1 = 1.0f;
+ const float2 vf12 = float2(1.0f, 2.0f);
+ const float3 vf123 = float3(1.0f, 2.0f, 3.0f);
+ const float4 vf1234 = float4(1.0f, 2.0f, 3.0f, 4.0f);
+ const int i1 = 1;
+ const int2 vi12 = int2(1, 2);
+ const int3 vi123 = int3(1, 2, 3);
+ const int4 vi1234 = int4(1, 2, 3, 4);
+ const uint u1 = 1u;
+ const uint2 vu12 = uint2(1u, 2u);
+ const uint3 vu123 = uint3(1u, 2u, 3u);
+ const uint4 vu1234 = uint4(1u, 2u, 3u, 4u);
+ const int2 offsets2d = int2(3, 4);
+ const float4 x_99 = float4(x_20.Load(int3(vi12, 0), i1).x, 0.0f, 0.0f, 0.0f);
+ return;
+}
+
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.msl b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.msl
new file mode 100644
index 0000000..a41a176
--- /dev/null
+++ b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.msl
@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void main_1(depth2d_ms<float, access::read> tint_symbol_1) {
+ float const f1 = 1.0f;
+ float2 const vf12 = float2(1.0f, 2.0f);
+ float3 const vf123 = float3(1.0f, 2.0f, 3.0f);
+ float4 const vf1234 = float4(1.0f, 2.0f, 3.0f, 4.0f);
+ int const i1 = 1;
+ int2 const vi12 = int2(1, 2);
+ int3 const vi123 = int3(1, 2, 3);
+ int4 const vi1234 = int4(1, 2, 3, 4);
+ uint const u1 = 1u;
+ uint2 const vu12 = uint2(1u, 2u);
+ uint3 const vu123 = uint3(1u, 2u, 3u);
+ uint4 const vu1234 = uint4(1u, 2u, 3u, 4u);
+ int2 const offsets2d = int2(3, 4);
+ float4 const x_99 = float4(tint_symbol_1.read(uint2(vi12), i1), 0.0f, 0.0f, 0.0f);
+ return;
+}
+
+fragment void tint_symbol(depth2d_ms<float, access::read> tint_symbol_2 [[texture(1)]]) {
+ main_1(tint_symbol_2);
+ return;
+}
+
diff --git a/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.spvasm b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.spvasm
new file mode 100644
index 0000000..a49c62c
--- /dev/null
+++ b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 50
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main"
+ OpExecutionMode %main OriginUpperLeft
+ OpName %x_20 "x_20"
+ OpName %main_1 "main_1"
+ OpName %main "main"
+ OpDecorate %x_20 DescriptorSet 2
+ OpDecorate %x_20 Binding 1
+ %float = OpTypeFloat 32
+ %3 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+ %x_20 = OpVariable %_ptr_UniformConstant_3 UniformConstant
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %float_1 = OpConstant %float 1
+ %v2float = OpTypeVector %float 2
+ %float_2 = OpConstant %float 2
+ %12 = OpConstantComposite %v2float %float_1 %float_2
+ %v3float = OpTypeVector %float 3
+ %float_3 = OpConstant %float 3
+ %15 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %v4float = OpTypeVector %float 4
+ %float_4 = OpConstant %float 4
+ %18 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
+ %int = OpTypeInt 32 1
+ %int_1 = OpConstant %int 1
+ %v2int = OpTypeVector %int 2
+ %int_2 = OpConstant %int 2
+ %23 = OpConstantComposite %v2int %int_1 %int_2
+ %v3int = OpTypeVector %int 3
+ %int_3 = OpConstant %int 3
+ %26 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+ %v4int = OpTypeVector %int 4
+ %int_4 = OpConstant %int 4
+ %29 = OpConstantComposite %v4int %int_1 %int_2 %int_3 %int_4
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+ %v2uint = OpTypeVector %uint 2
+ %uint_2 = OpConstant %uint 2
+ %34 = OpConstantComposite %v2uint %uint_1 %uint_2
+ %v3uint = OpTypeVector %uint 3
+ %uint_3 = OpConstant %uint 3
+ %37 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+ %v4uint = OpTypeVector %uint 4
+ %uint_4 = OpConstant %uint 4
+ %40 = OpConstantComposite %v4uint %uint_1 %uint_2 %uint_3 %uint_4
+ %41 = OpConstantComposite %v2int %int_3 %int_4
+ %float_0 = OpConstant %float 0
+ %main_1 = OpFunction %void None %5
+ %8 = OpLabel
+ %44 = OpLoad %3 %x_20
+ %43 = OpImageFetch %v4float %44 %23 Sample %int_1
+ %42 = OpCompositeExtract %float %43 0
+ %46 = OpCompositeConstruct %v4float %42 %float_0 %float_0 %float_0
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %5
+ %48 = OpLabel
+ %49 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.wgsl b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.wgsl
new file mode 100644
index 0000000..1f17b47
--- /dev/null
+++ b/test/unittest/reader/spirv/ImageFetch_DepthMultisampled_SpvParserHandleTest_ImageAccessTest_Variable_0.spvasm.expected.wgsl
@@ -0,0 +1,24 @@
+[[group(2), binding(1)]] var x_20 : texture_depth_multisampled_2d;
+
+fn main_1() {
+ let f1 : f32 = 1.0;
+ let vf12 : vec2<f32> = vec2<f32>(1.0, 2.0);
+ let vf123 : vec3<f32> = vec3<f32>(1.0, 2.0, 3.0);
+ let vf1234 : vec4<f32> = vec4<f32>(1.0, 2.0, 3.0, 4.0);
+ let i1 : i32 = 1;
+ let vi12 : vec2<i32> = vec2<i32>(1, 2);
+ let vi123 : vec3<i32> = vec3<i32>(1, 2, 3);
+ let vi1234 : vec4<i32> = vec4<i32>(1, 2, 3, 4);
+ let u1 : u32 = 1u;
+ let vu12 : vec2<u32> = vec2<u32>(1u, 2u);
+ let vu123 : vec3<u32> = vec3<u32>(1u, 2u, 3u);
+ let vu1234 : vec4<u32> = vec4<u32>(1u, 2u, 3u, 4u);
+ let offsets2d : vec2<i32> = vec2<i32>(3, 4);
+ let x_99 : vec4<f32> = vec4<f32>(textureLoad(x_20, vi12, i1), 0.0, 0.0, 0.0);
+ return;
+}
+
+[[stage(fragment)]]
+fn main() {
+ main_1();
+}
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm
index 55e55ee..bae6f59 100644
--- a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm
@@ -12,7 +12,7 @@
OpDecorate %1 Location 0
OpDecorate %2 Location 0
OpDecorate %3 Location 30
-OpDecorate %4 Location 40
+OpDecorate %4 Location 6
%void = OpTypeVoid
%7 = OpTypeFunction %void
%bool = OpTypeBool
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.hlsl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.hlsl
index 7017765..d2eadc7 100644
--- a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.hlsl
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
static uint x_1 = 0u;
static uint x_2 = 0u;
static uint x_3 = 0u;
@@ -19,7 +17,7 @@
};
struct tint_symbol_2 {
uint x_2_1 : SV_Target0;
- uint x_4_1 : SV_Target40;
+ uint x_4_1 : SV_Target6;
};
tint_symbol_2 main(tint_symbol_1 tint_symbol) {
@@ -32,11 +30,3 @@
const tint_symbol_2 tint_symbol_4 = {tint_symbol_3.x_2_1, tint_symbol_3.x_4_1};
return tint_symbol_4;
}
-warning: DXIL.dll not found. Resulting DXIL will not be signed for use in release environments.
-
-error: validation errors
-error: SV_Target semantic index exceeds maximum (7).
-Validation failed.
-
-
-
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.spvasm b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.spvasm
index 07d96dd..85a1b1b 100644
--- a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.spvasm
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.spvasm
@@ -27,7 +27,7 @@
OpDecorate %tint_symbol_1 Location 30
OpDecorate %tint_symbol_1 Flat
OpDecorate %tint_symbol_3 Location 0
- OpDecorate %tint_symbol_4 Location 40
+ OpDecorate %tint_symbol_4 Location 6
OpMemberDecorate %main_out 0 Offset 0
OpMemberDecorate %main_out 1 Offset 4
%uint = OpTypeInt 32 0
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.wgsl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.wgsl
index 542ae52..59ed023 100644
--- a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.wgsl
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_EntryPointWrapping_IOLocations.spvasm.expected.wgsl
@@ -13,7 +13,7 @@
struct main_out {
[[location(0)]]
x_2_1 : u32;
- [[location(40)]]
+ [[location(6)]]
x_4_1 : u32;
};
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.hlsl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.hlsl
new file mode 100644
index 0000000..ef0855d
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.hlsl
@@ -0,0 +1,10 @@
+RWByteAddressBuffer myvar : register(u0, space0);
+
+void main_1() {
+ return;
+}
+
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.msl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.msl
new file mode 100644
index 0000000..b2a51ad
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.msl
@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_padded_array_element {
+ /* 0x0000 */ packed_float2 el;
+ /* 0x0008 */ int8_t tint_pad[56];
+};
+struct tint_array_wrapper {
+ /* 0x0000 */ tint_padded_array_element arr[3];
+};
+struct S {
+ /* 0x0000 */ tint_array_wrapper field0;
+};
+
+void main_1() {
+ return;
+}
+
+fragment void tint_symbol() {
+ main_1();
+ return;
+}
+
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.spvasm b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.spvasm
new file mode 100644
index 0000000..33921e1
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.spvasm
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main"
+ OpExecutionMode %main OriginUpperLeft
+ OpName %S "S"
+ OpMemberName %S 0 "field0"
+ OpName %myvar "myvar"
+ OpName %main_1 "main_1"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %_arr_v2float_uint_3 ArrayStride 64
+ OpDecorate %myvar DescriptorSet 0
+ OpDecorate %myvar Binding 0
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+ %uint = OpTypeInt 32 0
+ %uint_3 = OpConstant %uint 3
+%_arr_v2float_uint_3 = OpTypeArray %v2float %uint_3
+ %S = OpTypeStruct %_arr_v2float_uint_3
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %myvar = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %main_1 = OpFunction %void None %9
+ %12 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %9
+ %14 = OpLabel
+ %15 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.wgsl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.wgsl
new file mode 100644
index 0000000..996fa56
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration.spvasm.expected.wgsl
@@ -0,0 +1,15 @@
+[[block]]
+struct S {
+ field0 : [[stride(64)]] array<vec2<f32>, 3>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> myvar : S;
+
+fn main_1() {
+ return;
+}
+
+[[stage(fragment)]]
+fn main() {
+ main_1();
+}
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.hlsl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.hlsl
new file mode 100644
index 0000000..ef0855d
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.hlsl
@@ -0,0 +1,10 @@
+RWByteAddressBuffer myvar : register(u0, space0);
+
+void main_1() {
+ return;
+}
+
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.msl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.msl
new file mode 100644
index 0000000..ecdf927
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.msl
@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ float3x2 field0;
+};
+
+void main_1() {
+ return;
+}
+
+fragment void tint_symbol() {
+ main_1();
+ return;
+}
+
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.spvasm b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.spvasm
new file mode 100644
index 0000000..64446e5
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main"
+ OpExecutionMode %main OriginUpperLeft
+ OpName %S "S"
+ OpMemberName %S 0 "field0"
+ OpName %myvar "myvar"
+ OpName %main_1 "main_1"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 ColMajor
+ OpMemberDecorate %S 0 MatrixStride 8
+ OpDecorate %myvar DescriptorSet 0
+ OpDecorate %myvar Binding 0
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %S = OpTypeStruct %mat3v2float
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %myvar = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %main_1 = OpFunction %void None %7
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %7
+ %12 = OpLabel
+ %13 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.wgsl b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.wgsl
new file mode 100644
index 0000000..79b8f48
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvModuleScopeVarParserTest_MatrixStrideDecoration_Natural_Dropped.spvasm.expected.wgsl
@@ -0,0 +1,15 @@
+[[block]]
+struct S {
+ field0 : mat3x2<f32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> myvar : S;
+
+fn main_1() {
+ return;
+}
+
+[[stage(fragment)]]
+fn main() {
+ main_1();
+}
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.hlsl b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.hlsl
new file mode 100644
index 0000000..a889e78
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.hlsl
@@ -0,0 +1,32 @@
+void main_1() {
+ const uint u1 = 10u;
+ const uint u2 = 15u;
+ const uint u3 = 20u;
+ const int i1 = 30;
+ const int i2 = 35;
+ const int i3 = 40;
+ const float f1 = 50.0f;
+ const float f2 = 60.0f;
+ const float f3 = 70.0f;
+ const uint2 v2u1 = uint2(10u, 20u);
+ const uint2 v2u2 = uint2(20u, 10u);
+ const uint2 v2u3 = uint2(15u, 15u);
+ const int2 v2i1 = int2(30, 40);
+ const int2 v2i2 = int2(40, 30);
+ const int2 v2i3 = int2(35, 35);
+ const float2 v2f1 = float2(50.0f, 60.0f);
+ const float2 v2f2 = float2(60.0f, 50.0f);
+ const float2 v2f3 = float2(70.0f, 70.0f);
+ const float3 v3f1 = float3(50.0f, 60.0f, 70.0f);
+ const float3 v3f2 = float3(60.0f, 70.0f, 50.0f);
+ const float4 v4f1 = float4(50.0f, 50.0f, 50.0f, 50.0f);
+ const float4 v4f2 = v4f1;
+ const float x_1 = ldexp(f1, int(u1));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.msl b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.msl
new file mode 100644
index 0000000..9484dda
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void main_1() {
+ uint const u1 = 10u;
+ uint const u2 = 15u;
+ uint const u3 = 20u;
+ int const i1 = 30;
+ int const i2 = 35;
+ int const i3 = 40;
+ float const f1 = 50.0f;
+ float const f2 = 60.0f;
+ float const f3 = 70.0f;
+ uint2 const v2u1 = uint2(10u, 20u);
+ uint2 const v2u2 = uint2(20u, 10u);
+ uint2 const v2u3 = uint2(15u, 15u);
+ int2 const v2i1 = int2(30, 40);
+ int2 const v2i2 = int2(40, 30);
+ int2 const v2i3 = int2(35, 35);
+ float2 const v2f1 = float2(50.0f, 60.0f);
+ float2 const v2f2 = float2(60.0f, 50.0f);
+ float2 const v2f3 = float2(70.0f, 70.0f);
+ float3 const v3f1 = float3(50.0f, 60.0f, 70.0f);
+ float3 const v3f2 = float3(60.0f, 70.0f, 50.0f);
+ float4 const v4f1 = float4(50.0f, 50.0f, 50.0f, 50.0f);
+ float4 const v4f2 = v4f1;
+ float const x_1 = ldexp(f1, int(u1));
+ return;
+}
+
+kernel void tint_symbol() {
+ main_1();
+ return;
+}
+
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.spvasm b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.spvasm
new file mode 100644
index 0000000..2eab6d4
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.spvasm
@@ -0,0 +1,54 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+ OpCapability Shader
+ %35 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main_1 "main_1"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_10 = OpConstant %uint 10
+ %uint_15 = OpConstant %uint 15
+ %uint_20 = OpConstant %uint 20
+ %int = OpTypeInt 32 1
+ %int_30 = OpConstant %int 30
+ %int_35 = OpConstant %int 35
+ %int_40 = OpConstant %int 40
+ %float = OpTypeFloat 32
+ %float_50 = OpConstant %float 50
+ %float_60 = OpConstant %float 60
+ %float_70 = OpConstant %float 70
+ %v2uint = OpTypeVector %uint 2
+ %18 = OpConstantComposite %v2uint %uint_10 %uint_20
+ %19 = OpConstantComposite %v2uint %uint_20 %uint_10
+ %20 = OpConstantComposite %v2uint %uint_15 %uint_15
+ %v2int = OpTypeVector %int 2
+ %22 = OpConstantComposite %v2int %int_30 %int_40
+ %23 = OpConstantComposite %v2int %int_40 %int_30
+ %24 = OpConstantComposite %v2int %int_35 %int_35
+ %v2float = OpTypeVector %float 2
+ %26 = OpConstantComposite %v2float %float_50 %float_60
+ %27 = OpConstantComposite %v2float %float_60 %float_50
+ %28 = OpConstantComposite %v2float %float_70 %float_70
+ %v3float = OpTypeVector %float 3
+ %30 = OpConstantComposite %v3float %float_50 %float_60 %float_70
+ %31 = OpConstantComposite %v3float %float_60 %float_70 %float_50
+ %v4float = OpTypeVector %float 4
+ %33 = OpConstantComposite %v4float %float_50 %float_50 %float_50 %float_50
+ %main_1 = OpFunction %void None %1
+ %4 = OpLabel
+ %36 = OpBitcast %int %uint_10
+ %34 = OpExtInst %float %35 Ldexp %float_50 %36
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.wgsl b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.wgsl
new file mode 100644
index 0000000..45e0784
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Scalar_Float_Uint.spvasm.expected.wgsl
@@ -0,0 +1,31 @@
+fn main_1() {
+ let u1 : u32 = 10u;
+ let u2 : u32 = 15u;
+ let u3 : u32 = 20u;
+ let i1 : i32 = 30;
+ let i2 : i32 = 35;
+ let i3 : i32 = 40;
+ let f1 : f32 = 50.0;
+ let f2 : f32 = 60.0;
+ let f3 : f32 = 70.0;
+ let v2u1 : vec2<u32> = vec2<u32>(10u, 20u);
+ let v2u2 : vec2<u32> = vec2<u32>(20u, 10u);
+ let v2u3 : vec2<u32> = vec2<u32>(15u, 15u);
+ let v2i1 : vec2<i32> = vec2<i32>(30, 40);
+ let v2i2 : vec2<i32> = vec2<i32>(40, 30);
+ let v2i3 : vec2<i32> = vec2<i32>(35, 35);
+ let v2f1 : vec2<f32> = vec2<f32>(50.0, 60.0);
+ let v2f2 : vec2<f32> = vec2<f32>(60.0, 50.0);
+ let v2f3 : vec2<f32> = vec2<f32>(70.0, 70.0);
+ let v3f1 : vec3<f32> = vec3<f32>(50.0, 60.0, 70.0);
+ let v3f2 : vec3<f32> = vec3<f32>(60.0, 70.0, 50.0);
+ let v4f1 : vec4<f32> = vec4<f32>(50.0, 50.0, 50.0, 50.0);
+ let v4f2 : vec4<f32> = v4f1;
+ let x_1 : f32 = ldexp(f1, i32(u1));
+ return;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ main_1();
+}
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.hlsl b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.hlsl
new file mode 100644
index 0000000..8968843
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.hlsl
@@ -0,0 +1,32 @@
+void main_1() {
+ const uint u1 = 10u;
+ const uint u2 = 15u;
+ const uint u3 = 20u;
+ const int i1 = 30;
+ const int i2 = 35;
+ const int i3 = 40;
+ const float f1 = 50.0f;
+ const float f2 = 60.0f;
+ const float f3 = 70.0f;
+ const uint2 v2u1 = uint2(10u, 20u);
+ const uint2 v2u2 = uint2(20u, 10u);
+ const uint2 v2u3 = uint2(15u, 15u);
+ const int2 v2i1 = int2(30, 40);
+ const int2 v2i2 = int2(40, 30);
+ const int2 v2i3 = int2(35, 35);
+ const float2 v2f1 = float2(50.0f, 60.0f);
+ const float2 v2f2 = float2(60.0f, 50.0f);
+ const float2 v2f3 = float2(70.0f, 70.0f);
+ const float3 v3f1 = float3(50.0f, 60.0f, 70.0f);
+ const float3 v3f2 = float3(60.0f, 70.0f, 50.0f);
+ const float4 v4f1 = float4(50.0f, 50.0f, 50.0f, 50.0f);
+ const float4 v4f2 = v4f1;
+ const float2 x_1 = ldexp(v2f1, int2(v2u1));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.msl b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.msl
new file mode 100644
index 0000000..937bf3d
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void main_1() {
+ uint const u1 = 10u;
+ uint const u2 = 15u;
+ uint const u3 = 20u;
+ int const i1 = 30;
+ int const i2 = 35;
+ int const i3 = 40;
+ float const f1 = 50.0f;
+ float const f2 = 60.0f;
+ float const f3 = 70.0f;
+ uint2 const v2u1 = uint2(10u, 20u);
+ uint2 const v2u2 = uint2(20u, 10u);
+ uint2 const v2u3 = uint2(15u, 15u);
+ int2 const v2i1 = int2(30, 40);
+ int2 const v2i2 = int2(40, 30);
+ int2 const v2i3 = int2(35, 35);
+ float2 const v2f1 = float2(50.0f, 60.0f);
+ float2 const v2f2 = float2(60.0f, 50.0f);
+ float2 const v2f3 = float2(70.0f, 70.0f);
+ float3 const v3f1 = float3(50.0f, 60.0f, 70.0f);
+ float3 const v3f2 = float3(60.0f, 70.0f, 50.0f);
+ float4 const v4f1 = float4(50.0f, 50.0f, 50.0f, 50.0f);
+ float4 const v4f2 = v4f1;
+ float2 const x_1 = ldexp(v2f1, int2(v2u1));
+ return;
+}
+
+kernel void tint_symbol() {
+ main_1();
+ return;
+}
+
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.spvasm b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.spvasm
new file mode 100644
index 0000000..69e8594
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.spvasm
@@ -0,0 +1,54 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+ OpCapability Shader
+ %35 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main_1 "main_1"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_10 = OpConstant %uint 10
+ %uint_15 = OpConstant %uint 15
+ %uint_20 = OpConstant %uint 20
+ %int = OpTypeInt 32 1
+ %int_30 = OpConstant %int 30
+ %int_35 = OpConstant %int 35
+ %int_40 = OpConstant %int 40
+ %float = OpTypeFloat 32
+ %float_50 = OpConstant %float 50
+ %float_60 = OpConstant %float 60
+ %float_70 = OpConstant %float 70
+ %v2uint = OpTypeVector %uint 2
+ %18 = OpConstantComposite %v2uint %uint_10 %uint_20
+ %19 = OpConstantComposite %v2uint %uint_20 %uint_10
+ %20 = OpConstantComposite %v2uint %uint_15 %uint_15
+ %v2int = OpTypeVector %int 2
+ %22 = OpConstantComposite %v2int %int_30 %int_40
+ %23 = OpConstantComposite %v2int %int_40 %int_30
+ %24 = OpConstantComposite %v2int %int_35 %int_35
+ %v2float = OpTypeVector %float 2
+ %26 = OpConstantComposite %v2float %float_50 %float_60
+ %27 = OpConstantComposite %v2float %float_60 %float_50
+ %28 = OpConstantComposite %v2float %float_70 %float_70
+ %v3float = OpTypeVector %float 3
+ %30 = OpConstantComposite %v3float %float_50 %float_60 %float_70
+ %31 = OpConstantComposite %v3float %float_60 %float_70 %float_50
+ %v4float = OpTypeVector %float 4
+ %33 = OpConstantComposite %v4float %float_50 %float_50 %float_50 %float_50
+ %main_1 = OpFunction %void None %1
+ %4 = OpLabel
+ %36 = OpBitcast %v2int %18
+ %34 = OpExtInst %v2float %35 Ldexp %26 %36
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.wgsl b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.wgsl
new file mode 100644
index 0000000..100a2a9
--- /dev/null
+++ b/test/unittest/reader/spirv/SpvParserTest_GlslStd450_Ldexp_Vector_Floatvec_Uintvec.spvasm.expected.wgsl
@@ -0,0 +1,31 @@
+fn main_1() {
+ let u1 : u32 = 10u;
+ let u2 : u32 = 15u;
+ let u3 : u32 = 20u;
+ let i1 : i32 = 30;
+ let i2 : i32 = 35;
+ let i3 : i32 = 40;
+ let f1 : f32 = 50.0;
+ let f2 : f32 = 60.0;
+ let f3 : f32 = 70.0;
+ let v2u1 : vec2<u32> = vec2<u32>(10u, 20u);
+ let v2u2 : vec2<u32> = vec2<u32>(20u, 10u);
+ let v2u3 : vec2<u32> = vec2<u32>(15u, 15u);
+ let v2i1 : vec2<i32> = vec2<i32>(30, 40);
+ let v2i2 : vec2<i32> = vec2<i32>(40, 30);
+ let v2i3 : vec2<i32> = vec2<i32>(35, 35);
+ let v2f1 : vec2<f32> = vec2<f32>(50.0, 60.0);
+ let v2f2 : vec2<f32> = vec2<f32>(60.0, 50.0);
+ let v2f3 : vec2<f32> = vec2<f32>(70.0, 70.0);
+ let v3f1 : vec3<f32> = vec3<f32>(50.0, 60.0, 70.0);
+ let v3f2 : vec3<f32> = vec3<f32>(60.0, 70.0, 50.0);
+ let v4f1 : vec4<f32> = vec4<f32>(50.0, 50.0, 50.0, 50.0);
+ let v4f2 : vec4<f32> = v4f1;
+ let x_1 : vec2<f32> = ldexp(v2f1, vec2<i32>(v2u1));
+ return;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ main_1();
+}
