How Do You Solve a Problem Like MSL?

[billhollings]

@HansKristian-Work raises valid concerns about MSL source being sufficiently different in structure and limitations, when compared to the other source backends. This is exacerbated because real-world runtime client apps such as MoltenVK specialize the way they interface with MSL in order to accomplish non-Metal functionality driven by outside needs.

I'm opening this issue up as a forum to discuss how best to handle this.

1. @HansKristian-Work has suggested that it might be time to separate MSL from SPIRV-Cross. We can discuss how this might work. Is there any need to maintain compatibility with other backends, or should MSL just be separated into a different repo, in a way that it can leverage SPIRV-Cross to date, but the entire repository can evolve only to support MSL, without worrying about how changes might affect other backends?

2. Can we refactor MSL generation to allow specialized functions to be output the handle much of the MSL differences? These functions in turn, might be managed in a SPIRV-branch, or separate repo.

3. With the latest version, Metal is starting to open up to allow runtime functionality to load and compile Metal LLVM IR code, which possibly opens up the idea of compiling from SPIR-V to Metal IR format directly, in a completely separate code conversion library. Doing so would bypass the need to generate and compile MSL source code, which might streamline code generation logic and runtime performance. The missing ingredient to this would be reverse-engineering Metal's specific LLVM IR representation in order to generate that IR code.

[HansKristian-Work]

separate MSL from SPIRV-Cross

That's not in the cards. Plenty of developers rely on SPIRV-Cross outside MoltenVK. They emit MSL and write their own Metal backend for example.

Can we refactor MSL generation to allow specialized functions to be output the handle much of the MSL differences?

I can see an approach where complicated transforms are done through stubbing rather than ramming all that code into SPIRV-Cross (which I will continue to push back on). SPIR-V can express functions which need to link to external code just fine, and there's no reason SPIRV-Cross cannot emit partially complete MSL or whatever and the user can copy in the helper code that fills in the function definitions.

Exactly how it would work in practice is another question, but I'm open to ideas in this area. I think #2204 is a good candidate to use as a testing ground.

Some light massaging of SPIR-V will likely be necessary, but going forward, when you consider how to link multiple shader stages together, whatever solution must be able to do some glue work in the SPIR-V domain.

but the entire repository can evolve only to support MSL, without worrying about how changes might affect other backends

Having a MoltenVK branch that is effectively a shallow fork so that bugfixes can migrate nicely might work, but it's not a great solution long term to deal with the extra maintenance burden, and I consider that a bad outcome. Of course, nothing is stopping anyone from forking SPIRV-Cross.

which possibly opens up the idea of compiling from SPIR-V to Metal IR format directly

If this becomes viable and matures at some point, then SPIRV-Cross has outlived its purpose, which is fine. Main concern is that the IR format is not open, and is subject to breaking changes. Knowing Apple's track record for compat, I'd be wary. From experience, emitting DXIL - which has an open source implementation, but poor documentation - is extremely hairy to emit and the drivers are highly temperamental in what they accept. Given Apple has one driver, I expect it to be extremely tied to their toolchain.

[etang-cw]

As I mentioned on #2204 I also think the best way forward is some sort of hybrid approach where MoltenVK does some codegen and some SPIR-V modifications, while leaving the more involved SPIR-V processing to SPIRV-Cross.

#2204 can get by with relatively small changes, as SPIRV-Cross could potentially just generate the exact same vertex shader it currently does as a static function without the [[xxx]] binding tags on the inputs, and then supply those tags to MoltenVK to apply to its actual shader (where MoltenVK could then replace the one tagged [[stage_in]] with its own vertex loader, or stick it in a kernel for tesselation, etc).

Example code

For this glsl shader:

layout (attribute = 0) in vec3 in0;
layout (attribute = 1) in vec4 in1;
layout (set = 0, binding = 0) uniform texture2d tex;
layout (set = 0, binding = 1, std140) uniform constants { vec4 constant0; };
layout (attribute = 0) out vec4 out0;

void main() {
    // Do stuff
}

SPIRV-Cross would generate

struct constants {
    float4 constant0;
};
struct main0_in {
    float3 in0 [[attribute(0)]];
    float4 in1 [[attribute(1)]];
};
struct main0_out {
    float4 position [[position]];
    float4 color [[user(attr0)]];
};
static main0_out main0(thread const main0_in& in, thread const texture2d<float>& tex, constant constants& _1) {
    // Do stuff
}

Then MoltenVK could generate

