Skip to content

Instantly share code, notes, and snippets.

@wakita
Last active July 21, 2023 08:53
Show Gist options
  • Save wakita/f4915757c6c6c128c05c8680cd859e1a to your computer and use it in GitHub Desktop.
Save wakita/f4915757c6c6c128c05c8680cd859e1a to your computer and use it in GitHub Desktop.
Metal compute shader example
#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;
}
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))
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
@rasky
Copy link

rasky commented Apr 22, 2020

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.

@rcorin
Copy link

rcorin commented Oct 31, 2020

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

@rasky
Copy link

rasky commented Nov 1, 2020

No sorry @rcorin

@DmitrySkripkin
Copy link

Thanks a lot, official Apple documentation lacks Swift examples.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment