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...
This commit is contained in:
Eryn Wells 2018-11-07 16:47:07 -05:00
parent 76f0065d8b
commit 11b9cba8ac
6 changed files with 66 additions and 6 deletions

View file

@ -7,6 +7,7 @@
objects = { objects = {
/* Begin PBXBuildFile section */ /* 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 */; }; C08C58A0218F46F000EAFC2D /* Algorithms.swift in Sources */ = {isa = PBXBuildFile; fileRef = C08C589F218F46F000EAFC2D /* Algorithms.swift */; };
C08C58A2218F474E00EAFC2D /* TerrainAlgorithms.metal in Sources */ = {isa = PBXBuildFile; fileRef = C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */; }; C08C58A2218F474E00EAFC2D /* TerrainAlgorithms.metal in Sources */ = {isa = PBXBuildFile; fileRef = C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */; };
C0C15A8E218DDD85007494E2 /* AppDelegate.swift in Sources */ = {isa = PBXBuildFile; fileRef = C0C15A8D218DDD85007494E2 /* AppDelegate.swift */; }; C0C15A8E218DDD85007494E2 /* AppDelegate.swift in Sources */ = {isa = PBXBuildFile; fileRef = C0C15A8D218DDD85007494E2 /* AppDelegate.swift */; };
@ -29,6 +30,7 @@
/* End PBXBuildFile section */ /* End PBXBuildFile section */
/* Begin PBXFileReference section */ /* Begin PBXFileReference section */
C019C8502191CE7100EAD5BB /* Uniforms.m */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.objc; path = Uniforms.m; sourceTree = "<group>"; };
C08C589F218F46F000EAFC2D /* Algorithms.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Algorithms.swift; sourceTree = "<group>"; }; C08C589F218F46F000EAFC2D /* Algorithms.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Algorithms.swift; sourceTree = "<group>"; };
C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = TerrainAlgorithms.metal; sourceTree = "<group>"; }; C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = TerrainAlgorithms.metal; sourceTree = "<group>"; };
C0C15A8A218DDD85007494E2 /* Terrain.app */ = {isa = PBXFileReference; explicitFileType = wrapper.application; includeInIndex = 0; path = Terrain.app; sourceTree = BUILT_PRODUCTS_DIR; }; C0C15A8A218DDD85007494E2 /* Terrain.app */ = {isa = PBXFileReference; explicitFileType = wrapper.application; includeInIndex = 0; path = Terrain.app; sourceTree = BUILT_PRODUCTS_DIR; };
@ -133,6 +135,7 @@
C0C15AB8218E2A90007494E2 /* Shaders.metal */, C0C15AB8218E2A90007494E2 /* Shaders.metal */,
C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */, C08C58A1218F474E00EAFC2D /* TerrainAlgorithms.metal */,
C0C15ABA218E2A90007494E2 /* ShaderTypes.h */, C0C15ABA218E2A90007494E2 /* ShaderTypes.h */,
C019C8502191CE7100EAD5BB /* Uniforms.m */,
C0C15ABB218E2A90007494E2 /* Assets.xcassets */, C0C15ABB218E2A90007494E2 /* Assets.xcassets */,
C0C15ABD218E2A90007494E2 /* Main.storyboard */, C0C15ABD218E2A90007494E2 /* Main.storyboard */,
C0C15AC0218E2A90007494E2 /* Info.plist */, C0C15AC0218E2A90007494E2 /* Info.plist */,
@ -260,6 +263,7 @@
C0C15AC6218E32B3007494E2 /* Terrain.swift in Sources */, C0C15AC6218E32B3007494E2 /* Terrain.swift in Sources */,
C0C15AB7218E2A90007494E2 /* Renderer.swift in Sources */, C0C15AB7218E2A90007494E2 /* Renderer.swift in Sources */,
C0C15AB3218E2A90007494E2 /* AppDelegate.swift in Sources */, C0C15AB3218E2A90007494E2 /* AppDelegate.swift in Sources */,
C019C8512191CE7100EAD5BB /* Uniforms.m in Sources */,
C08C58A0218F46F000EAFC2D /* Algorithms.swift in Sources */, C08C58A0218F46F000EAFC2D /* Algorithms.swift in Sources */,
); );
runOnlyForDeploymentPostprocessing = 0; runOnlyForDeploymentPostprocessing = 0;

View file

