Open vili-1 opened 2 years ago
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).
metal
, that can take a MSL file and try to compile it - see here.glslangValidator
that can be used to validate GLSL.@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:
metal
tool to validate the MSLSimilarly, 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.
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.
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..
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
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?
Yes, spirv-cross. Will check for newer versions..
Same problem with the latest version. @afd Shall I file the spirv-cross issue at https://github.com/KhronosGroup/SPIRV-Cross or GitLab?
Please file it at https://github.com/KhronosGroup/SPIRV-Cross
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
Without passing the
-d
option, then no error, however when I run it I get:Doing some digging I found at some forums that:
However couldn't find some more official statement about this. Running
vulkaninfo
I can see thatVK_LAYER_KHRONOS_validation
is in. @Jack-Clark did you get any output like this? @afd does this ring any bells?