Vulkan: missing validation layers: VK_LAYER_LUNARG_standard_validation

[vili-1]

When I run the following command under MoltenVK (with option -- Disable validation layers) for the amber-file below, I get:

amber -d -t spv1.3 -v 1.1 test_0_2417849846392727996.amber

[mvk-error] VK_ERROR_INITIALIZATION_FAILED: Shader library compile failed (Error code 3):
Compilation failed: 

program_source:28:13: error: use of undeclared identifier '_50'
    _13._m0[_50] = 8u;
            ^
.
[mvk-error] VK_ERROR_INVALID_SHADER_NV: Compute shader function could not be compiled into pipeline. See previous logged error.
test_0_2417849846392727996.amber: Vulkan::Calling vkCreateComputePipelines Fail

Without passing the -d option, then no error, however when I run it I get:

Vulkan: missing validation layers:
		VK_LAYER_LUNARG_standard_validation,
		
Sample: extensions of validation layers are not supported

Doing some digging I found at some forums that:

VK_LAYER_LUNARG_standard_validation is no longer available, it has been replaced by VK_LAYER_KHRONOS_validation

However couldn't find some more official statement about this. Running vulkaninfo I can see that VK_LAYER_KHRONOS_validation is in. @Jack-Clark did you get any output like this? @afd does this ring any bells?

#!amber

SHADER compute compute_shader SPIRV-ASM

; Follow the path:
; 8 -> <9> -> <13> -> 11 -> <14> -> edge_0 -> 16 -> <15> -> 19 -> <17> -> <9> -> 12
;
; 5 CFG nodes have OpBranchConditional or OpSwitch as their terminators (denoted <n>): 9, 13, 14, 15 and 17.
;
; To follow this path, we need to make these decisions each time we reach 9, 13, 14, 15 or 17.
; This path was generated with the seed 2417849846392727996 and has length 11.
;
; We equip the shader with 5+1 storage buffers:
; - An input storage buffer with the directions for each node 9, 13, 14, 15 or 17
; - An output storage buffer that records the blocks that are executed

