Add SPVASM test from clspv.

reference/shaders-hlsl/asm/comp/storage-buffer-basic.asm.comp

@@ -0,0 +1,26 @@
+const uint _3 = 1u;
+const uint _4 = 3u;
+const uint3 gl_WorkGroupSize = uint3(_3, 2u, _4);
+
+RWByteAddressBuffer _8 : register(u0);
+RWByteAddressBuffer _9 : register(u1);
+
+static uint3 gl_WorkGroupID;
+struct SPIRV_Cross_Input
+{
+    uint3 gl_WorkGroupID : SV_GroupID;
+};
+
+static uint3 _22 = gl_WorkGroupSize;
+
+void comp_main()
+{
+    _8.Store(gl_WorkGroupID.x * 4 + 0, asuint(asfloat(_9.Load(gl_WorkGroupID.x * 4 + 0)) + asfloat(_8.Load(gl_WorkGroupID.x * 4 + 0))));
+}
+
+[numthreads(1, 2, 3)]
+void main(SPIRV_Cross_Input stage_input)
+{
+    gl_WorkGroupID = stage_input.gl_WorkGroupID;
+    comp_main();
+}

reference/shaders-msl/asm/comp/storage-buffer-basic.asm.comp

@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+constant uint _3_tmp [[function_constant(0)]];
+constant uint _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 1u;
+constant uint _4_tmp [[function_constant(2)]];
+constant uint _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 3u;
+constant uint3 gl_WorkGroupSize = uint3(_3, 2u, _4);
+
+struct _6
+{
+    float _m0[];
+};
+
+kernel void main0(uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], device _6& _8 [[buffer(0)]], device _6& _9 [[buffer(1)]])
+{
+    uint3 _22 = gl_WorkGroupSize;
+    _8._m0[gl_WorkGroupID.x] = _9._m0[gl_WorkGroupID.x] + _8._m0[gl_WorkGroupID.x];
+}
+

reference/shaders/asm/comp/bitcast_iadd.asm.comp

@@ -1,13 +1,13 @@
 #version 310 es
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-layout(binding = 0, std430) restrict buffer _3
+layout(binding = 0, std430) restrict buffer _3_5
 {
     ivec4 _m0;
     uvec4 _m1;
 } _5;
 
