diff --git a/README.md b/README.md index b71a78e..5f4115d 100644 --- a/README.md +++ b/README.md @@ -25,10 +25,11 @@ Repository to accompany the following blog posts: - [Shadows in Metal part 1](http://metalkit.org/2017/01/31/shadows-in-metal-part-1.html) - [Shadows in Metal part 2](http://metalkit.org/2017/02/28/shadows-in-metal-part-2.html) - [Ambient Occlusion in Metal](http://metalkit.org/2017/03/22/ambient-occlusion-in-metal.html) -- [Working with memory in Metal](http://metalkit.org/2017/04/30/working-with-memory-in-metal.html) +- [Working with memory in Metal part 1](http://metalkit.org/2017/04/30/working-with-memory-in-metal.html) - [Working with memory in Metal part 2](http://metalkit.org/2017/05/26/working-with-memory-in-metal-part-2.html) - [Introducing Metal 2](http://metalkit.org/2017/06/30/introducing-metal-2.html) -- [Using ARKit with Metal](http://metalkit.org/2017/07/29/using-arkit-with-metal.html) +- [Using ARKit with Metal part 1](http://metalkit.org/2017/07/29/using-arkit-with-metal.html) - [Using ARKit with Metal part 2](http://metalkit.org/2017/08/31/using-arkit-with-metal-part-2.html) -- [Working with Particles in Metal](http://metalkit.org/2017/09/30/working-with-particles-in-metal.html) +- [Working with Particles in Metal part 1](http://metalkit.org/2017/09/30/working-with-particles-in-metal.html) - [Working with Particles in Metal part 2](http://metalkit.org/2017/10/31/working-with-particles-in-metal-part-2.html) +- [Working with Particles in Metal part 3](http://metalkit.org/2017/11/30/working-with-particles-in-metal-part-3.html) \ No newline at end of file diff --git a/particles/particle3.playground/Contents.swift b/particles/particle3.playground/Contents.swift new file mode 100644 index 0000000..fcee78a --- /dev/null +++ b/particles/particle3.playground/Contents.swift @@ -0,0 +1,9 @@ + +import MetalKit +import PlaygroundSupport + +let frame = NSRect(x: 0, y: 0, width: 600, height: 600) +let mView = MetalView() +let view = MTKView(frame: frame, device: mView.device) +view.delegate = mView +PlaygroundPage.current.liveView = view diff --git a/particles/particle3.playground/Resources/Shaders.metal b/particles/particle3.playground/Resources/Shaders.metal new file mode 100644 index 0000000..8a6345c --- /dev/null +++ b/particles/particle3.playground/Resources/Shaders.metal @@ -0,0 +1,35 @@ + +#include +using namespace metal; + +struct Particle { + float2 position; + float2 velocity; +}; + +kernel void firstPass(texture2d output [[texture(0)]], + uint2 id [[thread_position_in_grid]]) { + output.write(half4(0., 0., 0., 1.), id); +} + +kernel void secondPass(texture2d output [[texture(0)]], + device Particle *particles [[buffer(0)]], + uint id [[thread_position_in_grid]]) { + Particle particle = particles[id]; + float2 position = particle.position; + float2 velocity = particle.velocity; + int width = output.get_width(); + int height = output.get_height(); + if (position.x < 0 || position.x > width) { velocity.x *= -1; } + if (position.y < 0 || position.y > height) { velocity.y *= -1; } + position += velocity; + particle.position = position; + particle.velocity = velocity; + particles[id] = particle; + uint2 pos = uint2(position.x, position.y); + output.write(half4(1.), pos); + output.write(half4(1.), pos + uint2( 1, 0)); + output.write(half4(1.), pos + uint2( 0, 1)); + output.write(half4(1.), pos - uint2( 1, 0)); + output.write(half4(1.), pos - uint2( 0, 1)); +} diff --git a/particles/particle3.playground/Sources/MetalView.swift b/particles/particle3.playground/Sources/MetalView.swift new file mode 100644 index 0000000..fc1f7b6 --- /dev/null +++ b/particles/particle3.playground/Sources/MetalView.swift @@ -0,0 +1,74 @@ + +import MetalKit + +struct Particle { + var position: float2 + var velocity: float2 +} + +public class MetalView: NSObject, MTKViewDelegate { + + public var device: MTLDevice! + var queue: MTLCommandQueue! + var firstState: MTLComputePipelineState! + var secondState: MTLComputePipelineState! + var particleBuffer: MTLBuffer! + let particleCount = 10000 + var particles = [Particle]() + let side = 1200 + + override public init() { + super.init() + initializeMetal() + initializeBuffers() + } + + func initializeBuffers() { + for _ in 0 ..< particleCount { + let particle = Particle(position: float2(Float(arc4random() % UInt32(side)), Float(arc4random() % UInt32(side))), velocity: float2((Float(arc4random() % 10) - 5) / 10, (Float(arc4random() % 10) - 5) / 10)) + particles.append(particle) + } + let size = particles.count * MemoryLayout.size + particleBuffer = device.makeBuffer(bytes: &particles, length: size, options: []) + } + + func initializeMetal() { + device = MTLCreateSystemDefaultDevice() + queue = device.makeCommandQueue() + guard let path = Bundle.main.path(forResource: "Shaders", ofType: "metal") else { return } + do { + let input = try String(contentsOfFile: path, encoding: String.Encoding.utf8) + let library = try device.makeLibrary(source: input, options: nil) + guard let firstPass = library.makeFunction(name: "firstPass") else { return } + firstState = try device.makeComputePipelineState(function: firstPass) + guard let secondPass = library.makeFunction(name: "secondPass") else { return } + secondState = try device.makeComputePipelineState(function: secondPass) + } catch let e { print(e) } + } + + public func mtkView(_ view: MTKView, drawableSizeWillChange size: CGSize) {} + + public func draw(in view: MTKView) { + if let drawable = view.currentDrawable, + let commandBuffer = queue.makeCommandBuffer(), + let commandEncoder = commandBuffer.makeComputeCommandEncoder() { + // first pass + commandEncoder.setComputePipelineState(firstState) + commandEncoder.setTexture(drawable.texture, index: 0) + let w = firstState.threadExecutionWidth + let h = firstState.maxTotalThreadsPerThreadgroup / w + let threadsPerGroup = MTLSizeMake(w, h, 1) + var threadsPerGrid = MTLSizeMake(side, side, 1) + commandEncoder.dispatchThreads(threadsPerGrid, threadsPerThreadgroup: threadsPerGroup) + // second pass + commandEncoder.setComputePipelineState(secondState) + commandEncoder.setTexture(drawable.texture, index: 0) + commandEncoder.setBuffer(particleBuffer, offset: 0, index: 0) + threadsPerGrid = MTLSizeMake(particleCount, 1, 1) + commandEncoder.dispatchThreads(threadsPerGrid, threadsPerThreadgroup: threadsPerGroup) + commandEncoder.endEncoding() + commandBuffer.present(drawable) + commandBuffer.commit() + } + } +} diff --git a/particles/particle3.playground/contents.xcplayground b/particles/particle3.playground/contents.xcplayground new file mode 100644 index 0000000..a93d484 --- /dev/null +++ b/particles/particle3.playground/contents.xcplayground @@ -0,0 +1,4 @@ + + + + \ No newline at end of file diff --git a/particles/particle3.playground/playground.xcworkspace/contents.xcworkspacedata b/particles/particle3.playground/playground.xcworkspace/contents.xcworkspacedata new file mode 100644 index 0000000..919434a --- /dev/null +++ b/particles/particle3.playground/playground.xcworkspace/contents.xcworkspacedata @@ -0,0 +1,7 @@ + + + + + diff --git a/particles/particle3.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate b/particles/particle3.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate new file mode 100644 index 0000000..29c0411 Binary files /dev/null and b/particles/particle3.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate differ diff --git a/particles/particle3.playground/timeline.xctimeline b/particles/particle3.playground/timeline.xctimeline new file mode 100644 index 0000000..bf468af --- /dev/null +++ b/particles/particle3.playground/timeline.xctimeline @@ -0,0 +1,6 @@ + + + + +