Metal compute shader example

compute.metal

#include <metal_stdlib>
using namespace metal;

kernel void add(const device float2 *in [[ buffer(0) ]],
                device float  *out [[ buffer(1) ]],
                uint id [[ thread_position_in_grid ]]) {
  out[id] = in[id].x + in[id].y;
}

compute.swift

import MetalKit

let device = MTLCreateSystemDefaultDevice()!
let commandQueue = device.makeCommandQueue()!
let library = try device.makeLibrary(filepath: "compute.metallib")

let commandBuffer = commandQueue.makeCommandBuffer()!
let encoder = commandBuffer.makeComputeCommandEncoder()!
encoder.setComputePipelineState(try device.makeComputePipelineState(function: library.makeFunction(name: "add")!))

let input: [Float] = [1.0, 2.0]
encoder.setBuffer(device.makeBuffer(bytes: input as [Float], length: MemoryLayout<Float>.stride * input.count, options: []),
                  offset: 0, index: 0)
let outputBuffer = device.makeBuffer(length: MemoryLayout<Float>.stride, options: [])!
encoder.setBuffer(outputBuffer, offset: 0, index: 1)

let numThreadgroups = MTLSize(width: 1, height: 1, depth: 1)
let threadsPerThreadgroup = MTLSize(width: 1, height: 1, depth: 1)
encoder.dispatchThreadgroups(numThreadgroups, threadsPerThreadgroup: threadsPerThreadgroup)
encoder.endEncoding()

commandBuffer.commit()
commandBuffer.waitUntilCompleted()

let result = outputBuffer.contents().load(as: Float.self)
print(String(format: "%f + %f = %f", input[0], input[1], result))

Makefile

SDK = xcrun -sdk macosx

all: compute.metallib compute

compute.metallib: Compute.metal
	# Metal intermediate representation (.air)
	$(SDK) metal -c -Wall -Wextra -std=osx-metal2.0 -o /tmp/Compute.air $^
	# Metal library (.metallib)
	$(SDK) metallib -o $@ /tmp/Compute.air

compute: compute.swift
	$(SDK) swiftc -o $@ $^

clean:
	rm -f compute.metallib compute

Hi, trying to compile and run this, I get a binary that segfaults:

 thread #1, queue = 'com.apple.main-thread', stop reason = EXC_BAD_ACCESS (code=1, address=0x8)
  * frame #0: 0x00007fff40c3f9aa AppleIntelSKLGraphicsMTLDriver`CHAL_INTEL::ChalContext::IsValidCurbeEntry(unsigned long long) + 18
    frame #1: 0x00007fff40bdce1b AppleIntelSKLGraphicsMTLDriver`CHAL_INTEL::ChalContext::PatchScalarCurbeEntryWithStartOffset(unsigned char*, unsigned long long, unsigned int, unsigned int) + 37
    frame #2: 0x00007fff40bdce89 AppleIntelSKLGraphicsMTLDriver`CHAL_INTEL::ChalContext::PatchLocalWorkSizeCurbeEntry(unsigned char*, unsigned int*) + 43
    frame #3: 0x00007fff40bdcbea AppleIntelSKLGraphicsMTLDriver`CHAL_INTEL::ChalContext::PatchCurbeData(CHAL_INTEL::tagCHALCommandPacket*) + 232
    frame #4: 0x00007fff40bdcaba AppleIntelSKLGraphicsMTLDriver`CHAL_INTEL::ChalContext::PatchDescriptors(CHAL_INTEL::tagCHALCommandPacket*) + 346
    frame #5: 0x00007fff40c3a0d5 AppleIntelSKLGraphicsMTLDriver`CHAL_INTEL::ChalContext::ChalExecuteKernel(CHAL_INTEL::tagDecodedBinaryInfo const*, MTLSize&, MTLSize&, MTLSize&, IGBufferMemory<31ul>&, sTextureArgument const*, sSamplerArgument const*, unsigned int const*, MTLRegion, MTLIGAccelBuffer*, unsigned int, MTLIGAccelBuffer*, unsigned int) + 941
    frame #6: 0x00007fff40bef6f0 AppleIntelSKLGraphicsMTLDriver`IGAccelComputeCommandEncoder::dispatchThreadgroups(MTLSize, MTLSize) + 172
    frame #7: 0x00007fff40bef594 AppleIntelSKLGraphicsMTLDriver`-[MTLIGAccelComputeCommandEncoder dispatchThreadgroups:threadsPerThreadgroup:] + 234
    frame #8: 0x00000001000023b8 compute`main at compute.swift:19:9
    frame #9: 0x00007fff724dd3d5 libdyld.dylib`start + 1
    frame #10: 0x00007fff724dd3d5 libdyld.dylib`start + 1

Any idea? I'm running Mojave 10.14.6 with Xcode 11.3.1.

hi @rasky, i am chasing the same bug with a different metal library, did you find the cause to this "IsValidCurbeEntry" bug?
thanks
Ricardo

No sorry @rcorin

Thanks a lot, official Apple documentation lacks Swift examples.