From db5d49f04a7c3e81d1d82d76cd8d67b0c7648f86 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Fri, 29 Sep 2017 11:07:11 +0200 Subject: [PATCH] Add SPVASM testing support for HLSL/MSL. --- ...specialization-constant-workgroup.asm.comp | 16 +++++++ .../specialization-constant-workgroup.comp | 16 ------- ...specialization-constant-workgroup.asm.comp | 21 +++++++++ .../specialization-constant-workgroup.comp | 21 --------- ...specialization-constant-workgroup.asm.comp | 47 +++++++++++++++++++ .../specialization-constant-workgroup.comp | 12 ----- ...specialization-constant-workgroup.asm.comp | 47 +++++++++++++++++++ .../specialization-constant-workgroup.comp | 12 ----- test_shaders.py | 24 +++++++--- 9 files changed, 149 insertions(+), 67 deletions(-) create mode 100644 reference/shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp delete mode 100644 reference/shaders-hlsl/comp/specialization-constant-workgroup.comp create mode 100644 reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp delete mode 100644 reference/shaders-msl/comp/specialization-constant-workgroup.comp create mode 100644 shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp delete mode 100644 shaders-hlsl/comp/specialization-constant-workgroup.comp create mode 100644 shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp delete mode 100644 shaders-msl/comp/specialization-constant-workgroup.comp diff --git a/reference/shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp b/reference/shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp new file mode 100644 index 00000000..c8ebaa8b --- /dev/null +++ b/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(); +} diff --git a/reference/shaders-hlsl/comp/specialization-constant-workgroup.comp b/reference/shaders-hlsl/comp/specialization-constant-workgroup.comp deleted file mode 100644 index 61a55659..00000000 --- a/reference/shaders-hlsl/comp/specialization-constant-workgroup.comp +++ /dev/null @@ -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(); -} diff --git a/reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp b/reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp new file mode 100644 index 00000000..5802ddac --- /dev/null +++ b/reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp @@ -0,0 +1,21 @@ +#include +#include + +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; +} + diff --git a/reference/shaders-msl/comp/specialization-constant-workgroup.comp b/reference/shaders-msl/comp/specialization-constant-workgroup.comp deleted file mode 100644 index be827fea..00000000 --- a/reference/shaders-msl/comp/specialization-constant-workgroup.comp +++ /dev/null @@ -1,21 +0,0 @@ -#include -#include - -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; -} - diff --git a/shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp b/shaders-hlsl/asm/comp/specialization-constant-workgroup.asm.comp new file mode 100644 index 00000000..188e3fec --- /dev/null +++ b/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 diff --git a/shaders-hlsl/comp/specialization-constant-workgroup.comp b/shaders-hlsl/comp/specialization-constant-workgroup.comp deleted file mode 100644 index 3beeb8c1..00000000 --- a/shaders-hlsl/comp/specialization-constant-workgroup.comp +++ /dev/null @@ -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; -} diff --git a/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp b/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp new file mode 100644 index 00000000..188e3fec --- /dev/null +++ b/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 diff --git a/shaders-msl/comp/specialization-constant-workgroup.comp b/shaders-msl/comp/specialization-constant-workgroup.comp deleted file mode 100644 index 3beeb8c1..00000000 --- a/shaders-msl/comp/specialization-constant-workgroup.comp +++ /dev/null @@ -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; -} diff --git a/test_shaders.py b/test_shaders.py index f41c8f1a..3574ff24 100755 --- a/test_shaders.py +++ b/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)