I'm trying to follow WWDC2020/21 in making simplest use of Metal's binary linked functions.
I attach a visible function with CompileToBinary
to the binary linked functions of a compute pipeline descriptor.
I'm getting a
segmentation fault (I think), which occurs as the program suspending
with Thread 1: EXC_BAD_ACCESS (code=1, address=0x10)
where the device creates the pipeline (the program not reaching the NSLog
).
I can't figure out how to address this with memory debugging tools. Some points:
- Checking all relevant objects right before making the call (
[device newComputePipeline...]
) yields expected results. - I get the same error when making the call asynchronously.
- Adding
visible_func
to linkedfunctions
instead ofbinaryFunctions
works. - I've also been following the FunctionPointersAndStitching example code (for rendering rather than compute) to make sure I'm setting needed properties and creating things in the right order.
- I can't even begin to use these linked functions (creating a visible function table), until I create the pipeline, so stopping short after creating the pipeline doesn't seem a candidate issue.
#import <Foundation/Foundation.h>
#import <Metal/Metal.h>
int main() {
@autoreleasepool {
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
id<MTLLibrary> lib = [device newDefaultLibrary];
id<MTLFunction> kernel_func = [lib newFunctionWithName:@"kernel_func"];
MTLFunctionDescriptor * visible_func_desc = [MTLFunctionDescriptor functionDescriptor];
visible_func_desc.name = @"visible_func";
visible_func_desc.options = MTLFunctionOptionCompileToBinary;
NSError * error = nil;
id<MTLFunction> visible_func = [lib newFunctionWithDescriptor:visible_func_desc error: &error];
MTLLinkedFunctions *linked_funcs = [MTLLinkedFunctions linkedFunctions];
linked_funcs.binaryFunctions = @[visible_func];
MTLComputePipelineDescriptor * pipeline_desc = [MTLComputePipelineDescriptor new];
pipeline_desc.linkedFunctions = linked_funcs;
pipeline_desc.computeFunction = kernel_func;
MTLAutoreleasedComputePipelineReflection * reflection = nil;
id<MTLComputePipelineState> pipeline = [device newComputePipelineStateWithDescriptor:pipeline_desc options:MTLPipelineOptionNone reflection:reflection error:&error];
if(pipeline == nil) {
NSLog(@"Error making pipeline: %@",error);
}
}
}
#include <metal_stdlib>
using namespace metal;
[[visible]] int visible_func(int x, int y) {
return x + y;
}
[[kernel]] void kernel_func(uint gid [[ thread_position_in_grid ]])
{
gid;
}