Vulkan features available with Metal 3?

[billhollings]

There have been a few posts asking about Metal 3 support.

Since MoltenVK is a Vulkan implementation, we need to prioritize and develop Vulkan features. So I'd like to recast the question as which Vulkan features can we identify as being made possible or better with the capabilities of Metal 3.

With that in mind, I'm posting this here to crowd-source ideas and comments based on Apple's info on Metal 3.

Feel free to either edit the table here, or add separate comments below for discussion about capabilities and priorities. Some features are fairly straight-forward. Others will require significant development.

Vulkan Feature	Via Metal 3 Capability	Status
VK_EXT_shader_atomic_float VK_EXT_shader_atomic_float2	MSL atomic<float> (but no support for half or double)
VK_EXT_mesh_shader VK_NV_mesh_shader	mesh shaders
Geometry shaders	mesh or compute shaders
Potentially more Vulkan-efficient Tessellation shaders	mesh shaders
VK_KHR_acceleration_structure	ray tracing (previously available but enhanced now)
VK_KHR_ray_tracing_maintenance1	ray tracing (previously available but enhanced now)
VK_KHR_ray_tracing_pipeline	ray tracing (previously available but enhanced now)
VK_KHR_buffer_device_address	MTLBuffer.gpuAddress	Implemented in PR #1619.

[JuliaAgostino]

What's new in Metal 3

MetalFX Upscaling
Render complex scenes in less time per frame with high-performance upscaling and anti-aliasing. Choose a combination of temporal or spatial algorithms to help boost performance.

Fast resource loading
Optimally stream asset data to Metal textures and buffers directly from storage using asynchronous I/O.

Offline shader compilation
The compiler can generate GPU binaries at project build time to eliminate in-app shader compilation, helping games improve performance and reduce load times.

Mesh shaders
This new geometry pipeline replaces vertex shaders with two new shader stages — object and mesh — that enable more flexible culling and LOD selection, and more efficient geometry shading and generation.

Metal backend for PyTorch
The new Metal backend in PyTorch version 1.12 enables high-performance, GPU-accelerated training using MPS Graph and the Metal Performance Shaders primitives.

New Ray Tracing features
The latest advancements in Metal Ray Tracing mean less GPU time is spent building acceleration structures, work like culling can move to the GPU to reduce CPU overhead, and both intersection and shading can be optimized with direct access to primitive data.

[zaptrem]

Hardware-accelerated by the GPU but no dedicated ray-tracing silicon.

[iMonZ]

Hardware-accelerated by the GPU but no dedicated ray-tracing silicon.

Thanks but how does this then work?
Can it be compared to RTX?

[K0bin]

@iMonZ

It's most comparable to running ray tracing on a Nvidia GTX 1070 or 1080. You can run the same content but there's no special hardware to make it fast. On top of that, Mac GPUs (aside from the M1 Max and Ultra) are somewhat slow compared to big fat desktop GPUs, so performance is probably gonna be an issue.

[msitarzewski]

"probably gonna" Nice and scientific. :)

[JuliaAgostino]

We won't know the performance just yet.

[italomandara]

[JuliaAgostino]

Yeah, that would be great!

[billhollings]

Is this enabled by or related to specific Metal3 functionality?

[italomandara]

Is this enabled by or related to specific Metal3 functionality?

@billhollings no it's not it can be implemented in metal2 #1588

[expenses]

https://www.khronos.org/registry/vulkan/specs/1.3-extensions/man/html/VK_EXT_shader_atomic_float2.html should also be implementable - just not for float16 or float64 (afaik)

[billhollings]

I would try to implement VK_NV_mesh_shader if you don't mind.

Definitely! Please do so and submit a PR!

Please follow existing MoltenVK design and syntax forms, so we can keep the MoltenVK codebase consistent.

I've created issue #1618 to give you a place to discuss design options, progress, or to request help or answers.

[Try]

Hey, everyone!
Come here from HLSL side of spirv-cross, so maybe my understanding of msl is not accurate.
Anyway:
In MSL we have void metal::mesh<>::set_vertex(uint I, V v), and similar for primitives. And no get_* kind of method.
But in Vulkan use can define multiple per-verstex and per-primitive outputs. So, cross compiling code, like:

