Last active
July 21, 2022 13:12
-
-
Save icosahedron/e2a5c2c87744bd82a0357725645acf9c to your computer and use it in GitHub Desktop.
Metal Compute Shader from Command Line
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
#!/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 |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
//-*- 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]; | |
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
// 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)") | |
} |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Updated version of https://gist.github.com/bellbind/431cafee38210a684991.
This has been tested with Swift 5.1 and Xcode 13.