; SPIR-V
; Version: 1.3
; Generator: Khronos Glslang Reference Front End; 8
; Bound: 15
; Schema: 0

               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %7 "main"
               OpExecutionMode %7 LocalSize 1 1 1
               
               ; Below, we declare various types and variables for storage buffers.
               ; These decorations tell SPIR-V that the types and variables relate to storage buffers


               OpDecorate %size_1_struct_type BufferBlock
               OpMemberDecorate %size_1_struct_type 0 Offset 0
               OpDecorate %size_1_array_type ArrayStride 4

               OpDecorate %size_2_struct_type BufferBlock
               OpMemberDecorate %size_2_struct_type 0 Offset 0
               OpDecorate %size_2_array_type ArrayStride 4

               OpDecorate %output_struct_type BufferBlock
               OpMemberDecorate %output_struct_type 0 Offset 0
               OpDecorate %output_array_type ArrayStride 4

               OpDecorate %directions_9_variable DescriptorSet 0
               OpDecorate %directions_9_variable Binding 0

               OpDecorate %directions_13_variable DescriptorSet 0
               OpDecorate %directions_13_variable Binding 1

               OpDecorate %directions_14_variable DescriptorSet 0
               OpDecorate %directions_14_variable Binding 2

               OpDecorate %directions_15_variable DescriptorSet 0
               OpDecorate %directions_15_variable Binding 3

               OpDecorate %directions_17_variable DescriptorSet 0
               OpDecorate %directions_17_variable Binding 4

               OpDecorate %output_variable DescriptorSet 0
               OpDecorate %output_variable Binding 5


          %1 = OpTypeVoid
          %2 = OpTypeFunction %1
          %3 = OpTypeBool
          %4 = OpTypeInt 32 0
          %5 = OpConstantTrue %3
          %6 = OpConstant %4 0
          

               %constant_0 = OpConstant %4 0
               %constant_1 = OpConstant %4 1
               %constant_2 = OpConstant %4 2
               %constant_8 = OpConstant %4 8
               %constant_9 = OpConstant %4 9
               %constant_10 = OpConstant %4 10
               %constant_11 = OpConstant %4 11
               %constant_12 = OpConstant %4 12
               %constant_13 = OpConstant %4 13
               %constant_14 = OpConstant %4 14
               %constant_15 = OpConstant %4 15
               %constant_16 = OpConstant %4 16
               %constant_17 = OpConstant %4 17
               %constant_18 = OpConstant %4 18
               %constant_19 = OpConstant %4 19
               %constant_20 = OpConstant %4 20


               ; Declaration of storage buffers for the 5 directions and the output
               

               %size_1_array_type = OpTypeArray %4 %constant_1
               %size_1_struct_type = OpTypeStruct %size_1_array_type
               %size_1_pointer_type = OpTypePointer Uniform %size_1_struct_type
               %directions_17_variable = OpVariable %size_1_pointer_type Uniform
               %directions_13_variable = OpVariable %size_1_pointer_type Uniform
               %directions_14_variable = OpVariable %size_1_pointer_type Uniform
               %directions_15_variable = OpVariable %size_1_pointer_type Uniform

               %size_2_array_type = OpTypeArray %4 %constant_2
               %size_2_struct_type = OpTypeStruct %size_2_array_type
               %size_2_pointer_type = OpTypePointer Uniform %size_2_struct_type
               %directions_9_variable = OpVariable %size_2_pointer_type Uniform

               %output_array_type = OpTypeArray %4 %constant_11
               %output_struct_type = OpTypeStruct %output_array_type
               %output_pointer_type = OpTypePointer Uniform %output_struct_type
               %output_variable = OpVariable %output_pointer_type Uniform

               ; Pointer type for declaring local variables of int type
               %local_int_ptr = OpTypePointer Function %4

               ; Pointer type for integer data in a storage buffer
               %storage_buffer_int_ptr = OpTypePointer Uniform %4


          %7 = OpFunction %1 None %2

          %8 = OpLabel ; validCFG/StructurallyReachableBlock$8
               %output_index = OpVariable %local_int_ptr Function %constant_0
               %directions_9_index = OpVariable %local_int_ptr Function %constant_0
               %directions_13_index = OpVariable %local_int_ptr Function %constant_0
               %directions_14_index = OpVariable %local_int_ptr Function %constant_0
               %directions_15_index = OpVariable %local_int_ptr Function %constant_0
               %directions_17_index = OpVariable %local_int_ptr Function %constant_0


   %temp_8_0 = OpLoad %4 %output_index
   %temp_8_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_8_0
               OpStore %temp_8_1 %constant_8
   %temp_8_2 = OpIAdd %4 %temp_8_0 %constant_1
               OpStore %output_index %temp_8_2
               OpBranch %9


          %9 = OpLabel ; validCFG/LoopHeader$0
   %temp_9_0 = OpLoad %4 %output_index
   %temp_9_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_9_0
               OpStore %temp_9_1 %constant_9
   %temp_9_2 = OpIAdd %4 %temp_9_0 %constant_1
               OpStore %output_index %temp_9_2
   %temp_9_3 = OpLoad %4 %directions_9_index
   %temp_9_4 = OpAccessChain %storage_buffer_int_ptr %directions_9_variable %constant_0 %temp_9_3
   %temp_9_5 = OpLoad %4 %temp_9_4
   %temp_9_6 = OpIEqual %3 %temp_9_5 %constant_1
   %temp_9_7 = OpIAdd %4 %temp_9_3 %constant_1
               OpStore %directions_9_index %temp_9_7
               OpLoopMerge %10 %11 None
               OpBranchConditional %temp_9_6 %12 %13


         %12 = OpLabel ; validCFG/StructurallyReachableBlock$4
  %temp_12_0 = OpLoad %4 %output_index
  %temp_12_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_12_0
               OpStore %temp_12_1 %constant_12
  %temp_12_2 = OpIAdd %4 %temp_12_0 %constant_1
               OpStore %output_index %temp_12_2
               OpReturn


         %13 = OpLabel ; validCFG/StructurallyReachableBlock$3
  %temp_13_0 = OpLoad %4 %output_index
  %temp_13_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_13_0
               OpStore %temp_13_1 %constant_13
  %temp_13_2 = OpIAdd %4 %temp_13_0 %constant_1
               OpStore %output_index %temp_13_2
  %temp_13_3 = OpLoad %4 %directions_13_index
  %temp_13_4 = OpAccessChain %storage_buffer_int_ptr %directions_13_variable %constant_0 %temp_13_3
  %temp_13_5 = OpLoad %4 %temp_13_4
  %temp_13_6 = OpIEqual %3 %temp_13_5 %constant_1
  %temp_13_7 = OpIAdd %4 %temp_13_3 %constant_1
               OpStore %directions_13_index %temp_13_7
               OpBranchConditional %temp_13_6 %11 %10


         %11 = OpLabel ; validCFG/StructurallyReachableBlock$2
  %temp_11_0 = OpLoad %4 %output_index
  %temp_11_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_11_0
               OpStore %temp_11_1 %constant_11
  %temp_11_2 = OpIAdd %4 %temp_11_0 %constant_1
               OpStore %output_index %temp_11_2
               OpBranch %14


         %14 = OpLabel ; validCFG/SelectionHeader$1
  %temp_14_0 = OpLoad %4 %output_index
  %temp_14_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_14_0
               OpStore %temp_14_1 %constant_14
  %temp_14_2 = OpIAdd %4 %temp_14_0 %constant_1
               OpStore %output_index %temp_14_2
  %temp_14_3 = OpLoad %4 %directions_14_index
  %temp_14_4 = OpAccessChain %storage_buffer_int_ptr %directions_14_variable %constant_0 %temp_14_3
  %temp_14_5 = OpLoad %4 %temp_14_4
  %temp_14_7 = OpIAdd %4 %temp_14_3 %constant_1
               OpStore %directions_14_index %temp_14_7
               OpSelectionMerge %15 None
               OpSwitch %temp_14_5 %16


         %16 = OpLabel ; validCFG/StructurallyReachableBlock$1
  %temp_16_0 = OpLoad %4 %output_index
  %temp_16_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_16_0
               OpStore %temp_16_1 %constant_16
  %temp_16_2 = OpIAdd %4 %temp_16_0 %constant_1
               OpStore %output_index %temp_16_2
               OpBranch %15


         %15 = OpLabel ; validCFG/SelectionHeader$0
  %temp_15_0 = OpLoad %4 %output_index
  %temp_15_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_15_0
               OpStore %temp_15_1 %constant_15
  %temp_15_2 = OpIAdd %4 %temp_15_0 %constant_1
               OpStore %output_index %temp_15_2
  %temp_15_3 = OpLoad %4 %directions_15_index
  %temp_15_4 = OpAccessChain %storage_buffer_int_ptr %directions_15_variable %constant_0 %temp_15_3
  %temp_15_5 = OpLoad %4 %temp_15_4
  %temp_15_6 = OpIEqual %3 %temp_15_5 %constant_1
  %temp_15_7 = OpIAdd %4 %temp_15_3 %constant_1
               OpStore %directions_15_index %temp_15_7
               OpSelectionMerge %17 None
               OpBranchConditional %temp_15_6 %18 %19


         %18 = OpLabel ; validCFG/StructurallyReachableBlock$0
               OpBranch %17


         %19 = OpLabel ; validCFG/StructurallyReachableBlock$7
  %temp_19_0 = OpLoad %4 %output_index
  %temp_19_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_19_0
               OpStore %temp_19_1 %constant_19
  %temp_19_2 = OpIAdd %4 %temp_19_0 %constant_1
               OpStore %output_index %temp_19_2
               OpBranch %17


         %17 = OpLabel ; validCFG/StructurallyReachableBlock$6
  %temp_17_0 = OpLoad %4 %output_index
  %temp_17_1 = OpAccessChain %storage_buffer_int_ptr %output_variable %constant_0 %temp_17_0
               OpStore %temp_17_1 %constant_17
  %temp_17_2 = OpIAdd %4 %temp_17_0 %constant_1
               OpStore %output_index %temp_17_2
  %temp_17_3 = OpLoad %4 %directions_17_index
  %temp_17_4 = OpAccessChain %storage_buffer_int_ptr %directions_17_variable %constant_0 %temp_17_3
  %temp_17_5 = OpLoad %4 %temp_17_4
  %temp_17_6 = OpIEqual %3 %temp_17_5 %constant_1
  %temp_17_7 = OpIAdd %4 %temp_17_3 %constant_1
               OpStore %directions_17_index %temp_17_7
               OpBranchConditional %temp_17_6 %9 %10


         %10 = OpLabel ; validCFG/StructurallyReachableBlock$5
               OpBranch %20


         %20 = OpLabel ; validCFG/Block$0
               OpReturn

               OpFunctionEnd

 END

 BUFFER directions_9 DATA_TYPE uint32 STD430 DATA 0 1 END
 BUFFER directions_13 DATA_TYPE uint32 STD430 DATA 1 END
 BUFFER directions_14 DATA_TYPE uint32 STD430 DATA 0 END
 BUFFER directions_15 DATA_TYPE uint32 STD430 DATA 0 END
 BUFFER directions_17 DATA_TYPE uint32 STD430 DATA 1 END

 BUFFER output DATA_TYPE uint32 STD430 SIZE 11 FILL 0

 PIPELINE compute pipeline
   ATTACH compute_shader

   BIND BUFFER directions_9 AS storage DESCRIPTOR_SET 0 BINDING 0
   BIND BUFFER directions_13 AS storage DESCRIPTOR_SET 0 BINDING 1
   BIND BUFFER directions_14 AS storage DESCRIPTOR_SET 0 BINDING 2
   BIND BUFFER directions_15 AS storage DESCRIPTOR_SET 0 BINDING 3
   BIND BUFFER directions_17 AS storage DESCRIPTOR_SET 0 BINDING 4

   BIND BUFFER output AS storage DESCRIPTOR_SET 0 BINDING 5
 END

 RUN pipeline 1 1 1

 EXPECT directions_9 IDX 0 EQ 0 1
 EXPECT directions_13 IDX 0 EQ 1
 EXPECT directions_14 IDX 0 EQ 0
 EXPECT directions_15 IDX 0 EQ 0
 EXPECT directions_17 IDX 0 EQ 1
 EXPECT output IDX 0 EQ 8 9 13 11 14 16 15 19 17 9 12