gl_MeshVerticesEXT[gl_LocalInvocationIndex].gl_Position = vec4(vec3(gl_GlobalInvocationID), 1.0);
... lots of code ...
vOut[gl_LocalInvocationIndex] = vec4(vec3(gl_GlobalInvocationID), 2.0);

Has to 1-to-1 translation. Only solution I see right now is to emit all per-vertex and per-primitive stuff as shared arrays, and add copy-pass in the end of shader.
Copy pass is not a great idea, but so far that the only one. Maybe you have good ideas?

[BeastLe9enD]

Yes, you are right afaik. There is no get function, so you cannot read the data once you set it. But gl_MeshPerVertexEXT is writeonly, so its only relevant in cases where the VSOut attributes are not filled in order. The only way I see right now is to analyze the code and see if the dates are all set one after the other. Then we can just fill a structure and call set once.

layout(location = 0) out vec2[] tex_coords;
layout(location = 1) out vec3[] colors;

void main() {
	uint idx = index();
	
	// All vertex attributes are set in order at the same index
	gl_MeshVerticesEXT[idx].gl_Position = position();
	tex_coords[idx] = tex_coord();
	colords[idx] = color();
}

So we could turn this into:

struct VSOut {
	float4 position [[position]];
	float2 tex_coord [[attribute(0)]];
	float3 color [[attribute(1)]];
};

void main(mesh<VSOut, void, NUM_VERTS, NUM_PRIMS, topology::triangle> m) {
        uint idx = index();

	VSOut out;
	out.position = position();
	out.tex_coord = tex_coord();
	out.color = color();
	
	m.set_vertex(idx, out);
}

In cases where the pattern does not match, we need to fallback to shared arrays in my opinion.
I don't think it would be a problem in practice, since almost all mesh shaders actually set all attributes in order.

[Try]

For fast-path something like prefersLocalInvocationVertexOutput/prefersLocalInvocationPrimitiveOutput probably useful to expose.
Basically, if cross-compiler can detect that all OpAccessChain, accessing output arrays only with gl_LocalInvocationIndex, then VSOut can be decelerated as global variable.
Shared array for slow-path :(

[Try]

I've polished my PR with MSL-mesh shader: KhronosGroup/SPIRV-Cross#2074

Implementations is not 100% conformant(see PR), varying's that are "array of structure of arrays" probably wont work. Yet, still it's should be solid enough for initial testing.

Now bad news: in my engine I'm using Metal3(straight) + spirv-cross, and performance test shows that task-mesh-shader culling is not viable for MacOS. Works about 2x slower than brute-force draw-calls.

[spnda]

I would also like to put this table here; it shows which devices actually support Metal 3.

In code, however, its fairly easy to check if its supported nowadays:

if ([device supportsFamily:MTLGPUFamilyMetal3]) {}

[karnnypicture]

As far as I know it didn't support argument buffers either. So it didn't support a lot of features added with Metal 2.0 through Metal 2.4.

does this include A12Z?

[spnda]

MTLGPUFamilyApple5 is This feature set was introduced with Apple A12 GPUs.. So given the A12Z is still a A12 series SoC, I would assume so, yes.

[billhollings]

As far as I know it didn't support argument buffers either.

Argument buffers were introduced in iOS 11, and are supported on all Apple GPUs. However, A13+ supports Tier2 argument buffers.

[expenses]

I've installed macOS Ventura on my 2017 MacBook Pro, and while this official mesh shading sample works, other mesh shader code like @BeastLe9enD's
https://github.com/ProjectKML/metal_3_example does not. This could be either be down to a driver bug or differing levels of support for mesh shaders across hardware.

[kode54]

@expenses This may be relevant, from the change log of Ventura beta 4, which just came out today:

Fixed: When using the new Metal mesh shaders feature, render pipeline state objects (PSOs) created with a mesh shader stage but without a object shader stage can fail to compile or fail to work correctly on some devices. (89836551)

Fixed: When using mesh shaders on Apple Silicon, mesh shader invocations or geometry output can be unexpectedly dropped on large workloads when calling set_threadgroups_per_grid from only a single thread in the threadgroup in a object shader or when calling set_primitive_count from only a single thread in the threadgroup in a mesh shader. (94376391)

Edit: Even more things:

Fixed: When using Metal mesh shaders on a device with Apple silicon, object or mesh shaders producing large amounts of output data can result in incorrect rendering and device instability. (93797539)

Fixed: When using Metal mesh shaders on a device with Apple silicon, vertex buffers set at indices 29 or 30 can become unset after calling any of these methods: drawMeshThreadgroups, drawMeshThreads, drawMeshThreadgroupsWithIndirectBuffer. (93797408)

Fixed: When using Metal mesh shaders on a device with Apple silicon, graphics corruption can occur when emitting a metal::mesh with a large number of primitives and using the viewport_array_index, render_target_array_index, or primitive_culled per-primitive attributes. (93797184)

[nolrinale]

I think it would be wise to focus the development efforts on the rest of features but the raytracing capabilities, since theres no dedicated RTX hardware inside the M1/M2 or A processors GPUs at this point.

Mesh Shaders, Upscaling, Resources Loading sounds like better choices to improve games and 3D apps with the available hardware.

[kode54]

Well, whether or not they're using the ray tracing opcode (in RDNA it's an opcode), the performance difference is definitely noticeable. It's a function that's the difference between a single opcode function call, and several kilobytes of polyfill shader code that does a whole lot more work for the same result. Oh, and in case you're feeling bored, Mesa has a branch, not sure if it got merged, that implements this polyfill so you can test out ray tracing on anything as old as Polaris. The result is that the Vulkan RT demos run like 45 fps for the reflections demo, 26 fps for the lighted models demo, and Quake 2 RTX runs like 2fps, if you can get it to run at all.

