Update heights of plane geometry in a shader kernel

Terrain2/Renderer.swift

@@ -19,6 +19,7 @@ let maxBuffersInFlight = 3
 
 enum RendererError: Error {
     case badVertexDescriptor
+    case badComputeFunction
 }
 
 class Renderer: NSObject, MTKViewDelegate {
@@ -31,6 +32,8 @@ class Renderer: NSObject, MTKViewDelegate {
     var depthState: MTLDepthStencilState
     var colorMap: MTLTexture
 
+    var updateGeometryHeightsPipeline: MTLComputePipelineState
+
     let inFlightSemaphore = DispatchSemaphore(value: maxBuffersInFlight)
     let regenerationSemaphore = DispatchSemaphore(value: 1)
 
@@ -46,6 +49,7 @@ class Renderer: NSObject, MTKViewDelegate {
     var terrain: Terrain
 
     private var iterateTerrainAlgorithm = true
+    private var didUpdateTerrain = false
 
     init?(metalKitView: MTKView) {
         self.device = metalKitView.device!
@@ -94,6 +98,13 @@ class Renderer: NSObject, MTKViewDelegate {
             return nil
         }
 
+        do {
+            updateGeometryHeightsPipeline = try Renderer.buildUpdateGeometryPipeline(withDevice: device, library: library)
+        } catch {
+            print("Unable to create update geometry pipeline. Error: \(error)")
+            return nil
+        }
+
         super.init()
 
     }
@@ -121,6 +132,13 @@ class Renderer: NSObject, MTKViewDelegate {
         return try device.makeRenderPipelineState(descriptor: pipelineDescriptor)
     }
 
+    class func buildUpdateGeometryPipeline(withDevice device: MTLDevice, library: MTLLibrary) throws -> MTLComputePipelineState {
+        guard let computeFunction = library.makeFunction(name: "updateGeometryHeights") else {
+            throw RendererError.badComputeFunction
+        }
+        return try device.makeComputePipelineState(function: computeFunction)
+    }
+
     class func loadTexture(device: MTLDevice,
                            textureName: String) throws -> MTLTexture {
         /// Load texture data with optimal parameters for sampling
@@ -146,6 +164,7 @@ class Renderer: NSObject, MTKViewDelegate {
                 print("Rendering terrain...")
                 self.terrain.generator.render()
                 print("Rendering terrain...complete!")
+                self.didUpdateTerrain = true
             }
         }
         regenerationSemaphore.signal()
@@ -175,7 +194,10 @@ class Renderer: NSObject, MTKViewDelegate {
         let modelMatrix = matrix4x4_rotation(radians: rotation, axis: rotationAxis)
         let viewMatrix = matrix4x4_translation(0.0, -2.0, -8.0)
         uniforms[0].modelViewMatrix = simd_mul(viewMatrix, modelMatrix)
-        rotation += 0.0025
+        rotation += 0.003
+
+        uniforms[0].terrainDimensions = terrain.dimensions
+        uniforms[0].terrainSegments = terrain.segments
     }
 
     func draw(in view: MTKView) {
@@ -192,6 +214,9 @@ class Renderer: NSObject, MTKViewDelegate {
                 if didScheduleAlgorithmIteration && self.iterateTerrainAlgorithm {
                     self.iterateTerrainAlgorithm = false
                 }
+                if self.didUpdateTerrain {
+                    self.didUpdateTerrain = false
+                }
                 regenSem.signal()
                 inFlightSem.signal()
             }
@@ -209,6 +234,16 @@ class Renderer: NSObject, MTKViewDelegate {
                 computeEncoder.endEncoding()
                 didScheduleAlgorithmIteration = true
             }
+
+            if didScheduleAlgorithmIteration || didUpdateTerrain, let computeEncoder = commandBuffer.makeComputeCommandEncoder() {
+                print("Scheduling update geometry iteration")
+                computeEncoder.label = "Geometry Heights Encoder"
+                computeEncoder.pushDebugGroup("Update Geometry: Heights")
+                computeEncoder.setComputePipelineState(updateGeometryHeightsPipeline)
+                computeEncoder.dispatchThreads(MTLSize(width: Int(terrain.segments.x), height: Int(terrain.segments.y), depth: 1), threadsPerThreadgroup: MTLSize(width: 10, height: 10, depth: 1))
+                computeEncoder.popDebugGroup()
+                computeEncoder.endEncoding()
+            }
             
             /// Delay getting the currentRenderPassDescriptor until we absolutely need it to avoid
             ///   holding onto the drawable and blocking the display pipeline any longer than necessary

Terrain2/ShaderTypes.h

@@ -41,16 +41,23 @@ typedef NS_ENUM(NSInteger, TextureIndex)
     TextureIndexColor    = 0,
 };
 
-typedef NS_ENUM(NSInteger, GeneratorTextureIndex)
-{
+typedef NS_ENUM(NSInteger, GeneratorBufferIndex) {
+    GeneratorBufferIndexVertexes = 0,
+    GeneratorBufferIndexTexCoords = 1,
+    GeneratorBufferIndexIndexes = 2,
+    GeneratorBufferIndexUniforms = 3,
+};
+
+typedef NS_ENUM(NSInteger, GeneratorTextureIndex) {
     GeneratorTextureIndexIn  = 0,
     GeneratorTextureIndexOut = 1,
 };
 
-typedef struct
-{
+typedef struct {
     matrix_float4x4 projectionMatrix;
     matrix_float4x4 modelViewMatrix;
+    packed_float2 terrainDimensions;
+    packed_uint2 terrainSegments;
 } Uniforms;
 
 #define kRandomAlgorithmUniforms_RandomCount (41)

Terrain2/Shaders/Shaders.metal

@@ -34,15 +34,14 @@ vertex ColorInOut vertexShader(Vertex in [[stage_in]],
                                texture2d<float> heights [[texture(0)]],
                                constant Uniforms & uniforms [[buffer(BufferIndexUniforms)]])
 {
-    constexpr sampler s(coord::normalized, address::clamp_to_zero, filter::linear);
+//    constexpr sampler s(coord::normalized, address::clamp_to_zero, filter::linear);
 
     ColorInOut out;
 
-    float4 height = heights.sample(s, in.texCoord);
+//    float4 height = heights.sample(s, in.texCoord);
 
     // Replace the y coordinate with the height we read from the texture.
-    float4 position(in.position.x, height.r, in.position.z, 1.0);
-    out.position = uniforms.projectionMatrix * uniforms.modelViewMatrix * position;
+    out.position = uniforms.projectionMatrix * uniforms.modelViewMatrix * float4(in.position, 1.0);
     out.normal = in.normal;
     out.texCoord = in.texCoord;
 

Terrain2/Shaders/TerrainAlgorithms.metal

@@ -43,13 +43,42 @@ private:
     uint mSeed;
 };
 
+kernel void updateGeometryHeights(texture2d<float> texture [[texture(GeneratorTextureIndexIn)]],
+                                  constant float2 *texCoords [[buffer(GeneratorBufferIndexTexCoords)]],
+                                  constant Uniforms &uniforms [[buffer(GeneratorBufferIndexUniforms)]],
+                                  device float3 *vertexes [[buffer(GeneratorBufferIndexVertexes)]],
+                                  uint2 tid [[thread_position_in_grid]])
+{
+    constexpr sampler s(coord::normalized, address::clamp_to_zero, filter::linear);
+
+    const uint vIdx = tid.y * uniforms.terrainSegments.x + tid.x;
+
+    // Get the height from the texture.
+    float2 texCoord = texCoords[vIdx];
+    float4 height = texture.sample(s, texCoord);
+
+    // Update the vertex data.
+    vertexes[vIdx].y = height.r;
+}
+
+kernel void updateGeometryNormals(texture2d<float> texture [[texture(GeneratorTextureIndexIn)]],
+                                  constant float3 *vertexes [[buffer(GeneratorBufferIndexVertexes)]],
+                                  constant float2 *texCoords [[buffer(GeneratorBufferIndexTexCoords)]],
+                                  constant uint *indexes [[buffer(GeneratorBufferIndexIndexes)]],
+                                  constant Uniforms &uniforms [[buffer(GeneratorBufferIndexUniforms)]],
+                                  uint2 tid [[thread_position_in_grid]])
+{
+}
+
+#pragma mark - ZeroGenerator
+
 kernel void zeroKernel(texture2d<float, access::write> outTexture [[texture(GeneratorTextureIndexOut)]],
                        uint2 tid [[thread_position_in_grid]])
 {
     outTexture.write(0, tid);
 }
 
-#pragma mark - RandomAlgorithm
+#pragma mark - RandomGenerator
 
 kernel void randomKernel(texture2d<float, access::write> outTexture [[texture(GeneratorTextureIndexOut)]],
                          constant RandomAlgorithmUniforms &uniforms [[buffer(0)]],