@ -18,6 +18,7 @@ protocol Algorithm {
var name: String { get } var name: String { get }
var outTexture: MTLTexture { get } var outTexture: MTLTexture { get }
func updateUniforms()
func encode(in encoder: MTLComputeCommandEncoder) func encode(in encoder: MTLComputeCommandEncoder)
} }
@ -33,6 +34,7 @@ class Kernel {
let pipeline: MTLComputePipelineState let pipeline: MTLComputePipelineState
let textures: [MTLTexture] let textures: [MTLTexture]
let uniformBuffer: MTLBuffer?
var outTexture: MTLTexture { var outTexture: MTLTexture {
return textures[textureIndexes.out] return textures[textureIndexes.out]
@ -40,7 +42,7 @@ class Kernel {
private(set) var textureIndexes: (`in`: Int, out: Int) = (in: 0, out: 1) 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 { guard let computeFunction = library.makeFunction(name: functionName) else {
throw KernelError.badFunction throw KernelError.badFunction
} }
@ -56,12 +58,15 @@ class Kernel {
textures.append(tex) textures.append(tex)
} }
self.textures = textures self.textures = textures
self.uniformBuffer = uniformBuffer
} }
func encode(in encoder: MTLComputeCommandEncoder) { func encode(in encoder: MTLComputeCommandEncoder) {
encoder.setComputePipelineState(pipeline) encoder.setComputePipelineState(pipeline)
encoder.setTexture(textures[textureIndexes.in], index: textureIndexes.in) encoder.setTexture(textures[textureIndexes.in], index: textureIndexes.in)
encoder.setTexture(textures[textureIndexes.out], index: textureIndexes.out) 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)) encoder.dispatchThreads(Kernel.textureSize, threadsPerThreadgroup: MTLSize(width: 8, height: 8, depth: 1))
} }
} }
@ -78,19 +83,39 @@ class ZeroAlgorithm: Kernel, Algorithm {
return nil return nil
} }
} }
// MARK: Algorithm
func updateUniforms() { }
} }
/// Randomly generate heights that are independent of all others. /// Randomly generate heights that are independent of all others.
class RandomAlgorithm: Kernel, Algorithm { class RandomAlgorithm: Kernel, Algorithm {
let name = "Random" let name = "Random"
private var uniforms: UnsafeMutablePointer<RandomAlgorithmUniforms>
init?(device: MTLDevice, library: MTLLibrary) { init?(device: MTLDevice, library: MTLLibrary) {
let bufferSize = (MemoryLayout<RandomAlgorithmUniforms>.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 { do {
try super.init(device: device, library: library, functionName: "randomKernel") try super.init(device: device, library: library, functionName: "randomKernel", uniformBuffer: buffer)
} catch let e { } catch let e {
print("Couldn't create compute kernel. Error: \(e)") print("Couldn't create compute kernel. Error: \(e)")
return nil return nil
} }
updateUniforms()
}
func updateUniforms() {
RandomAlgorithmUniforms_refreshRandoms(uniforms)
} }
} }

View file

@ -151,6 +151,9 @@ class Renderer: NSObject, MTKViewDelegate {
private func updateGameState() { private func updateGameState() {
/// Update any game state before rendering /// Update any game state before rendering
if iterateTerrainAlgorithm {
terrain.algorithm.updateUniforms()
}
uniforms[0].projectionMatrix = projectionMatrix uniforms[0].projectionMatrix = projectionMatrix

View file

@ -53,5 +53,15 @@ typedef struct
matrix_float4x4 modelViewMatrix; matrix_float4x4 modelViewMatrix;
} Uniforms; } 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 */ #endif /* ShaderTypes_h */

View file

@ -49,12 +49,14 @@ kernel void zeroKernel(texture2d<float, access::write> outTexture [[texture(Gene
outTexture.write(0, tid); outTexture.write(0, tid);
} }
#pragma mark - RandomAlgorithm
kernel void randomKernel(texture2d<float, access::write> outTexture [[texture(GeneratorTextureIndexOut)]], kernel void randomKernel(texture2d<float, access::write> outTexture [[texture(GeneratorTextureIndexOut)]],
uint2 tid [[thread_position_in_grid]], constant RandomAlgorithmUniforms &uniforms [[buffer(0)]],
uint2 tgid [[threadgroup_position_in_grid]]) uint2 tid [[thread_position_in_grid]])
{ {
PRNG rng(tid.x * tid.y); PRNG rng(uniforms.randoms[(tid.x * tid.y) % kRandomAlgorithmUniforms_RandomCount]);
uint r = rng.lcg(); uint r = rng.xorShift();
float x = float(r * (1.0 / float(UINT_MAX))) * 0.5f; float x = float(r * (1.0 / float(UINT_MAX))) * 0.5f;
outTexture.write(x, tid); outTexture.write(x, tid);
} }

16
Terrain2/Uniforms.m Normal file
View file

@ -0,0 +1,16 @@
//
// Uniforms.c
// Terrain2
//
// Created by Eryn Wells on 11/6/18.
// Copyright © 2018 Eryn Wells. All rights reserved.
//
#include <stdio.h>
#include "ShaderTypes.h"
void RandomAlgorithmUniforms_refreshRandoms(RandomAlgorithmUniforms *uniforms) {
for (int i = 0; i < kRandomAlgorithmUniforms_RandomCount; i++) {
uniforms->randoms[i] = arc4random();
}
}