[spnda]

If you look into the Metal drivers you'll find something like a raytracing.metallib. That seems to contain shader code that emulates what RDNA, Turing, and Ampere do on the hardware. So none of it is on hardware yet and all of it is implemented in compute.

[mr-mobster]

@kode54 I am a bit puzzled about your mention of "several kilobytes of polyfill shader code"... RDNA2 offers two RT instructions — ray triangle intersection and ray-box-node intersection. Both only need a few dozen instructions at most to be implemented "in software". It is also interesting to note that AMD's solution does not perform as well as Nvidia's because AMD only accelerates the interaction math, while Nvidia seems to be also reordering execution to improve the memory access pattern.

[kode54]

Funny you should say it only requires a few instructions. Look at what Joshua Ashton wrote to implement the ray intersections here: https://gitlab.freedesktop.org/JoshuaAshton/mesa/-/commit/9e7733d4d0a871f98f23350a6bc2d0e5200e7346

[kode54]

Okay, maybe I exaggerated, but that is a few dozen instructions. I had no idea that it was so slow on AMD's GPUs if that lowering has such an overhead. Guess I'll buy a 3090 in the future.

[mr-mobster]

Metal 3 introduces a revamp of the binding model. MTLArgumentEncoder is now deprecated and one populates bindings by simply writing 64bit GPU addresses into the bindings buffer (argument buffer). These changes should make it much easier to emulate Vulkan binding model efficiently, as you don't have to fight your way around the typed argument encoders anymore, you just memcpy() the GPU address to where you need it in the table.

Another note on supported limits. It has been pointed out that Metal Tier 2 limits of 500,000 buffers and textures is lower than some of the advanced versions of Vulkan and DirectX 12, at least on certain hardware. Upon clarification with Apple engineers, it seems that this is a poor wording in the docs. This limit is the number of resources you can realistically use in a single shader while still expecting reasonable performance. This is not the limit of binding points/descriptors you can have. Since Metal has no concept of special descriptor pools — resource descriptors reside in regular GPU buffer memory, it seems that the upper limit on the number of descriptors is simply the amount of available buffer memory. I had no issues creating a test application that declares, binds and actually uses up to five million resource descriptors.

[billhollings]

Metal 3 introduces a revamp of the binding model.

