diff --git a/README.md b/README.md index 563835a..4662ffb 100644 --- a/README.md +++ b/README.md @@ -22,4 +22,5 @@ Repository to accompany the following blog posts: - [Using MetalKit part 18](http://metalkit.org/2016/10/01/using-metalkit-part-2-3-2.html) - [Raymarching in Metal](http://metalkit.org/2016/12/30/raymarching-in-metal.html) - [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) \ No newline at end of file +- [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) \ No newline at end of file diff --git a/ambient_occlusion/ao.playground/Contents.swift b/ambient_occlusion/ao.playground/Contents.swift new file mode 100644 index 0000000..694d377 --- /dev/null +++ b/ambient_occlusion/ao.playground/Contents.swift @@ -0,0 +1,9 @@ + +import MetalKit +import PlaygroundSupport + +let frame = NSRect(x: 0, y: 0, width: 400, height: 400) +let delegate = MetalView() +let view = MTKView(frame: frame, device: delegate.device) +view.delegate = delegate +PlaygroundPage.current.liveView = view diff --git a/ambient_occlusion/ao.playground/Resources/Shaders.metal b/ambient_occlusion/ao.playground/Resources/Shaders.metal new file mode 100644 index 0000000..9843452 --- /dev/null +++ b/ambient_occlusion/ao.playground/Resources/Shaders.metal @@ -0,0 +1,156 @@ + +#include +using namespace metal; + +struct Ray { + float3 origin; + float3 direction; + Ray(float3 o, float3 d) { + origin = o; + direction = d; + } +}; + +struct Sphere { + float3 center; + float radius; + Sphere(float3 c, float r) { + center = c; + radius = r; + } +}; + +struct Plane { + float yCoord; + Plane(float y) { + yCoord = y; + } +}; + +struct Box { + float3 center; + float size; + Box(float3 c, float s) { + center = c; + size = s; + } +}; + +struct Camera { + float3 position; + Ray ray = Ray(float3(0), float3(0)); + float rayDivergence; + Camera(float3 pos, Ray r, float div) { + position = pos; + ray = r; + rayDivergence = div; + } +}; + +float unionOp(float d0, float d1) { + return min(d0, d1); +} + +float differenceOp(float d0, float d1) { + return max(d0, -d1); +} + +float distToSphere(Ray ray, Sphere s) { + return length(ray.origin - s.center) - s.radius; +} + +float distToPlane(Ray ray, Plane plane) { + return ray.origin.y - plane.yCoord; +} + +float distToBox(Ray r, Box b) { + float3 d = abs(r.origin - b.center) - float3(b.size); + return min(max(d.x, max(d.y, d.z)), 0.0) + length(max(d, 0.0)); +} + +float distToScene(Ray r) { + Plane p = Plane(0.0); + float d2p = distToPlane(r, p); + Sphere s1 = Sphere(float3(0.0, 0.5, 0.0), 8.0); + Sphere s2 = Sphere(float3(0.0, 0.5, 0.0), 6.0); + Sphere s3 = Sphere(float3(10., -5., -10.), 15.0); + Box b = Box(float3(1., 1., -4.), 1.); + float dtb = distToBox(r, b); + float d2s1 = distToSphere(r, s1); + float d2s2 = distToSphere(r, s2); + float d2s3 = distToSphere(r, s3); + float dist = differenceOp(d2s1, d2s2); + dist = differenceOp(dist, d2s3); + dist = unionOp(dist, dtb); + dist = unionOp(d2p, dist); + return dist; +} + +float3 getNormal(Ray ray) { + float2 eps = float2(0.001, 0.0); + float3 n = float3(distToScene(Ray(ray.origin + eps.xyy, ray.direction)) - + distToScene(Ray(ray.origin - eps.xyy, ray.direction)), + distToScene(Ray(ray.origin + eps.yxy, ray.direction)) - + distToScene(Ray(ray.origin - eps.yxy, ray.direction)), + distToScene(Ray(ray.origin + eps.yyx, ray.direction)) - + distToScene(Ray(ray.origin - eps.yyx, ray.direction))); + return normalize(n); +} + +float ao(float3 pos, float3 n) { + float eps = 0.01; + pos += n * eps * 2.0; + float occlusion = 0.0; + for (float i=1.0; i<10.0; i++) { + float d = distToScene(Ray(pos, float3(0))); + float coneWidth = 2.0 * eps; + float occlusionAmount = max(coneWidth - d, 0.); + float occlusionFactor = occlusionAmount / coneWidth; + occlusionFactor *= 1.0 - (i / 10.0); + occlusion = max(occlusion, occlusionFactor); + eps *= 2.0; + pos += n * eps; + } + return max(0.0, 1.0 - occlusion); +} + +Camera setupCam(float3 pos, float3 target, float fov, float2 uv, int x) { + uv *= fov; + float3 cw = normalize(target - pos ); + float3 cp = float3(0.0, 1.0, 0.0); + float3 cu = normalize(cross(cw, cp)); + float3 cv = normalize(cross(cu, cw)); + Ray ray = Ray(pos, normalize(uv.x * cu + uv.y * cv + 0.5 * cw)); + Camera cam = Camera(pos, ray, fov / float(x)); + return cam; +} + +kernel void compute(texture2d output [[texture(0)]], + constant float &time [[buffer(0)]], + uint2 gid [[thread_position_in_grid]]) { + int width = output.get_width(); + int height = output.get_height(); + float2 uv = float2(gid) / float2(width, height); + uv = uv * 2.0 - 1.0; + uv.y = -uv.y; + float3 camPos = float3(sin(time) * 10., 3., cos(time) * 10.); + Camera cam = setupCam(camPos, float3(0), 1.25, uv, width); + float3 col = float3(1.0); + bool hit = false; + for (int i=0; i<200; i++) { + float dist = distToScene(cam.ray); + if (dist < 0.001) { + hit = true; + break; + } + cam.ray.origin += cam.ray.direction * dist; + } + if (!hit) { + col = float3(0.5); + } else { + float3 n = getNormal(cam.ray); + float o = ao(cam.ray.origin, n); + col = col * o; + } + output.write(float4(col, 1.0), gid); +} diff --git a/ambient_occlusion/ao.playground/Sources/MetalView.swift b/ambient_occlusion/ao.playground/Sources/MetalView.swift new file mode 100644 index 0000000..bf69602 --- /dev/null +++ b/ambient_occlusion/ao.playground/Sources/MetalView.swift @@ -0,0 +1,56 @@ + +import MetalKit + +public class MetalView: NSObject, MTKViewDelegate { + + public var device: MTLDevice! = nil + var queue: MTLCommandQueue! = nil + var cps: MTLComputePipelineState! = nil + var timerBuffer: MTLBuffer! = nil + var timer: Float = 0 + + override public init() { + super.init() + device = MTLCreateSystemDefaultDevice() + queue = device.makeCommandQueue() + registerShaders() + } + + func registerShaders() { + 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 kernel = library.makeFunction(name: "compute") else { return } + cps = try device.makeComputePipelineState(function: kernel) + } catch let e { + Swift.print("\(e)") + } + timerBuffer = device.makeBuffer(length: MemoryLayout.size, options: []) + } + + func update() { + timer += 0.01 + let bufferPointer = timerBuffer.contents() + memcpy(bufferPointer, &timer, MemoryLayout.size) + } + + 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() + commandEncoder.setComputePipelineState(cps) + commandEncoder.setTexture(drawable.texture, at: 0) + commandEncoder.setBuffer(timerBuffer, offset: 0, at: 0) + update() + let threadGroupCount = MTLSizeMake(8, 8, 1) + let threadGroups = MTLSizeMake(drawable.texture.width / threadGroupCount.width, drawable.texture.height / threadGroupCount.height, 1) + commandEncoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadGroupCount) + commandEncoder.endEncoding() + commandBuffer.present(drawable) + commandBuffer.commit() + } + } +} diff --git a/ambient_occlusion/ao.playground/contents.xcplayground b/ambient_occlusion/ao.playground/contents.xcplayground new file mode 100644 index 0000000..a93d484 --- /dev/null +++ b/ambient_occlusion/ao.playground/contents.xcplayground @@ -0,0 +1,4 @@ + + + + \ No newline at end of file diff --git a/ambient_occlusion/ao.playground/playground.xcworkspace/contents.xcworkspacedata b/ambient_occlusion/ao.playground/playground.xcworkspace/contents.xcworkspacedata new file mode 100644 index 0000000..919434a --- /dev/null +++ b/ambient_occlusion/ao.playground/playground.xcworkspace/contents.xcworkspacedata @@ -0,0 +1,7 @@ + + + + + diff --git a/ambient_occlusion/ao.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate b/ambient_occlusion/ao.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate new file mode 100644 index 0000000..1ea18a6 Binary files /dev/null and b/ambient_occlusion/ao.playground/playground.xcworkspace/xcuserdata/marius.xcuserdatad/UserInterfaceState.xcuserstate differ diff --git a/ambient_occlusion/ao.playground/playground.xcworkspace/xcuserdata/mhorga.xcuserdatad/UserInterfaceState.xcuserstate b/ambient_occlusion/ao.playground/playground.xcworkspace/xcuserdata/mhorga.xcuserdatad/UserInterfaceState.xcuserstate new file mode 100644 index 0000000..501a0aa Binary files /dev/null and b/ambient_occlusion/ao.playground/playground.xcworkspace/xcuserdata/mhorga.xcuserdatad/UserInterfaceState.xcuserstate differ diff --git a/ambient_occlusion/ao.playground/timeline.xctimeline b/ambient_occlusion/ao.playground/timeline.xctimeline new file mode 100644 index 0000000..bf468af --- /dev/null +++ b/ambient_occlusion/ao.playground/timeline.xctimeline @@ -0,0 +1,6 @@ + + + + +