diff --git a/README.md b/README.md index f36d39d..a7f010c 100644 --- a/README.md +++ b/README.md @@ -35,3 +35,4 @@ Repository to accompany the following blog posts: - [Working with Particles in Metal part 3](http://metalkit.org/2017/11/30/working-with-particles-in-metal-part-3.html) - [Metal By Tutorials book!](http://metalkit.org/2018/05/29/metal-by-tutorials-book.html) - [Raytracing with Metal](http://metalkit.org/2018/07/14/raytracing-with-metal.html) +- [Introduction to compute using Metal](http://metalkit.org/2019/01/31/intro-to-metal-compute.html) diff --git a/compute/compute.playground/Contents.swift b/compute/compute.playground/Contents.swift new file mode 100644 index 0000000..305228a --- /dev/null +++ b/compute/compute.playground/Contents.swift @@ -0,0 +1,9 @@ + +import MetalKit +import PlaygroundSupport + +let frame = NSRect(x: 0, y: 0, width: 800, height: 500) +let delegate = Renderer() +let view = MTKView(frame: frame, device: delegate.device) +view.delegate = delegate +PlaygroundPage.current.liveView = view diff --git a/compute/compute.playground/Resources/Shaders.metal b/compute/compute.playground/Resources/Shaders.metal new file mode 100644 index 0000000..79b4803 --- /dev/null +++ b/compute/compute.playground/Resources/Shaders.metal @@ -0,0 +1,11 @@ + +#include +using namespace metal; + +kernel void compute(texture2d input [[texture(0)]], + texture2d output [[texture(1)]], + uint2 id [[thread_position_in_grid]]) { + uint2 index = uint2((id.x / 5) * 5, (id.y / 5) * 5); + float4 color = input.read(index); + output.write(color, id); +} diff --git a/compute/compute.playground/Resources/nature.jpg b/compute/compute.playground/Resources/nature.jpg new file mode 100644 index 0000000..d480305 Binary files /dev/null and b/compute/compute.playground/Resources/nature.jpg differ diff --git a/compute/compute.playground/Sources/MetalView.swift b/compute/compute.playground/Sources/MetalView.swift new file mode 100644 index 0000000..6ba430a --- /dev/null +++ b/compute/compute.playground/Sources/MetalView.swift @@ -0,0 +1,57 @@ + +import MetalKit + +public class Renderer: NSObject, MTKViewDelegate { + + public var device: MTLDevice! + var queue: MTLCommandQueue! + var pipelineState: MTLComputePipelineState! + var image: MTLTexture! + + override public init() { + super.init() + initializeMetal() + } + + func initializeMetal() { + device = MTLCreateSystemDefaultDevice() + queue = device.makeCommandQueue() + let textureLoader = MTKTextureLoader(device: device) + let url = Bundle.main.url(forResource: "nature", withExtension: "jpg")! + guard let file = Bundle.main.path(forResource: "Shaders", ofType: "metal") else { return } + do { + let source = try String(contentsOfFile: file, encoding: String.Encoding.utf8) + let library = try device.makeLibrary(source: source, options: nil) + guard let function = library.makeFunction(name: "compute") else { return } + pipelineState = try device.makeComputePipelineState(function: function) + image = try textureLoader.newTexture(URL: url, options: [:]) + } catch let error { + print(error.localizedDescription) + } + } + + public func draw(in view: MTKView) { + guard let commandBuffer = queue.makeCommandBuffer(), + let commandEncoder = commandBuffer.makeComputeCommandEncoder(), + let drawable = view.currentDrawable else { + return + } + commandEncoder.setComputePipelineState(pipelineState) + commandEncoder.setTexture(image, index: 0) + commandEncoder.setTexture(drawable.texture, index: 1) + + var width = pipelineState.threadExecutionWidth + var height = pipelineState.maxTotalThreadsPerThreadgroup / width + let threadsPerGroup = MTLSizeMake(width, height, 1) + width = Int(view.drawableSize.width) + height = Int(view.drawableSize.height) + let threadsPerGrid = MTLSizeMake(width, height, 1) + commandEncoder.dispatchThreads(threadsPerGrid, threadsPerThreadgroup: threadsPerGroup) + + commandEncoder.endEncoding() + commandBuffer.present(drawable) + commandBuffer.commit() + } + + public func mtkView(_ view: MTKView, drawableSizeWillChange size: CGSize) {} +} diff --git a/compute/compute.playground/contents.xcplayground b/compute/compute.playground/contents.xcplayground new file mode 100644 index 0000000..b478c1d --- /dev/null +++ b/compute/compute.playground/contents.xcplayground @@ -0,0 +1,2 @@ + + \ No newline at end of file diff --git a/compute/compute.playground/playground.xcworkspace/contents.xcworkspacedata b/compute/compute.playground/playground.xcworkspace/contents.xcworkspacedata new file mode 100644 index 0000000..919434a --- /dev/null +++ b/compute/compute.playground/playground.xcworkspace/contents.xcworkspacedata @@ -0,0 +1,7 @@ + + + + + diff --git a/compute/compute.playground/playground.xcworkspace/xcshareddata/IDEWorkspaceChecks.plist b/compute/compute.playground/playground.xcworkspace/xcshareddata/IDEWorkspaceChecks.plist new file mode 100644 index 0000000..18d9810 --- /dev/null +++ b/compute/compute.playground/playground.xcworkspace/xcshareddata/IDEWorkspaceChecks.plist @@ -0,0 +1,8 @@ + + + + + IDEDidComputeMac32BitWarning + + + diff --git a/compute/compute.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate b/compute/compute.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate new file mode 100644 index 0000000..a413dde Binary files /dev/null and b/compute/compute.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate differ