[afd]

Two separate things going on here: (1) you don't have validation layers set up, and (2) it looks like there may be a spirv-cross bug.

Let me address (2) first.

MoltenVk works by translating each SPIR-V shader into Metal Shading Language, Apple's shading language. It does this using the spirv-cross tool. The resulting Metal Shading Langauge (MSL) code then gets compiled by Apple's downstream compiler.

From this output:

program_source:28:13: error: use of undeclared identifier '_50'
    _13._m0[_50] = 8u;

it looks like the Metal compiler is rejecting the program that spirv-cross has produced, which is probably due to spirv-cross having generated an incorrect program.

We should do some direct testing of spirv-cross to check that it produces valid code - not just for MSL, but for its other back-ends (GLSL and HLSL).

• It looks like Apple provides a command-line tool, metal, that can take a MSL file and try to compile it - see here.
• Similarly, Khronos provide glslangValidator that can be used to validate GLSL.
• Finally, Microsoft provide DXC which can b compiled to run on [Linux/MacOS].(https://github.com/microsoft/DirectXShaderCompiler/blob/master/docs/DxcOnUnix.rst), which can validate HLSL code.

@vili-1 and @Jack-Clark would you be able to work together to set up some tooling for testing spirv-cross against these validators? Something like the following:

• Given an Amber file, rip out the SPIR-V and assemble it
• Use spirv-cross to compile the SPIR-V into MSL
• Use the metal tool to validate the MSL
• Repeat using spirv-cross to compile to GLSL and HLSL, with associated validators

Similarly, naga can convert from SPIR-V into various formats - it would be great to use our fleshed examples to test that naga is generating valid code.

[afd]

Regarding (1): install the Lunar G Vulkan SDK somewhere. I'm not sure what you need to do on Mac to get validation layers working, but on Linux you source a .sh file and it sets up all the right environment variables - I'm sure there will be decent docs for Mac.

[vili-1]

Indeed, there's an environment setup .sh file which I've sourced into the shell, but no luck. There is no VK_LAYER_LUNARG_standard_validation in this file. Will keep digging..

[vili-1]

I compiled the amber above into MSL (see below), and used the metal tool to validate it - but the following error is generated (same error when running amber).

which makes sense because _62 is initialised one line after. I corrected the file and it works :-) the metal tool now compiles .metal file into .air file.

The original metal file:

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct _21
{
    uint _m0[1];
};

struct _23
{
    uint _m0[2];
};

struct _25
{
    uint _m0[11];
};

kernel void main0(device _21& _31 [[buffer(0)]], device _21& _28 [[buffer(1)]], device _21& _29 [[buffer(2)]], device _21& _30 [[buffer(3)]], device _23& _27 [[buffer(4)]], device _25& _32 [[buffer(5)]])
{
    uint _55 = 0u;
    uint _56 = 0u;
    uint _57 = 0u;
    uint _58 = 0u;
    uint _59 = 0u;
    _32._m0[_62] = 8u;
    uint _62 = _62 + 1u;
    uint _54 = _62;
    for (;;)
    {
        _32._m0[_54] = 9u;
        _54++;
        uint _66 = _55;
        _55 = _66 + 1u;
        if (_27._m0[_66] == 1u)
        {
            _32._m0[_54] = 12u;
            _54++;
            return;
        }
        else
        {
            _32._m0[_54] = 13u;
            _54++;
            uint _77 = _56;
            _56 = _77 + 1u;
            if (_28._m0[_77] == 1u)
            {
                _32._m0[_54] = 11u;
                _54++;
                _32._m0[_54] = 14u;
                _54++;
                _57++;
                do
                {
                    _32._m0[_54] = 16u;
                    _54++;
                    break;
                } while(false);
                _32._m0[_54] = 15u;
                _54++;
                uint _98 = _58;
                _58 = _98 + 1u;
                if (_30._m0[_98] == 1u)
                {
                }
                else
                {
                    _32._m0[_54] = 19u;
                    _54++;
                }
                _32._m0[_54] = 17u;
                _54++;
                uint _109 = _59;
                _59 = _109 + 1u;
                if (_31._m0[_109] == 1u)
                {
                    continue;
                }
                else
                {
                    break;
                }
            }
            else
            {
                break;
            }
        }
    }
}

