I'm interested in moving away from Xcode and manually compiling Metal shaders in a project for a mixed-language application.
I have no idea how to do this, though. Xcode hides the details of shader compilation and subsequent loading into the application at runtime (you just call device.newDefaultLibrary()). Is this even possible, or will I have to use runtime shader compilation for my purposes?
Generally, you have three ways to load a shader library in Metal:
Use runtime shader compilation from shader source code via the MTLDevice newLibraryWithSource:options:error: or newLibraryWithSource:options:completionHandler: methods. Although purists may shy away from runtime compilation, this option has minimal practical overhead, and so is completely viable. Your primary practical reason for avoiding this option might be to avoid making your shader source code available as part of your application, to protect your IP.
Load compiled binary libraries using the MTLLibrary newLibraryWithFile:error: or newLibraryWithData:error: methods. Follow the instructions in Using Command Line Utilities to Build a Library to create these individual binary libraries at build time.
Let Xcode compile your various *.metal files at build time into the default library available through MTLDevice newDefaultLibrary.
Here's actual code that creates vertex and fragment programs from a string; using it allows you to compile shaders at runtime (shown in the code following the shader string method).
To eliminate the need for the use of escape sequences (e.g., n...), I use the STRINGIFY macro. To workaround its limitation on the use of double quotes, I wrote a block that takes an array of header file names and creates import statements from them. It then inserts them into the shader at the appropriate place; I did the same for include statements. It simplifies and expedites the insertion of what are sometimes rather lengthy lists.
Incorporating this code will not only allow you to select a particular shader to use based on localization, but, if necessary, could also be used to update your app's shaders without having to update the app. You would simply create and ship a text file containing your shader code, which your app could be preprogrammed to reference as the shader source.
#if !defined(_STRINGIFY)
#define __STRINGIFY( _x ) # _x
#define _STRINGIFY( _x ) __STRINGIFY( _x )
#endif
typedef NSString *(^StringifyArrayOfIncludes)(NSArray <NSString *> *includes);
static NSString *(^stringifyHeaderFileNamesArray)(NSArray <NSString *> *) = ^(NSArray <NSString *> *includes) {
NSMutableString *importStatements = [NSMutableString new];
[includes enumerateObjectsUsingBlock:^(NSString * _Nonnull include, NSUInteger idx, BOOL * _Nonnull stop) {
[importStatements appendString:@"#include <"];
[importStatements appendString:include];
[importStatements appendString:@">\n"];
}];
return [NSString new];
};
typedef NSString *(^StringifyArrayOfHeaderFileNames)(NSArray <NSString *> *headerFileNames);
static NSString *(^stringifyIncludesArray)(NSArray *) = ^(NSArray *headerFileNames) {
NSMutableString *importStatements = [NSMutableString new];
[headerFileNames enumerateObjectsUsingBlock:^(NSString * _Nonnull headerFileName, NSUInteger idx, BOOL * _Nonnull stop) {
[importStatements appendString:@"#import "];
[importStatements appendString:@_STRINGIFY("")];
[importStatements appendString:headerFileName];
[importStatements appendString:@_STRINGIFY("")];
[importStatements appendString:@"\n"];
}];
return [NSString new];
};
- (NSString *)shader
{
NSString *includes = stringifyIncludesArray(@[@"metal_stdlib", @"simd/simd.h"]);
NSString *imports = stringifyHeaderFileNamesArray(@[@"ShaderTypes.h"]);
NSString *code = [NSString stringWithFormat:@"%s",
_STRINGIFY(
using namespace metal;
typedef struct {
float scale_factor;
float display_configuration;
} Uniforms;
typedef struct {
float4 renderedCoordinate [[position]];
float2 textureCoordinate;
} TextureMappingVertex;
vertex TextureMappingVertex mapTexture(unsigned int vertex_id [[ vertex_id ]],
constant Uniforms &uniform [[ buffer(1) ]])
{
float4x4 renderedCoordinates;
float4x2 textureCoordinates;
if (uniform.display_configuration == 0 ||
uniform.display_configuration == 2 ||
uniform.display_configuration == 4 ||
uniform.display_configuration == 6)
{
renderedCoordinates = float4x4(float4( -1.0, -1.0, 0.0, 1.0 ),
float4( 1.0, -1.0, 0.0, 1.0 ),
float4( -1.0, 1.0, 0.0, 1.0 ),
float4( 1.0, 1.0, 0.0, 1.0 ));
textureCoordinates = float4x2(float2( 0.0, 1.0 ),
float2( 2.0, 1.0 ),
float2( 0.0, 0.0 ),
float2( 2.0, 0.0 ));
} else if (uniform.display_configuration == 1 ||
uniform.display_configuration == 3 ||
uniform.display_configuration == 5 ||
uniform.display_configuration == 7)
{
renderedCoordinates = float4x4(float4( -1.0, -1.0, 0.0, 1.0 ),
float4( -1.0, 1.0, 0.0, 1.0 ),
float4( 1.0, -1.0, 0.0, 1.0 ),
float4( 1.0, 1.0, 0.0, 1.0 ));
if (uniform.display_configuration == 1 ||
uniform.display_configuration == 5)
{
textureCoordinates = float4x2(float2( 0.0, 1.0 ),
float2( 1.0, 1.0 ),
float2( 0.0, -1.0 ),
float2( 1.0, -1.0 ));
} else if (uniform.display_configuration == 3 ||
uniform.display_configuration == 7)
{
textureCoordinates = float4x2(float2( 0.0, 2.0 ),
float2( 1.0, 2.0 ),
float2( 0.0, 0.0 ),
float2( 1.0, 0.0 ));
}
}
TextureMappingVertex outVertex;
outVertex.renderedCoordinate = float4(uniform.scale_factor, uniform.scale_factor , 1.0f, 1.0f ) * renderedCoordinates[vertex_id];
outVertex.textureCoordinate = textureCoordinates[vertex_id];
return outVertex;
}
fragment half4 displayTexture(TextureMappingVertex mappingVertex [[ stage_in ]],
texture2d<float, access::sample> texture [[ texture(0) ]],
sampler samplr [[sampler(0)]],
constant Uniforms &uniform [[ buffer(1) ]]) {
if (uniform.display_configuration == 1 ||
uniform.display_configuration == 2 ||
uniform.display_configuration == 4 ||
uniform.display_configuration == 6 ||
uniform.display_configuration == 7)
{
mappingVertex.textureCoordinate.x = 1 - mappingVertex.textureCoordinate.x;
}
if (uniform.display_configuration == 2 ||
uniform.display_configuration == 6)
{
mappingVertex.textureCoordinate.y = 1 - mappingVertex.textureCoordinate.y;
}
if (uniform.scale_factor < 1.0)
{
mappingVertex.textureCoordinate.y += (texture.get_height(0) - (texture.get_height(0) * uniform.scale_factor));
}
half4 new_texture = half4(texture.sample(samplr, mappingVertex.textureCoordinate));
return new_texture;
}
)];
return [NSString stringWithFormat:@"%@\n%@", includes, imports, code];
}
/*
* Metal setup: Library
*/
__autoreleasing NSError *error = nil;
NSString* librarySrc = [self shader];
if(!librarySrc) {
[NSException raise:@"Failed to read shaders" format:@"%@", [error localizedDescription]];
}
_library = [_device newLibraryWithSource:librarySrc options:nil error:&error];
if(!_library) {
[NSException raise:@"Failed to compile shaders" format:@"%@", [error localizedDescription]];
}
id <MTLFunction> vertexProgram = [_library newFunctionWithName:@"mapTexture"];
id <MTLFunction> fragmentProgram = [_library newFunctionWithName:@"displayTexture"];
.
.
.
Adding my answer, as I found current replies unsatisfactory because they so-far have not addressed building metal GPU binaries (.metal files into .metallib binaries) without X-code installed, and without compiling them every time at runtime.
Commands such as
xcrun -sdk macosx metal MyLibrary.metal -o MyLibrary.air
Still require X-code installed. You can get around this.
First, configure your project to include the metal-cpp metal-cpp-extensions headers, available here: https://developer.apple.com/metal/LearnMetalCPP.zip
Second, write your shaders in .metal files, compile them once at runtime, and then build your metal GPU binaries using the headers from step 1.
Some sample code:
#define NS_PRIVATE_IMPLEMENTATION
#define MTL_PRIVATE_IMPLEMENTATION
#define MTK_PRIVATE_IMPLEMENTATION
#define CA_PRIVATE_IMPLEMENTATION
#include <Metal/Metal.hpp>
#include <MetalKit/MetalKit.hpp>
#include <AppKit/AppKit.hpp>
void saveBinary(MTL::RenderPipelineDescriptor* renderPipelineDescriptor) {
NS::Error* error = nullptr;
MTL::BinaryArchiveDescriptor* binaryArchiveDescriptor = MTL::BinaryArchiveDescriptor::alloc()->init();
MTL::BinaryArchive* binaryArchive = device->newBinaryArchive(binaryArchiveDescriptor, &error);
NS::URL* saveLocation = NS::URL::alloc()->initFileURLWithPath(NS::String::string("/binary.metallib", NS::StringEncoding::UTF8StringEncoding));
binaryArchive->addRenderPipelineFunctions(renderPipelineDescriptor, &error);
binaryArchive->serializeToURL(saveLocation, &serializeError);
}
The 'serializeToURL' function serializes to disk, while the binaryArchive object itself can be loaded with pipeline descriptors, including but not limited to the 'MTL::RenderPipelineDescriptor'. These descriptors themselves are initialised using your uncompiled kernel, fragment and vertex shader functions.
Upon initialisation of these descriptors, the first time, you compile the shader functions at runtime. But after you wrap up your development, you can compile the binary and skip compilation at runtime all together. Without any X-code involvement.
This process is essentially the same in swift as it would be in c++. I found this video helpful and my answer is based on this: https://developer.apple.com/videos/play/wwdc2020/10615/
You can use CMake to create a project that supports building containing source code in different languages (c/c++/objc/swift/etc). I give the following fragments as an example, which can be used as the basis for a more complex project.
CMake file:
set(METAL_TARGET_FLAGS "air64-apple-macos11.0")
set(CURRENT_KERNEL_TARGET_NAME "example")
set(CURRENT_TARGET_NAME "${CURRENT_KERNEL_TARGET_NAME}_metal")
set(CURRENT_TARGET_LIB ${CURRENT_TARGET_NAME}.metallib)
include_directories(${CMAKE_CURRENT_SOURCE_DIR})
file (GLOB DEPENDS_SOURCES
${CMAKE_CURRENT_SOURCE_DIR}/*.h
)
string(STRIP METAL_TARGET_FLAGS ${METAL_TARGET_FLAGS})
add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${CURRENT_KERNEL_TARGET_NAME}.air
DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/${CURRENT_KERNEL_TARGET_NAME}.metal ${DEPENDS_SOURCES} ${COMMON_DEPENDENCIES}
COMMAND xcrun -v -sdk macosx metal
${METAL_FLAGS}
-I .
-I ${CMAKE_CURRENT_SOURCE_DIR}
-I ${CMAKE_SOURCE_DIR}/include
-I ${CMAKE_SOURCE_DIR}/tests/shaders
-O3 -ffast-math
-c
-Wno-unused-variable
-target ${METAL_TARGET_FLAGS}
${CMAKE_CURRENT_SOURCE_DIR}/${CURRENT_KERNEL_TARGET_NAME}.metal
-o ${CMAKE_CURRENT_BINARY_DIR}/${CURRENT_KERNEL_TARGET_NAME}.air
VERBATIM
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
)
add_custom_target(
${CURRENT_TARGET_NAME}
DEPENDS ${DEPENDS_SOURCES}
${CMAKE_CURRENT_BINARY_DIR}/${CURRENT_KERNEL_TARGET_NAME}.air
COMMAND xcrun -sdk macosx metallib
${CMAKE_CURRENT_BINARY_DIR}/*.air
-o ${CURRENT_TARGET_LIB}
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
)
set_target_properties(${CURRENT_TARGET_NAME} PROPERTIES BUNDLE TRUE)
set_target_properties(${CURRENT_TARGET_NAME} PROPERTIES PREFIX "")
set_target_properties(${CURRENT_TARGET_NAME} PROPERTIES SUFFIX ".metallib")
Where example.metal makes simple copping pixel from source texture to destination:
#include "example.h"
/***
* Pass kernel
* @param source
* @param destination
* @return
*/
kernel void kernel_by_pass(
metal::texture2d<float, metal::access::sample> source,
metal::texture2d<float, metal::access::write> destination,
uint2 tid [[thread_position_in_grid]]
){
uint w = destination.get_width();
uint h = destination.get_height();
float2 coords = float2(tid.x,tid.y)/float2(w,h) ;
float4 color = source.sample(linear_normalized_sampler, coords);
destination.write(color,tid);
}
and example.h is a common source preferences:
#ifndef METAL_EXAMPLE_H
#define METAL_EXAMPLE_H
#include <metal_stdlib>
constexpr metal::sampler linear_normalized_sampler(metal::address::mirrored_repeat,
metal::filter::linear,
metal::coord::normalized,
metal::mag_filter::linear,
metal::min_filter::linear,
metal::mip_filter::linear
);
/* common prefs */
#endif //METAL_EXAMPLE_H
© 2022 - 2024 — McMap. All rights reserved.