-layout(binding = 1, std430) restrict buffer _4
+layout(binding = 1, std430) restrict buffer _4_6
 {
     uvec4 _m0;
     ivec4 _m1;

reference/shaders/asm/comp/bitcast_iequal.asm.comp

@@ -1,13 +1,13 @@
 #version 310 es
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-layout(binding = 0, std430) buffer _3
+layout(binding = 0, std430) buffer _3_5
 {
     ivec4 _m0;
     uvec4 _m1;
 } _5;
 
-layout(binding = 1, std430) buffer _4
+layout(binding = 1, std430) buffer _4_6
 {
     uvec4 _m0;
     ivec4 _m1;

reference/shaders/asm/comp/bitcast_sar.asm.comp

@@ -1,13 +1,13 @@
 #version 310 es
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-layout(binding = 0, std430) buffer _3
+layout(binding = 0, std430) buffer _3_5
 {
     ivec4 _m0;
     uvec4 _m1;
 } _5;
 
-layout(binding = 1, std430) buffer _4
+layout(binding = 1, std430) buffer _4_6
 {
     uvec4 _m0;
     ivec4 _m1;

reference/shaders/asm/comp/bitcast_sdiv.asm.comp

@@ -1,13 +1,13 @@
 #version 310 es
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-layout(binding = 0, std430) buffer _3
+layout(binding = 0, std430) buffer _3_5
 {
     ivec4 _m0;
     uvec4 _m1;
 } _5;
 
-layout(binding = 1, std430) buffer _4
+layout(binding = 1, std430) buffer _4_6
 {
     uvec4 _m0;
     ivec4 _m1;

reference/shaders/asm/comp/bitcast_slr.asm.comp

@@ -1,13 +1,13 @@
 #version 310 es
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-layout(binding = 0, std430) buffer _3
+layout(binding = 0, std430) buffer _3_5
 {
     ivec4 _m0;
     uvec4 _m1;
 } _5;
 
-layout(binding = 1, std430) buffer _4
+layout(binding = 1, std430) buffer _4_6
 {
     uvec4 _m0;
     ivec4 _m1;

reference/shaders/asm/comp/multiple-entry.asm.comp

@@ -1,13 +1,13 @@
 #version 310 es
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-layout(binding = 0, std430) restrict buffer _6
+layout(binding = 0, std430) restrict buffer _6_8
 {
     ivec4 _m0;
     uvec4 _m1;
 } _8;
 
-layout(binding = 1, std430) restrict buffer _7
+layout(binding = 1, std430) restrict buffer _7_9
 {
     uvec4 _m0;
     ivec4 _m1;

reference/shaders/asm/comp/name-alias.asm.invalid.comp

@@ -19,12 +19,12 @@ struct alias_2
     alias_1 alias_1;
 };
 
-layout(binding = 0, std430) buffer _10
+layout(binding = 0, std430) buffer _10_11
 {
     alias_2 alias;
 } alias_3;
 
-layout(binding = 1, std140) buffer _15
+layout(binding = 1, std140) buffer _15_16
 {
     alias_2 alias;
 } alias_4;

reference/shaders/asm/comp/storage-buffer-basic.asm.comp

@@ -0,0 +1,20 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 2, local_size_z = 3) in;
+
+layout(binding = 0, std430) buffer _6_8
+{
+    float _m0[];
+} _8;
+
+layout(binding = 1, std430) buffer _6_9
+{
+    float _m0[];
+} _9;
+
+uvec3 _22 = gl_WorkGroupSize;
+
+void main()
+{
+    _8._m0[gl_WorkGroupID.x] = _9._m0[gl_WorkGroupID.x] + _8._m0[gl_WorkGroupID.x];
+}
+

shaders-hlsl/asm/comp/storage-buffer-basic.asm.comp

@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Codeplay; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpCapability VariablePointers
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpExtension "SPV_KHR_variable_pointers"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %22 "main" %gl_WorkGroupID
+               OpSource OpenCL_C 120
+               OpDecorate %15 SpecId 0
+               ;OpDecorate %16 SpecId 1
+               OpDecorate %17 SpecId 2
+               OpDecorate %_runtimearr_float ArrayStride 4
+               OpMemberDecorate %_struct_4 0 Offset 0
+               OpDecorate %_struct_4 Block
+               OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId
+               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+               OpDecorate %20 DescriptorSet 0
+               OpDecorate %20 Binding 0
+               OpDecorate %21 DescriptorSet 0
+               OpDecorate %21 Binding 1
+      %float = OpTypeFloat 32
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+%_runtimearr_float = OpTypeRuntimeArray %float
+  %_struct_4 = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer__struct_4 = OpTypePointer StorageBuffer %_struct_4
+       %uint = OpTypeInt 32 0
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Private_v3uint = OpTypePointer Private %v3uint
+     %uint_0 = OpConstant %uint 0
+%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input
+         %15 = OpSpecConstant %uint 1
+         %16 = OpConstant %uint 2
+         %17 = OpSpecConstant %uint 3
+%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %15 %16 %17
+         %19 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize
+         %20 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
+         %21 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
+         %22 = OpFunction %void None %8
+         %23 = OpLabel
+         %24 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0
+         %25 = OpLoad %uint %24
+         %26 = OpAccessChain %_ptr_StorageBuffer_float %21 %uint_0 %25
+         %27 = OpLoad %float %26
+         %28 = OpAccessChain %_ptr_StorageBuffer_float %20 %uint_0 %25
+         %29 = OpLoad %float %28
+         %30 = OpFAdd %float %27 %29
+               OpStore %28 %30
+               OpReturn
+               OpFunctionEnd

shaders-msl/asm/comp/storage-buffer-basic.asm.comp

@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Codeplay; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpCapability VariablePointers
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpExtension "SPV_KHR_variable_pointers"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %22 "main" %gl_WorkGroupID
+               OpSource OpenCL_C 120
+               OpDecorate %15 SpecId 0
+               ;OpDecorate %16 SpecId 1
+               OpDecorate %17 SpecId 2
+               OpDecorate %_runtimearr_float ArrayStride 4
+               OpMemberDecorate %_struct_4 0 Offset 0
+               OpDecorate %_struct_4 Block
+               OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId
+               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+               OpDecorate %20 DescriptorSet 0
+               OpDecorate %20 Binding 0
+               OpDecorate %21 DescriptorSet 0
+               OpDecorate %21 Binding 1
+      %float = OpTypeFloat 32
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+%_runtimearr_float = OpTypeRuntimeArray %float
+  %_struct_4 = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer__struct_4 = OpTypePointer StorageBuffer %_struct_4
+       %uint = OpTypeInt 32 0
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Private_v3uint = OpTypePointer Private %v3uint
+     %uint_0 = OpConstant %uint 0
+%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input
+         %15 = OpSpecConstant %uint 1
+         %16 = OpConstant %uint 2
+         %17 = OpSpecConstant %uint 3
+%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %15 %16 %17
+         %19 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize
+         %20 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
+         %21 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
+         %22 = OpFunction %void None %8
+         %23 = OpLabel
+         %24 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0
+         %25 = OpLoad %uint %24
+         %26 = OpAccessChain %_ptr_StorageBuffer_float %21 %uint_0 %25
+         %27 = OpLoad %float %26
+         %28 = OpAccessChain %_ptr_StorageBuffer_float %20 %uint_0 %25
+         %29 = OpLoad %float %28
+         %30 = OpFAdd %float %27 %29
+               OpStore %28 %30
+               OpReturn
+               OpFunctionEnd

shaders/asm/comp/storage-buffer-basic.asm.comp

@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Codeplay; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpCapability VariablePointers
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpExtension "SPV_KHR_variable_pointers"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %22 "main" %gl_WorkGroupID
+               OpSource OpenCL_C 120
+               OpDecorate %15 SpecId 0
+               ;OpDecorate %16 SpecId 1
+               OpDecorate %17 SpecId 2
+               OpDecorate %_runtimearr_float ArrayStride 4
+               OpMemberDecorate %_struct_4 0 Offset 0
+               OpDecorate %_struct_4 Block
+               OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId
+               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+               OpDecorate %20 DescriptorSet 0
+               OpDecorate %20 Binding 0
+               OpDecorate %21 DescriptorSet 0
+               OpDecorate %21 Binding 1
+      %float = OpTypeFloat 32
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+%_runtimearr_float = OpTypeRuntimeArray %float
+  %_struct_4 = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer__struct_4 = OpTypePointer StorageBuffer %_struct_4
+       %uint = OpTypeInt 32 0
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Private_v3uint = OpTypePointer Private %v3uint
+     %uint_0 = OpConstant %uint 0
+%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input
+         %15 = OpSpecConstant %uint 1
+         %16 = OpConstant %uint 2
+         %17 = OpSpecConstant %uint 3
+%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %15 %16 %17
+         %19 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize
+         %20 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
+         %21 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
+         %22 = OpFunction %void None %8
+         %23 = OpLabel
+         %24 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0
+         %25 = OpLoad %uint %24
+         %26 = OpAccessChain %_ptr_StorageBuffer_float %21 %uint_0 %25
+         %27 = OpLoad %float %26
+         %28 = OpAccessChain %_ptr_StorageBuffer_float %20 %uint_0 %25
+         %29 = OpLoad %float %28
+         %30 = OpFAdd %float %27 %29
+               OpStore %28 %30
+               OpReturn
+               OpFunctionEnd