Since Metal has no concept of special descriptor pools — resource descriptors reside in regular GPU buffer memory, it seems that the upper limit on the number of descriptors is simply the amount of available buffer memory. I had no issues creating a test application that declares, binds and actually uses up to five million resource descriptors.

Thanks for the research, and design input.

The new binding model is good news. It should clean up how MoltenVK can map descriptor indexing to argument buffers, and hopefully allow us to finally support descriptor indexing on iOS.

[SpiritHi]

Does Metal 3 do away with the tier 2 argument buffer limitation?

[Try]

bindless resources

Can you go into detail about this? I was thinking that sm6.6 maps very well descriptor-buffer extension and the extension can be built on top of argument buffers, isn't it right? Any blockers?

[K0bin]

It does at first glance but the way Metal combines automatic barrier tracking with residency is a problem.

#214 (comment)

[billhollings]

descriptor-buffer extension and the extension can be built on top of argument buffers

There is an ongoing attempt to do just this, but resource tracking is currently an unsolved issue.

[K0bin]

Given that you can write descriptor buffers on the GPU in uncached memory or VRAM & PARTIALLY_BOUND applies, I don't see any way to implement that.

[SwampyApple]

No, it does not. DXVK is doable but VKD3D-Proton is blocked by support for bindless resources.

It's probably worth it to address this blocker over at the Apple developer forums. They might be willing to add some features in a future Metal 3.x version. The people over at CodeWeavers got in touch with Apple over some Rosetta 2 bugs which got fixed. They've also added a raw image mode for the Asahi Linux development, instead of needing Mach-O kernel files.

I'd offer to do this myself, but unfortunately I don't have the technical knowledge to communicate the requirements.

[rcoreilly]

I could really use the VK_EXT_shader_atomic_float extension. How hard would it be to get that going? I could try to do it myself with some minimal pointers to what might be required, if it is more-or-less a matter of extending atomic to accept and putting the right glue somewhere for atomicAdd to do the right thing.. I just did this for the HLSL parser in glslc and it seems to work..

[spnda]

Have a look at MoltenVK/MoltenVK/Layers/MVKExtensions.def. Then in MVKDevice.mm search for the code that enables properties and features through the vkGetPhysicalDeviceProperties calls. That should be enough to advertise support. Also look at other commits that add new extensions and read what they do to get an idea how the system works.

[rcoreilly]

Well, after spending way too much time rebooting after each build of the .dylib (see comments at end), it is now working! I submitted a PR to SPIRV-Cross: KhronosGroup/SPIRV-Cross#2105 -- otherwise it seems to work properly with just this one change to MoltenVK itself:

diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.def b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
index 17ed1f25..504567c5 100644
--- a/MoltenVK/MoltenVK/Layers/MVKExtensions.def
+++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
@@ -113,6 +113,7 @@ MVK_EXTENSION(EXT_scalar_block_layout,             EXT_SCALAR_BLOCK_LAYOUT,
 MVK_EXTENSION(EXT_separate_stencil_usage,          EXT_SEPARATE_STENCIL_USAGE,           DEVICE,   10.11,  8.0)
 MVK_EXTENSION(EXT_shader_stencil_export,           EXT_SHADER_STENCIL_EXPORT,            DEVICE,   10.14, 12.0)
 MVK_EXTENSION(EXT_shader_viewport_index_layer,     EXT_SHADER_VIEWPORT_INDEX_LAYER,      DEVICE,   10.11,  8.0)
+MVK_EXTENSION(EXT_shader_atomic_float,             EXT_SHADER_ATOMIC_FLOAT,              DEVICE,   10.11,  8.0)
 MVK_EXTENSION(EXT_subgroup_size_control,           EXT_SUBGROUP_SIZE_CONTROL,            DEVICE,   10.14, 13.0)
 MVK_EXTENSION(EXT_swapchain_colorspace,            EXT_SWAPCHAIN_COLOR_SPACE,            INSTANCE, 10.11,  9.0)
 MVK_EXTENSION(EXT_texel_buffer_alignment,          EXT_TEXEL_BUFFER_ALIGNMENT,           DEVICE,   10.13, 11.0)

The logic for requiring the extension to be enabled and dealing with the sub-options there seems to have already been built into SPIRV-Cross. I'll submit a PR for this one change along with additions to the UserGuide if someone provides the secret sauce to avoid rebooting.

The other critical thing that took a long time to figure out was where to set the metal compiler version to 3.0 -- finally found out that I needed to upgrade to Ventura (13.x) -- was on 12.x. I did get it compiling fine separately using spirv-cross and the metal compiler on 12.3 or whatever I was on:

$ spirv-cross gpu_sendspike.spv --metal --msl-version 30000 >ss.metal
$ xcrun -sdk macosx metal ss.metal -std=metal3.0

so perhaps now the defaults can be changed here:

diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
index 5ec47c4c..42738607 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
@@ -1775,7 +1775,7 @@
        }
 #endif
 #if MVK_XCODE_14
