From 55a134882d98f7b5884200f09f05af30fac40498 Mon Sep 17 00:00:00 2001 From: Eryn Wells Date: Sun, 4 Nov 2018 13:44:47 -0500 Subject: [PATCH] Proof of concept for Algorithm objects Next up compute kernel dispatch! --- Terrain.xcodeproj/project.pbxproj | 8 +++ Terrain2/Algorithms.swift | 104 ++++++++++++++++++++++++++++++ Terrain2/Renderer.swift | 19 ++++-- Terrain2/ShaderTypes.h | 6 ++ Terrain2/Terrain.swift | 34 ++-------- Terrain2/TerrainAlgorithms.metal | 23 +++++++ 6 files changed, 160 insertions(+), 34 deletions(-) create mode 100644 Terrain2/Algorithms.swift create mode 100644 Terrain2/TerrainAlgorithms.metal diff --git a/Terrain.xcodeproj/project.pbxproj b/Terrain.xcodeproj/project.pbxproj index ffd4d83..8d4cd94 100644 --- a/Terrain.xcodeproj/project.pbxproj +++ b/Terrain.xcodeproj/project.pbxproj @@ -7,6 +7,8 @@ objects = { /* Begin PBXBuildFile section */ + C08C58A0218F46F000EAFC2D /* Algorithms.swift in Sources */ = {isa = PBXBuildFile; fileRef = C08C589F218F46F000EAFC2D /* Algorithms.swift */; }; + C08C58A2218F474E00EAFC2D /* TerrainAlgorithms.metal in Sources */ = {isa = PBXBuildFile; fileRef = C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */; }; C0C15A8E218DDD85007494E2 /* AppDelegate.swift in Sources */ = {isa = PBXBuildFile; fileRef = C0C15A8D218DDD85007494E2 /* AppDelegate.swift */; }; C0C15A90218DDD87007494E2 /* Assets.xcassets in Resources */ = {isa = PBXBuildFile; fileRef = C0C15A8F218DDD87007494E2 /* Assets.xcassets */; }; C0C15A93218DDD87007494E2 /* MainMenu.xib in Resources */ = {isa = PBXBuildFile; fileRef = C0C15A91218DDD87007494E2 /* MainMenu.xib */; }; @@ -27,6 +29,8 @@ /* End PBXBuildFile section */ /* Begin PBXFileReference section */ + C08C589F218F46F000EAFC2D /* Algorithms.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Algorithms.swift; sourceTree = ""; }; + C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = TerrainAlgorithms.metal; sourceTree = ""; }; C0C15A8A218DDD85007494E2 /* Terrain.app */ = {isa = PBXFileReference; explicitFileType = wrapper.application; includeInIndex = 0; path = Terrain.app; sourceTree = BUILT_PRODUCTS_DIR; }; C0C15A8D218DDD85007494E2 /* AppDelegate.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = AppDelegate.swift; sourceTree = ""; }; C0C15A8F218DDD87007494E2 /* Assets.xcassets */ = {isa = PBXFileReference; lastKnownFileType = folder.assetcatalog; path = Assets.xcassets; sourceTree = ""; }; @@ -125,7 +129,9 @@ C0C15AB4218E2A90007494E2 /* GameViewController.swift */, C0C15AB6218E2A90007494E2 /* Renderer.swift */, C0C15AC5218E32B2007494E2 /* Terrain.swift */, + C08C589F218F46F000EAFC2D /* Algorithms.swift */, C0C15AB8218E2A90007494E2 /* Shaders.metal */, + C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */, C0C15ABA218E2A90007494E2 /* ShaderTypes.h */, C0C15ABB218E2A90007494E2 /* Assets.xcassets */, C0C15ABD218E2A90007494E2 /* Main.storyboard */, @@ -248,11 +254,13 @@ isa = PBXSourcesBuildPhase; buildActionMask = 2147483647; files = ( + C08C58A2218F474E00EAFC2D /* TerrainAlgorithms.metal in Sources */, C0C15AB9218E2A90007494E2 /* Shaders.metal in Sources */, C0C15AB5218E2A90007494E2 /* GameViewController.swift in Sources */, C0C15AC6218E32B3007494E2 /* Terrain.swift in Sources */, C0C15AB7218E2A90007494E2 /* Renderer.swift in Sources */, C0C15AB3218E2A90007494E2 /* AppDelegate.swift in Sources */, + C08C58A0218F46F000EAFC2D /* Algorithms.swift in Sources */, ); runOnlyForDeploymentPostprocessing = 0; }; diff --git a/Terrain2/Algorithms.swift b/Terrain2/Algorithms.swift new file mode 100644 index 0000000..1391a82 --- /dev/null +++ b/Terrain2/Algorithms.swift @@ -0,0 +1,104 @@ +// +// Algorithms.swift +// Terrain2 +// +// Created by Eryn Wells on 11/4/18. +// Copyright © 2018 Eryn Wells. All rights reserved. +// + +import Foundation +import Metal + +enum KernelError: Error { + case badFunction + case textureCreationFailed +} + +protocol Algorithm { + static var name: String { get } + var outTexture: MTLTexture { get } +} + +class Kernel { + static let textureSize = MTLSize(width: 512, height: 512, depth: 1) + + class func buildTexture(device: MTLDevice, size: MTLSize) -> MTLTexture? { + let desc = MTLTextureDescriptor.texture2DDescriptor(pixelFormat: .r32Float, width: size.width, height: size.height, mipmapped: false) + desc.usage = [.shaderRead, .shaderWrite] + let tex = device.makeTexture(descriptor: desc) + return tex + } + + let pipeline: MTLComputePipelineState + let textures: [MTLTexture] + + var outTexture: MTLTexture { + return textures[indexes.out] + } + + private(set) var indexes: (`in`: Int, out: Int) = (in: 0, out: 1) + + init(device: MTLDevice, library: MTLLibrary, functionName: String) throws { + guard let computeFunction = library.makeFunction(name: functionName) else { + throw KernelError.badFunction + } + self.pipeline = try device.makeComputePipelineState(function: computeFunction) + + // Create our input and output textures + var textures = [MTLTexture]() + for i in 0..<2 { + guard let tex = Kernel.buildTexture(device: device, size: Kernel.textureSize) else { + print("Couldn't create heights texture i=\(i)") + throw KernelError.textureCreationFailed + } + textures.append(tex) + } + self.textures = textures + } + + func encode(in encoder: MTLComputeCommandEncoder, inTexture: MTLTexture?, outTexture: MTLTexture?) { + encoder.setComputePipelineState(pipeline) + encoder.setTexture(inTexture, index: GeneratorTextureIndex.in.rawValue) + encoder.setTexture(inTexture, index: GeneratorTextureIndex.out.rawValue) + encoder.dispatchThreads(Kernel.textureSize, threadsPerThreadgroup: MTLSize(width: 8, height: 8, depth: 1)) + } +} + +/// "Compute" zero for every value of the height map. +class ZeroAlgorithm: Kernel, Algorithm { + static let name = "Zero" + + init?(device: MTLDevice, library: MTLLibrary) { + do { + try super.init(device: device, library: library, functionName: "zeroKernel") + } catch let e { + print("Couldn't create compute kernel. Error: \(e)") + return nil + } + } +} + +/// Randomly generate heights that are independent of all others. +class RandomAlgorithm: Kernel, Algorithm { + static let name = "Random" + + init?(device: MTLDevice, library: MTLLibrary) { + do { + try super.init(device: device, library: library, functionName: "randomKernel") + } catch let e { + print("Couldn't create compute kernel. Error: \(e)") + return nil + } + } +} + +/// Implementation of the Diamond-Squares algorithm. +/// - https://en.wikipedia.org/wiki/Diamond-square_algorithm +//class DiamondSquareAlgorithm: Algorithm { +// static let name = "Diamond-Square" +//} + +/// Implementation of the Circles algorithm. +//class CirclesAlgorithm: Algorithm { +// static let name = "Circles" +//} diff --git a/Terrain2/Renderer.swift b/Terrain2/Renderer.swift index 6838a3b..18f4e79 100644 --- a/Terrain2/Renderer.swift +++ b/Terrain2/Renderer.swift @@ -24,6 +24,7 @@ enum RendererError: Error { class Renderer: NSObject, MTKViewDelegate { public let device: MTLDevice + let library: MTLLibrary let commandQueue: MTLCommandQueue var dynamicUniformBuffer: MTLBuffer var pipelineState: MTLRenderPipelineState @@ -61,10 +62,17 @@ class Renderer: NSObject, MTKViewDelegate { metalKitView.colorPixelFormat = MTLPixelFormat.bgra8Unorm_srgb metalKitView.sampleCount = 1 - terrain = Terrain(dimensions: float2(10, 10), segments: uint2(100, 100), device: device)! + guard let library = device.makeDefaultLibrary() else { + print("Unable to create default library") + return nil + } + self.library = library + + terrain = Terrain(dimensions: float2(10, 10), segments: uint2(100, 100), device: device, library: library)! do { pipelineState = try Renderer.buildRenderPipelineWithDevice(device: device, + library: library, metalKitView: metalKitView, mtlVertexDescriptor: terrain.vertexDescriptor) } catch { @@ -89,14 +97,13 @@ class Renderer: NSObject, MTKViewDelegate { } class func buildRenderPipelineWithDevice(device: MTLDevice, + library: MTLLibrary, metalKitView: MTKView, mtlVertexDescriptor: MTLVertexDescriptor) throws -> MTLRenderPipelineState { /// Build a render state pipeline object - let library = device.makeDefaultLibrary() - - let vertexFunction = library?.makeFunction(name: "vertexShader") - let fragmentFunction = library?.makeFunction(name: "fragmentShader") + let vertexFunction = library.makeFunction(name: "vertexShader") + let fragmentFunction = library.makeFunction(name: "fragmentShader") let pipelineDescriptor = MTLRenderPipelineDescriptor() pipelineDescriptor.label = "RenderPipeline" @@ -203,7 +210,7 @@ class Renderer: NSObject, MTKViewDelegate { } } - renderEncoder.setVertexTexture(terrain.heightMap, index: 0) + renderEncoder.setVertexTexture(terrain.algorithm.outTexture, index: 0) renderEncoder.setFragmentTexture(colorMap, index: TextureIndex.color.rawValue) for submesh in terrain.mesh.submeshes { diff --git a/Terrain2/ShaderTypes.h b/Terrain2/ShaderTypes.h index eff3ef0..ff21037 100644 --- a/Terrain2/ShaderTypes.h +++ b/Terrain2/ShaderTypes.h @@ -41,6 +41,12 @@ typedef NS_ENUM(NSInteger, TextureIndex) TextureIndexColor = 0, }; +typedef NS_ENUM(NSInteger, GeneratorTextureIndex) +{ + GeneratorTextureIndexIn = 0, + GeneratorTextureIndexOut = 1, +}; + typedef struct { matrix_float4x4 projectionMatrix; diff --git a/Terrain2/Terrain.swift b/Terrain2/Terrain.swift index 12ee6eb..b412fd4 100644 --- a/Terrain2/Terrain.swift +++ b/Terrain2/Terrain.swift @@ -72,36 +72,14 @@ class Terrain: NSObject { private static let heightMapSize = MTLSize(width: 512, height: 512, depth: 1) - class func buildHeightsTexture(device: MTLDevice) -> MTLTexture? { - let heightsDesc = MTLTextureDescriptor.texture2DDescriptor(pixelFormat: .r32Float, width: heightMapSize.width, height: heightMapSize.height, mipmapped: false) - heightsDesc.usage = [.shaderRead, .shaderWrite] - - let tex = device.makeTexture(descriptor: heightsDesc) - - if let tex = tex { - var initialHeights = [Float]() - let numberOfHeights = tex.height * tex.width - initialHeights.reserveCapacity(numberOfHeights) - for _ in 0...stride * tex.width - tex.replace(region: region, mipmapLevel: 0, withBytes: initialHeights, bytesPerRow: bytesPerRow) - } - - return tex - } - let dimensions: float2 let segments: uint2 let vertexDescriptor: MTLVertexDescriptor let mesh: MTKMesh - let heightMap: MTLTexture - init?(dimensions dim: float2, segments seg: uint2, device: MTLDevice) { + var algorithm: Algorithm + + init?(dimensions dim: float2, segments seg: uint2, device: MTLDevice, library: MTLLibrary) { dimensions = dim segments = seg vertexDescriptor = Terrain.buildVertexDescriptor() @@ -113,11 +91,11 @@ class Terrain: NSObject { return nil } - guard let tex = Terrain.buildHeightsTexture(device: device) else { - print("Couldn't create heights texture") + guard let alg = ZeroAlgorithm(device: device, library: library) else { + print("Couldn't create algorithm") return nil } - heightMap = tex + algorithm = alg super.init() } diff --git a/Terrain2/TerrainAlgorithms.metal b/Terrain2/TerrainAlgorithms.metal new file mode 100644 index 0000000..1c9ce3f --- /dev/null +++ b/Terrain2/TerrainAlgorithms.metal @@ -0,0 +1,23 @@ +// +// TerrainAlgorithms.metal +// Terrain2 +// +// Created by Eryn Wells on 11/4/18. +// Copyright © 2018 Eryn Wells. All rights reserved. +// + +#include +#include "ShaderTypes.h" +using namespace metal; + +kernel void zeroKernel(texture2d outTexture [[texture(GeneratorTextureIndexOut)]], + uint2 tid [[thread_position_in_grid]]) +{ + outTexture.write(0, tid); +} + +kernel void randomKernel(texture2d outTexture [[texture(GeneratorTextureIndexOut)]], + uint2 tid [[thread_position_in_grid]]) +{ + +}