The corrected metal file:

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct _21
{
    uint _m0[1];
};

struct _23
{
    uint _m0[2];
};

struct _25
{
    uint _m0[11];
};

kernel void main0(device _21& _31 [[buffer(0)]], device _21& _28 [[buffer(1)]], device _21& _29 [[buffer(2)]], device _21& _30 [[buffer(3)]], device _23& _27 [[buffer(4)]], device _25& _32 [[buffer(5)]])
{
    uint _55 = 0u;
    uint _56 = 0u;
    uint _57 = 0u;
    uint _58 = 0u;
    uint _59 = 0u;
    uint _62 = 0u;
    _62 = _62 + 1u;
    _32._m0[_62] = 8u;
    uint _54 = _62;
    for (;;)
    {
        _32._m0[_54] = 9u;
        _54++;
        uint _66 = _55;
        _55 = _66 + 1u;
        if (_27._m0[_66] == 1u)
        {
            _32._m0[_54] = 12u;
            _54++;
            return;
        }
        else
        {
            _32._m0[_54] = 13u;
            _54++;
            uint _77 = _56;
            _56 = _77 + 1u;
            if (_28._m0[_77] == 1u)
            {
                _32._m0[_54] = 11u;
                _54++;
                _32._m0[_54] = 14u;
                _54++;
                _57++;
                do
                {
                    _32._m0[_54] = 16u;
                    _54++;
                    break;
                } while(false);
                _32._m0[_54] = 15u;
                _54++;
                uint _98 = _58;
                _58 = _98 + 1u;
                if (_30._m0[_98] == 1u)
                {
                }
                else
                {
                    _32._m0[_54] = 19u;
                    _54++;
                }
                _32._m0[_54] = 17u;
                _54++;
                uint _109 = _59;
                _59 = _109 + 1u;
                if (_31._m0[_109] == 1u)
                {
                    continue;
                }
                else
                {
                    break;
                }
            }
            else
            {
                break;
            }
        }
    }
}

@afd @Jack-Clark @johnwickerson FYI

[afd]

Thanks. I assume you used spirv-cross to do the translation to MSL. Can you check whether this reproduces with the latest build of spirv-cross? (I believe it's straightforward to build using CMake.) If it does, can you file a spirv-cross issue?

[vili-1]

Yes, spirv-cross. Will check for newer versions..

[vili-1]

Same problem with the latest version. @afd Shall I file the spirv-cross issue at https://github.com/KhronosGroup/SPIRV-Cross or GitLab?

[afd]

Please file it at https://github.com/KhronosGroup/SPIRV-Cross