From 11b9cba8ac29b4cf045965572609121996887a7e Mon Sep 17 00:00:00 2001 From: Eryn Wells Date: Wed, 7 Nov 2018 16:47:07 -0500 Subject: [PATCH] Add a Uniform object for RandomAlgorithm This object contains several random numbers for the GPU kernel to use when generating random points. This isn't right yet (there are still patterns in the generated data) but it's time to move on... --- Terrain.xcodeproj/project.pbxproj | 4 ++++ Terrain2/Algorithms.swift | 29 +++++++++++++++++++++++++++-- Terrain2/Renderer.swift | 3 +++ Terrain2/ShaderTypes.h | 10 ++++++++++ Terrain2/TerrainAlgorithms.metal | 10 ++++++---- Terrain2/Uniforms.m | 16 ++++++++++++++++ 6 files changed, 66 insertions(+), 6 deletions(-) create mode 100644 Terrain2/Uniforms.m diff --git a/Terrain.xcodeproj/project.pbxproj b/Terrain.xcodeproj/project.pbxproj index 8d4cd94..fe80c61 100644 --- a/Terrain.xcodeproj/project.pbxproj +++ b/Terrain.xcodeproj/project.pbxproj @@ -7,6 +7,7 @@ objects = { /* Begin PBXBuildFile section */ + C019C8512191CE7100EAD5BB /* Uniforms.m in Sources */ = {isa = PBXBuildFile; fileRef = C019C8502191CE7100EAD5BB /* Uniforms.m */; }; 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 */; }; @@ -29,6 +30,7 @@ /* End PBXBuildFile section */ /* Begin PBXFileReference section */ + C019C8502191CE7100EAD5BB /* Uniforms.m */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.objc; path = Uniforms.m; sourceTree = ""; }; 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; }; @@ -133,6 +135,7 @@ C0C15AB8218E2A90007494E2 /* Shaders.metal */, C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */, C0C15ABA218E2A90007494E2 /* ShaderTypes.h */, + C019C8502191CE7100EAD5BB /* Uniforms.m */, C0C15ABB218E2A90007494E2 /* Assets.xcassets */, C0C15ABD218E2A90007494E2 /* Main.storyboard */, C0C15AC0218E2A90007494E2 /* Info.plist */, @@ -260,6 +263,7 @@ C0C15AC6218E32B3007494E2 /* Terrain.swift in Sources */, C0C15AB7218E2A90007494E2 /* Renderer.swift in Sources */, C0C15AB3218E2A90007494E2 /* AppDelegate.swift in Sources */, + C019C8512191CE7100EAD5BB /* Uniforms.m in Sources */, C08C58A0218F46F000EAFC2D /* Algorithms.swift in Sources */, ); runOnlyForDeploymentPostprocessing = 0; diff --git a/Terrain2/Algorithms.swift b/Terrain2/Algorithms.swift index ea147a5..7bf1321 100644 --- a/Terrain2/Algorithms.swift +++ b/Terrain2/Algorithms.swift @@ -18,6 +18,7 @@ protocol Algorithm { var name: String { get } var outTexture: MTLTexture { get } + func updateUniforms() func encode(in encoder: MTLComputeCommandEncoder) } @@ -33,6 +34,7 @@ class Kernel { let pipeline: MTLComputePipelineState let textures: [MTLTexture] + let uniformBuffer: MTLBuffer? var outTexture: MTLTexture { return textures[textureIndexes.out] @@ -40,7 +42,7 @@ class Kernel { private(set) var textureIndexes: (`in`: Int, out: Int) = (in: 0, out: 1) - init(device: MTLDevice, library: MTLLibrary, functionName: String) throws { + init(device: MTLDevice, library: MTLLibrary, functionName: String, uniformBuffer: MTLBuffer? = nil) throws { guard let computeFunction = library.makeFunction(name: functionName) else { throw KernelError.badFunction } @@ -56,12 +58,15 @@ class Kernel { textures.append(tex) } self.textures = textures + + self.uniformBuffer = uniformBuffer } func encode(in encoder: MTLComputeCommandEncoder) { encoder.setComputePipelineState(pipeline) encoder.setTexture(textures[textureIndexes.in], index: textureIndexes.in) encoder.setTexture(textures[textureIndexes.out], index: textureIndexes.out) + encoder.setBuffer(uniformBuffer, offset: 0, index: 0) encoder.dispatchThreads(Kernel.textureSize, threadsPerThreadgroup: MTLSize(width: 8, height: 8, depth: 1)) } } @@ -78,19 +83,39 @@ class ZeroAlgorithm: Kernel, Algorithm { return nil } } + + // MARK: Algorithm + + func updateUniforms() { } } /// Randomly generate heights that are independent of all others. class RandomAlgorithm: Kernel, Algorithm { let name = "Random" + private var uniforms: UnsafeMutablePointer + init?(device: MTLDevice, library: MTLLibrary) { + let bufferSize = (MemoryLayout.stride & ~0xFF) + 0x100; + guard let buffer = device.makeBuffer(length: bufferSize, options: [.storageModeShared]) else { + print("Couldn't create uniform buffer") + return nil + } + + uniforms = UnsafeMutableRawPointer(buffer.contents()).bindMemory(to: RandomAlgorithmUniforms.self, capacity:1) + do { - try super.init(device: device, library: library, functionName: "randomKernel") + try super.init(device: device, library: library, functionName: "randomKernel", uniformBuffer: buffer) } catch let e { print("Couldn't create compute kernel. Error: \(e)") return nil } + + updateUniforms() + } + + func updateUniforms() { + RandomAlgorithmUniforms_refreshRandoms(uniforms) } } diff --git a/Terrain2/Renderer.swift b/Terrain2/Renderer.swift index 8497345..5cbcd72 100644 --- a/Terrain2/Renderer.swift +++ b/Terrain2/Renderer.swift @@ -151,6 +151,9 @@ class Renderer: NSObject, MTKViewDelegate { private func updateGameState() { /// Update any game state before rendering + if iterateTerrainAlgorithm { + terrain.algorithm.updateUniforms() + } uniforms[0].projectionMatrix = projectionMatrix diff --git a/Terrain2/ShaderTypes.h b/Terrain2/ShaderTypes.h index ff21037..4172884 100644 --- a/Terrain2/ShaderTypes.h +++ b/Terrain2/ShaderTypes.h @@ -53,5 +53,15 @@ typedef struct matrix_float4x4 modelViewMatrix; } Uniforms; +#define kRandomAlgorithmUniforms_RandomCount (41) + +typedef struct { + uint randoms[kRandomAlgorithmUniforms_RandomCount]; +} RandomAlgorithmUniforms; + +#ifndef __METAL_VERSION__ +extern void RandomAlgorithmUniforms_refreshRandoms(RandomAlgorithmUniforms *uniforms); +#endif + #endif /* ShaderTypes_h */ diff --git a/Terrain2/TerrainAlgorithms.metal b/Terrain2/TerrainAlgorithms.metal index 80deb6d..75f1db1 100644 --- a/Terrain2/TerrainAlgorithms.metal +++ b/Terrain2/TerrainAlgorithms.metal @@ -49,12 +49,14 @@ kernel void zeroKernel(texture2d outTexture [[texture(Gene outTexture.write(0, tid); } +#pragma mark - RandomAlgorithm + kernel void randomKernel(texture2d outTexture [[texture(GeneratorTextureIndexOut)]], - uint2 tid [[thread_position_in_grid]], - uint2 tgid [[threadgroup_position_in_grid]]) + constant RandomAlgorithmUniforms &uniforms [[buffer(0)]], + uint2 tid [[thread_position_in_grid]]) { - PRNG rng(tid.x * tid.y); - uint r = rng.lcg(); + PRNG rng(uniforms.randoms[(tid.x * tid.y) % kRandomAlgorithmUniforms_RandomCount]); + uint r = rng.xorShift(); float x = float(r * (1.0 / float(UINT_MAX))) * 0.5f; outTexture.write(x, tid); } diff --git a/Terrain2/Uniforms.m b/Terrain2/Uniforms.m new file mode 100644 index 0000000..0aca68e --- /dev/null +++ b/Terrain2/Uniforms.m @@ -0,0 +1,16 @@ +// +// Uniforms.c +// Terrain2 +// +// Created by Eryn Wells on 11/6/18. +// Copyright © 2018 Eryn Wells. All rights reserved. +// + +#include +#include "ShaderTypes.h" + +void RandomAlgorithmUniforms_refreshRandoms(RandomAlgorithmUniforms *uniforms) { + for (int i = 0; i < kRandomAlgorithmUniforms_RandomCount; i++) { + uniforms->randoms[i] = arc4random(); + } +}