-       if ( mvkOSVersionIsAtLeast(13.0) ) {
+       if ( mvkOSVersionIsAtLeast(12.0) ) {
                _metalFeatures.mslVersionEnum = MTLLanguageVersion3_0;
        }
 #endif

when I tried this, it gave me warnings about the OS version so somewhere else I didn't track down seems to also have this constraint.

Finally, it might be useful to have an explicit recipe in https://github.com/KhronosGroup/MoltenVK/blob/main/Docs/MoltenVK_Runtime_UserGuide.md for how to build and install a modified .dylib into the /usr/local/lib location used by the Vulkan SDK, as I would assume that is perhaps the most likely way in which people will get started on trying to modify MoltenVK. My strategy of copying the .dylib from the Package made by make macos:

$ sudo cp Package/Release/MoltenVK/dylib/macOS/libMoltenVK.dylib /usr/local/lib

resulted in inexplicable "Killed: 9" errors that left me scratching my head for a while before seeing the suggestion to reboot. Seriously!? And I still can't get around that issue despite attempts to turn off gatekeeper and sign the .dylib. Also, after figuring that out, the next hurdle was this MoltenVK_icd.json file in /usr/local/share/vulkan/icd.d/MoltenVK_icd.json that I only got working once I edited it to specify the full path:

{
    "file_format_version" : "1.0.0",
    "ICD": {
        "library_path": "/usr/local/lib/libMoltenVK.dylib",
        "api_version" : "1.2.0",
        "is_portability_driver" : true
    }
}

So much time was spent battling these issues relative to the almost trivial changes needed to get this to work!

[cdavis5e]

MVK_EXTENSION(EXT_shader_atomic_float,             EXT_SHADER_ATOMIC_FLOAT,              DEVICE,   10.11,  8.0)

This needs to be:

MVK_EXTENSION(EXT_shader_atomic_float,             EXT_SHADER_ATOMIC_FLOAT,              DEVICE,   13.0 , 16.0)

because the required support in MSL wasn't added until Metal 3, which was introduced in macOS 13 and iOS 16.

My strategy of copying the .dylib from the Package made by make macos:

$ sudo cp Package/Release/MoltenVK/dylib/macOS/libMoltenVK.dylib /usr/local/lib

resulted in inexplicable "Killed: 9" errors that left me scratching my head for a while before seeing the suggestion to reboot.

That's because the kernel caches code signatures by vnode. (The kernel will terminate with extreme prejudice any process that loads an image with the wrong code signature.) If instead you remove the old version with:

$ sudo rm /usr/local/lib/libMoltenVK.dylib

before installing the new one, this will force the kernel to make a new vnode, and then that in turn will force it to reload the code signature. Then you won't have to restart in order to try out a new version.

[rcoreilly]

Wow thanks for the tip! I can't believe I didn't come across that in my googling -- won't soon forget it (I hope!). I'll go ahead and update the versions accordingly and submit a PR..

[spnda]

otherwise it seems to work properly with just this one change to MoltenVK itself:

Well yes, but that doesn't cover extension properties and features that will currently all be disabled. You will need to disable all the float64 properties anyway as Metal doesn't support that, but enable all the float32 properties that are applicable. Unless you did that and just didn't mention it? :)

[onlyalexicon]

How is mesh shading support going along for moltenVK?