Add SPVASM testing support for HLSL/MSL.

reference/shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp

@@ -0,0 +1,16 @@
+const uint _5 = 9u;
+const uint _6 = 4u;
+const uint3 gl_WorkGroupSize = uint3(_5, 20u, _6);
+
+RWByteAddressBuffer _4 : register(u0);
+
+void comp_main()
+{
+    _4.Store(0, asuint(asfloat(_4.Load(0)) + 1.0f));
+}
+
+[numthreads(9, 20, 4)]
+void main()
+{
+    comp_main();
+}

reference/shaders-hlsl/comp/specialization-constant-workgroup.comp

@@ -1,16 +0,0 @@
-const uint _13 = 1u;
-const uint _15 = 1u;
-const uint3 gl_WorkGroupSize = uint3(_13, 10u, _15);
-
-RWByteAddressBuffer _10 : register(u0);
-
-void comp_main()
-{
-    _10.Store3(0, gl_WorkGroupSize);
-}
-
-[numthreads(1, 10, 1)]
-void main()
-{
-    comp_main();
-}

reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+constant uint _5_tmp [[function_constant(10)]];
+constant uint _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 9u;
+constant uint _6_tmp [[function_constant(12)]];
+constant uint _6 = is_function_constant_defined(_6_tmp) ? _6_tmp : 4u;
+constant uint3 gl_WorkGroupSize = uint3(_5, 20u, _6);
+
+struct SSBO
+{
+    float a;
+};
+
+kernel void main0(device SSBO& _4 [[buffer(0)]])
+{
+    _4.a += 1.0;
+}
+

reference/shaders-msl/comp/specialization-constant-workgroup.comp

@@ -1,21 +0,0 @@
-#include <metal_stdlib>
-#include <simd/simd.h>
-
-using namespace metal;
-
-constant uint _13_tmp [[function_constant(3)]];
-constant uint _13 = is_function_constant_defined(_13_tmp) ? _13_tmp : 1u;
-constant uint _15_tmp [[function_constant(20)]];
-constant uint _15 = is_function_constant_defined(_15_tmp) ? _15_tmp : 1u;
-constant uint3 gl_WorkGroupSize = uint3(_13, 10u, _15);
-
-struct SSBO
-{
-    uint3 wg_size;
-};
-
-kernel void main0(device SSBO& _10 [[buffer(0)]])
-{
-    _10.wg_size = gl_WorkGroupSize;
-}
-

shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 1
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 20 1
+               OpSource ESSL 310
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "a"
+               OpName %_ ""
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+               OpDecorate %19 SpecId 10
+               OpDecorate %21 SpecId 12
+               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+       %SSBO = OpTypeStruct %float
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+    %float_1 = OpConstant %float 1
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+       %uint = OpTypeInt 32 0
+         %19 = OpSpecConstant %uint 9
+    %uint_20 = OpConstant %uint 20
+         %21 = OpSpecConstant %uint 4
+     %v3uint = OpTypeVector %uint 3
+%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %19 %uint_20 %21
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_float %_ %int_0
+         %15 = OpLoad %float %14
+         %16 = OpFAdd %float %15 %float_1
+         %17 = OpAccessChain %_ptr_Uniform_float %_ %int_0
+               OpStore %17 %16
+               OpReturn
+               OpFunctionEnd

shaders-hlsl/comp/specialization-constant-workgroup.comp

@@ -1,12 +0,0 @@
-#version 310 es
-layout(local_size_x_id = 3, local_size_y = 10, local_size_z_id = 20) in;
-
-layout(binding = 0) buffer SSBO
-{
-	uvec3 wg_size;
-};
-
-void main()
-{
-	wg_size = gl_WorkGroupSize;
-}

shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 1
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 20 1
+               OpSource ESSL 310
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "a"
+               OpName %_ ""
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+               OpDecorate %19 SpecId 10
+               OpDecorate %21 SpecId 12
+               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+       %SSBO = OpTypeStruct %float
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+    %float_1 = OpConstant %float 1
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+       %uint = OpTypeInt 32 0
+         %19 = OpSpecConstant %uint 9
+    %uint_20 = OpConstant %uint 20
+         %21 = OpSpecConstant %uint 4
+     %v3uint = OpTypeVector %uint 3
+%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %19 %uint_20 %21
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_float %_ %int_0
+         %15 = OpLoad %float %14
+         %16 = OpFAdd %float %15 %float_1
+         %17 = OpAccessChain %_ptr_Uniform_float %_ %int_0
+               OpStore %17 %16
+               OpReturn
+               OpFunctionEnd

shaders-msl/comp/specialization-constant-workgroup.comp

@@ -1,12 +0,0 @@
-#version 310 es
-layout(local_size_x_id = 3, local_size_y = 10, local_size_z_id = 20) in;
-
-layout(binding = 0) buffer SSBO
-{
-	uvec3 wg_size;
-};
-
-void main()
-{
-	wg_size = gl_WorkGroupSize;
-}

test_shaders.py

@@ -86,12 +86,17 @@ def validate_shader_msl(shader):
         print('Error compiling Metal shader: ' + msl_path)
         sys.exit(1)
 
-def cross_compile_msl(shader):
+def cross_compile_msl(shader, spirv):
     spirv_f, spirv_path = tempfile.mkstemp()
     msl_f, msl_path = tempfile.mkstemp(suffix = os.path.basename(shader))
     os.close(spirv_f)
     os.close(msl_f)
-    subprocess.check_call(['glslangValidator', '-V', '-o', spirv_path, shader])
+
+    if spirv:
+        subprocess.check_call(['spirv-as', '-o', spirv_path, shader])
+    else:
+        subprocess.check_call(['glslangValidator', '-V', '-o', spirv_path, shader])
+
     spirv_cross_path = './spirv-cross'
     subprocess.check_call([spirv_cross_path, '--entry', 'main', '--output', msl_path, spirv_path, '--msl'])
     subprocess.check_call(['spirv-val', spirv_path])
@@ -110,12 +115,17 @@ def shader_to_sm(shader):
     else:
         return '50'
 
-def cross_compile_hlsl(shader):
+def cross_compile_hlsl(shader, spirv):
     spirv_f, spirv_path = tempfile.mkstemp()
     hlsl_f, hlsl_path = tempfile.mkstemp(suffix = os.path.basename(shader))
     os.close(spirv_f)
     os.close(hlsl_f)
-    subprocess.check_call(['glslangValidator', '-V', '-o', spirv_path, shader])
+
+    if spirv:
+        subprocess.check_call(['spirv-as', '-o', spirv_path, shader])
+    else:
+        subprocess.check_call(['glslangValidator', '-V', '-o', spirv_path, shader])
+
     spirv_cross_path = './spirv-cross'
 
     sm = shader_to_sm(shader)
@@ -298,7 +308,8 @@ def test_shader(stats, shader, update, keep):
 def test_shader_msl(stats, shader, update, keep):
     joined_path = os.path.join(shader[0], shader[1])
     print('\nTesting MSL shader:', joined_path)
-    spirv, msl = cross_compile_msl(joined_path)
+    is_spirv = shader_is_spirv(shader[1])
+    spirv, msl = cross_compile_msl(joined_path, is_spirv)
     regression_check(shader, msl, update, keep)
     os.remove(spirv)
 
@@ -308,7 +319,8 @@ def test_shader_msl(stats, shader, update, keep):
 def test_shader_hlsl(stats, shader, update, keep):
     joined_path = os.path.join(shader[0], shader[1])
     print('Testing HLSL shader:', joined_path)
-    spirv, msl = cross_compile_hlsl(joined_path)
+    is_spirv = shader_is_spirv(shader[1])
+    spirv, msl = cross_compile_hlsl(joined_path, is_spirv)
     regression_check(shader, msl, update, keep)
     os.remove(spirv)