Skip to content

Instantly share code, notes, and snippets.

@icosahedron
Last active July 21, 2022 13:12
Show Gist options
  • Star 2 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save icosahedron/e2a5c2c87744bd82a0357725645acf9c to your computer and use it in GitHub Desktop.
Save icosahedron/e2a5c2c87744bd82a0357725645acf9c to your computer and use it in GitHub Desktop.
Metal Compute Shader from Command Line
#!/bin/bash
xcrun -sdk macosx metal -Wall -Wextra -std=osx-metal1.1 Shaders.metal -c -o Shaders.air
xcrun -sdk macosx metal-ar rcs Shaders.metal-ar Shaders.air
xcrun -sdk macosx metallib -o default.metallib Shaders.metal-ar
xcrun -sdk macosx swiftc usingmetal.swift -o usingmetal
./usingmetal
//-*- mode: c++ -*-
// build:
// See run.sh
#include <metal_stdlib>
kernel void square(const device float* input [[buffer(0)]],
device float* output [[buffer(1)]],
metal::uint id [[thread_position_in_grid]]) {
output[id] = input[id] * input[id];
}
// Metal programming with commandline
// 1. build a ".metallib" file from shader sources ".metal"
// 2. run swift command with metal used swift file
import Metal
import Darwin
// input data
let count = 1024
var data = Array<Float>(repeating: 0, count: count)
var i = 0
for (i, _) in data.enumerated() {
data[i] = Float(i)
}
// metal preparation
let device = MTLCreateSystemDefaultDevice()
if device == nil {
print("This device is not supported Metal")
exit(1)
}
let dev = device!
// `dev.newDefaultLibrary()` applied only when compiled exe
// as "./default.metallib" or "./resources/defaut.metallib" relative from exe
//let lib = dev.newDefaultLibrary()!
let lib = try! dev.makeLibrary(filepath:"default.metallib")
let queue = dev.makeCommandQueue()
let commands = queue!.makeCommandBuffer()
// encoding begin
let encoder = commands!.makeComputeCommandEncoder()
// shader program
let fun = lib.makeFunction(name:"square")!;
let state = try! dev.makeComputePipelineState(function:fun)
encoder!.setComputePipelineState(state)
// as buffer
let size = count * MemoryLayout<Float>.size
let options : MTLResourceOptions = [] // MTLResourceOptions.CPUCacheModeDefaultCache
let input = dev.makeBuffer(bytes:&data, length: size, options: options)
let output = dev.makeBuffer(length: size, options: options)
encoder!.setBuffer(input, offset: 0, index: 0)
encoder!.setBuffer(output, offset: 0, index: 1)
// thread-group setting
let t = 32
let g = count / t + (count % t == 0 ? 0 : 1)
let threads = MTLSize(width: t, height: 1, depth: 1)
let groups = MTLSize(width: g, height: 1, depth: 1)
encoder!.dispatchThreadgroups(groups, threadsPerThreadgroup: threads)
// encoding end
encoder!.endEncoding()
// execute
commands!.commit()
commands!.waitUntilCompleted()
let presult = output!.contents().bindMemory(to: Float.self, capacity: count)
let result = Array(UnsafeBufferPointer(start: presult, count: count))
for (i, r) in result.enumerated() {
print("data: \(data[i]) square: \(r)")
}
@icosahedron
Copy link
Author

Updated version of https://gist.github.com/bellbind/431cafee38210a684991.

This has been tested with Swift 5.1 and Xcode 13.

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