vertex main0_out vs_main0(main0_in [[stage_in]], texture2d<float> tex [[texture(0)]], constant constants& _1 [[buffer(0)]]) {
    return main0(main0_in, tex, _1);
}

Or it could generate

kernel void vs_main0(uint2 gid [[thread_position_in_grid]], uint gsize [[grid_size]], main0_in [[stage_in]], texture2d<float> tex [[texture(0)]], constant constants& _1 [[buffer(0)]], device main0_out* mvkOut [[buffer(1)]]) {
    mvkOut[gid.y * gsize.x + gid.x] = main0(main0_in, tex, _1);
}

Or it could generate

struct vertexBuffer0 { /* ... */ };
static main0_in mvkLoadVertex(const device vertexBuffer0& buffer0) { /* ... */ }
vertex main0_out vs_main0(uint vid [[vertex_id]], const device vertexBuffer0* vb0 [[buffer(0)]], texture2d<float> tex [[texture(0)]], constant constants& _1 [[buffer(1)]]) {
    main0_in in = mvkLoadVertex(vertexBuffer0[vid]);
    return main0(in, tex, _1);
}

But for the more complicated things, I think we're going to have to answer a few more questions about who does what. In particular, who scans the SPIR-V to figure out what parts of what builtins/bindings are and aren't used? e.g. to emit a vertex into a geometry stream in a GS, you need to know what components make up a vertex. At the moment, SPIRV-Cross figures that out, which means MoltenVK wouldn't know the content of a struct that it has to declare as the input to its vertex emission function. Maybe we could do something like this:

MoltenVK adds the following to the GS:

%mvkGeometryStream = OpTypeStruct
OpName %mvkGeometryStream "mvkGeometryStream"
OpDecorate %mvkGeometryStream DeclareOnlySPIRVCROSS
%mvkVertexOut = OpTypeStruct
OpName %mvkVertexOut "mvkVertexOut"
OpDecorate %mvkVertexOut StageOutDataSPIRVCROSS
%_ptr_mvkGeometryStream = OpTypePointer Private %mvkGeometryStream
%_ptr_mvkVertexOut = OpTypePointer Function %mvkVertexOut
%_type_mvkEmitVertex = OpTypeFunction %void %_ptr_mvkGeometryStream %_ptr_mvkVertexOut

%mvkEmitVertex = OpFunction %void None %_type_mvkEmitVertex
%_param_geometryStream = OpFunctionParameter %_ptr_mvkGeometryStream
%_param_vertex = OpFunctionParameter %_ptr_mvkVertexIn
OpDecorate %_param_vertex NonWritable
OpFunctionEnd
OpDecorate %mvkEmitVertex LinkageAttributes "mvkEmitVertex" Import

%_ptr_Private_mvkGeometryStream = OpTypePointer Private %mvkGeometryStream
%_local_geometryStream = OpVariable %_ptr_Private_mvkGeometryStream Private

# At the beginning of the GS function
%_local_vertexOut OpVariable Function %_ptr_mvkVertexOut
# Replacing each OpEmitVertex
OpFunctionCall %void %mvkEmitVertex %_local_geometryStream %_local_vertexOut

and SPIRV-Cross populates the struct decorated StageOutData, generating something like

struct mvkVertexOut {
    float4 position [[position]];
    float4 color [[user(attr0)]];
};
struct mvkGeometryEmitter;
static void mvkEmitVertex(thread mvkGeometryEmitter&, const thread mvkVertexOut&);
static void main0(thread mvkGeometryEmitter& _1 /* , ... */) {
    mvkVertexOut out;
    // ...
    mvkEmitVertex(_1, out);
}

Or maybe this job should move to MoltenVK and it should be in charge of rewriting builtins and redirecting them to the [[stage_in]]/output structs?

[alyssarosenzweig]

IMHO, if you're going to be emulating big ticket features like geometry shaders, you need a real compiler, not a thin translator. Doing serious transforms on either SPIRV or MSL is not especially viable.

If I were building MoltenVK, I'd do it as a Mesa driver (with a NIR->MSL or NIR->AIR backend and Mesa's Vulkan runtime), like Microsoft did for Vulkan on D3D12.

[HansKristian-Work]

Doing transforms in SPIR-V domain is quite viable I think, and that's what SPIRV-Tools does.

OpDecorate %mvkEmitVertex LinkageAttributes "mvkEmitVertex" Import

I think first thing to look into would be supporting extern linking and see how that goes. Finding a way to express the weird and wonderful type system of Metal would be the second challenge, but it might be somewhat solvable with UserTypeGOOGLE. I'll look into it soon.