diff --git a/Terrain2/ShaderTypes.h b/Terrain2/ShaderTypes.h
index c7433bb..0b839fe 100644
--- a/Terrain2/ShaderTypes.h
+++ b/Terrain2/ShaderTypes.h
@@ -46,7 +46,8 @@ typedef NS_ENUM(NSInteger, GeneratorBufferIndex) {
     GeneratorBufferIndexTexCoords = 1,
     GeneratorBufferIndexIndexes = 2,
     GeneratorBufferIndexNormals = 3,
-    GeneratorBufferIndexUniforms = 4,
+    GeneratorBufferIndexFaceNormals = 4,
+    GeneratorBufferIndexUniforms = 5,
 };
 
 typedef NS_ENUM(NSInteger, GeneratorTextureIndex) {
diff --git a/Terrain2/Shaders/TerrainAlgorithms.metal b/Terrain2/Shaders/TerrainAlgorithms.metal
index f217aa9..2456946 100644
--- a/Terrain2/Shaders/TerrainAlgorithms.metal
+++ b/Terrain2/Shaders/TerrainAlgorithms.metal
@@ -46,7 +46,7 @@ private:
 kernel void updateGeometryHeights(texture2d<float> texture [[texture(GeneratorTextureIndexIn)]],
                                   constant float2 *texCoords [[buffer(GeneratorBufferIndexTexCoords)]],
                                   constant Uniforms &uniforms [[buffer(GeneratorBufferIndexUniforms)]],
-                                  device packed_float3 *vertexes [[buffer(GeneratorBufferIndexVertexes)]],
+                                  device packed_float3 *vertexes [[buffer(GeneratorBufferIndexMeshPositions)]],
                                   uint2 tid [[thread_position_in_grid]])
 {
     constexpr sampler s(coord::normalized, address::clamp_to_zero, filter::linear);
@@ -61,9 +61,9 @@ kernel void updateGeometryHeights(texture2d<float> texture [[texture(GeneratorTe
     vertexes[vIdx].y = height.r;
 }
 
-kernel void updateGeometryNormals(constant float3 *vertexes [[buffer(GeneratorBufferIndexVertexes)]],
+kernel void updateGeometryNormals(constant float3 *vertexes [[buffer(GeneratorBufferIndexMeshPositions)]],
                                   constant packed_uint3 *indexes [[buffer(GeneratorBufferIndexIndexes)]],
-                                  device packed_float3 *normals [[buffer(GeneratorBufferIndexNormals)]],
+                                  device packed_float3 *normals [[buffer(GeneratorBufferIndexFaceNormals)]],
                                   uint tid [[thread_position_in_grid]])
 {
     const uint3 triIdx = indexes[tid];
diff --git a/Terrain2/Terrain.swift b/Terrain2/Terrain.swift
index 74333d8..7e03d16 100644
--- a/Terrain2/Terrain.swift
+++ b/Terrain2/Terrain.swift
@@ -87,6 +87,7 @@ class Terrain: NSObject {
     let segments: uint2
     let vertexDescriptor: MTLVertexDescriptor
     let mesh: MTKMesh
+    let faceNormalsBuffer: MTLBuffer
 
     var generator: TerrainGenerator
 
@@ -117,6 +118,14 @@ class Terrain: NSObject {
             return nil
         }
 
+        // A normal is a float 3. Two triangles per segment, x * t segments.
+        let faceNormalsLength = MemoryLayout<float3>.stride * 2 * Int(segments.x * segments.y)
+        guard let faceNormalsBuf = device.makeBuffer(length: faceNormalsLength, options: .storageModePrivate) else {
+            print("Couldn't create buffer for face normals")
+            return nil
+        }
+        faceNormalsBuffer = faceNormalsBuf
+
         super.init()
     }
 
@@ -164,9 +173,8 @@ class Terrain: NSObject {
             computeEncoder.setBuffer(indexBuffer.buffer, offset: indexBuffer.offset, index: GeneratorBufferIndex.meshPositions.rawValue)
             let vertexBuffer = mesh.vertexBuffers[BufferIndex.meshPositions.rawValue]
             computeEncoder.setBuffer(vertexBuffer.buffer, offset: vertexBuffer.offset, index: GeneratorBufferIndex.indexes.rawValue)
-//            let normalBuffer = mesh.vertexBuffers[BufferIndex.faceNormals.rawValue]
-//            computeEncoder.setBuffer(normalBuffer.buffer, offset: normalBuffer.offset, index: BufferIndex.faceNormals.rawValue)
-            computeEncoder.dispatchThreads(MTLSize(width: mesh.vertexCount, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: 64, height: 1, depth: 1))
+            computeEncoder.setBuffer(faceNormalsBuffer, offset: 0, index: GeneratorBufferIndex.faceNormals.rawValue)
+            computeEncoder.dispatchThreads(MTLSize(width: 2 * Int(segments.x * segments.y), height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: 64, height: 1, depth: 1))
             computeEncoder.popDebugGroup()
             computeEncoder.endEncoding()
         }