Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
17 changes: 16 additions & 1 deletion Sources/CShaderTypes/ShaderTypes.h
Original file line number Diff line number Diff line change
Expand Up @@ -445,10 +445,12 @@ typedef enum {
} ARTextureIndices;

typedef enum{
gaussianSplatIndex,
gaussianEncodedSplatIndex,
gaussianUniformIndex,
gaussianNumberOfSplatsIndex,
gaussianIndicesIndex,
gaussianVisibleIndicesIndex,
gaussianVisibleCountIndex,
}GaussianDepthBufferIndices;

typedef struct{
Expand All @@ -470,12 +472,25 @@ typedef struct{
float _pad2;
}EncodedGaussianSplat;

// Higher-order SH coefficients are stored in a separate packed half buffer.
// Layout per splat: R[1...n], G[1...n], B[1...n]. The DC coefficient remains
// represented by EncodedGaussianSplat.color.
typedef struct{
uint degree;
uint coefficientsPerChannel;
uint higherOrderCoefficientsPerSplat;
uint _pad0;
}GaussianSHMetadata;

typedef enum{
gaussianTBDRRenderIndicesIndex = 0,
gaussianTBDRRenderSplatIndex,
gaussianTBDRRenderUniformIndex,
gaussianTBDRRenderViewPortIndex,
gaussianTBDRRenderReverseZIndex,
gaussianTBDRRenderSHIndex,
gaussianTBDRRenderSHMetadataIndex,
gaussianTBDRRenderLocalCameraIndex,
}GaussianTBDRRenderBufferIndices;

typedef enum{
Expand Down
7 changes: 6 additions & 1 deletion Sources/UntoldEngine/ECS/Components.swift
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@
// License, v. 2.0. If a copy of the MPL was not distributed with this
// file, You can obtain one at https://mozilla.org/MPL/2.0/.

import CShaderTypes
import Foundation
import Metal
import MetalKit
Expand Down Expand Up @@ -89,9 +90,13 @@ public class EntitySceneChannelsComponent: Component {
}

public class GaussianComponent: Component {
var splatData: MTLBuffer?
var encodedSplatData: MTLBuffer?
var sphericalHarmonicsData: MTLBuffer?
var sphericalHarmonicsMetadata: GaussianSHMetadata?
var gaussianSortedIndices: MTLBuffer?
var gaussianVisibleIndices: MTLBuffer?
var gaussianVisibleCount: MTLBuffer?
var visibleSplatCountForRendering: UInt = 0
public var spaceUniform: [MTLBuffer?] = Array(repeating: nil, count: totalPerMeshUniformBuffers())
var splatCount: UInt = 0

Expand Down
112 changes: 112 additions & 0 deletions Sources/UntoldEngine/Profiling/GaussianProfiling.swift
Original file line number Diff line number Diff line change
@@ -0,0 +1,112 @@
//
// GaussianProfiling.swift
// UntoldEngine
//
// Copyright (C) Untold Engine Studios
//
// This Source Code Form is subject to the terms of the Mozilla Public
// License, v. 2.0. If a copy of the MPL was not distributed with this
// file, You can obtain one at https://mozilla.org/MPL/2.0/.

import Foundation
import QuartzCore

struct GaussianProfileTotals {
var entityCount: Int = 0
var splatCount: Int = 0
var drawCallCount: Int = 0
var dispatchCount: Int = 0
var radixPassCount: Int = 0
var encodedBytes: Int = 0
var sortedIndexBytes: Int = 0
var visibleIndexBytes: Int = 0
var visibleCountBytes: Int = 0
var sphericalHarmonicsBytes: Int = 0
var uniformBytes: Int = 0
var scratchBytes: Int = 0
var maxSphericalHarmonicsDegree: UInt32 = 0
var higherOrderCoefficientsPerSplat: UInt32 = 0

var totalResidentBytes: Int {
encodedBytes + sortedIndexBytes + visibleIndexBytes + visibleCountBytes + sphericalHarmonicsBytes + uniformBytes + scratchBytes
}

mutating func include(component: GaussianComponent) {
entityCount += 1
splatCount += Int(component.splatCount)
encodedBytes += component.encodedSplatData?.length ?? 0
sortedIndexBytes += component.gaussianSortedIndices?.length ?? 0
visibleIndexBytes += component.gaussianVisibleIndices?.length ?? 0
visibleCountBytes += component.gaussianVisibleCount?.length ?? 0
if let shBuffer = component.sphericalHarmonicsData {
sphericalHarmonicsBytes += shBuffer.length
}
uniformBytes += component.spaceUniform.reduce(0) { total, buffer in
total + (buffer?.length ?? 0)
}
if let metadata = component.sphericalHarmonicsMetadata {
maxSphericalHarmonicsDegree = max(maxSphericalHarmonicsDegree, metadata.degree)
higherOrderCoefficientsPerSplat = max(
higherOrderCoefficientsPerSplat,
metadata.higherOrderCoefficientsPerSplat
)
}
}
}

@inline(__always)
func gaussianProfilingStartTime() -> CFTimeInterval? {
guard Logger.isEnabled(category: .gaussian) else { return nil }
return CACurrentMediaTime()
}

func logGaussianProfile(
stage: String,
startTime: CFTimeInterval?,
totals: GaussianProfileTotals,
extra: String = ""
) {
guard let startTime else { return }

let elapsedMs = (CACurrentMediaTime() - startTime) * 1000.0
let bytesPerSplat = totals.splatCount > 0
? Double(totals.totalResidentBytes) / Double(totals.splatCount)
: 0.0
let suffix = extra.isEmpty ? "" : " \(extra)"

Logger.log(
message: String(
format: "[Gaussian][%@] cpuEncodeMs=%.3f entities=%d splats=%d draws=%d dispatches=%d radixPasses=%d shDegree=%u shRestCoeffsPerSplat=%u memory=%@ bytesPerSplat=%.1f encoded=%@ sorted=%@ visible=%@ sh=%@ uniforms=%@ scratch=%@%@",
stage,
elapsedMs,
totals.entityCount,
totals.splatCount,
totals.drawCallCount,
totals.dispatchCount,
totals.radixPassCount,
totals.maxSphericalHarmonicsDegree,
totals.higherOrderCoefficientsPerSplat,
gaussianFormatBytes(totals.totalResidentBytes),
bytesPerSplat,
gaussianFormatBytes(totals.encodedBytes),
gaussianFormatBytes(totals.sortedIndexBytes),
gaussianFormatBytes(totals.visibleIndexBytes + totals.visibleCountBytes),
gaussianFormatBytes(totals.sphericalHarmonicsBytes),
gaussianFormatBytes(totals.uniformBytes),
gaussianFormatBytes(totals.scratchBytes),
suffix
),
category: LogCategory.gaussian.rawValue
)
}

func gaussianFormatBytes(_ bytes: Int) -> String {
let value = Double(bytes)
if bytes >= 1024 * 1024 {
return String(format: "%.2fMiB", value / 1_048_576.0)
}
if bytes >= 1024 {
return String(format: "%.2fKiB", value / 1024.0)
}
return "\(bytes)B"
}
46 changes: 45 additions & 1 deletion Sources/UntoldEngine/Renderer/RenderPasses.swift
Original file line number Diff line number Diff line change
Expand Up @@ -3993,6 +3993,10 @@ public enum RenderPasses {
}

public static let gaussianExecution: RenderPassExecution = { commandBuffer in
let profileStart = gaussianProfilingStartTime()
var profileTotals = GaussianProfileTotals()
var activeSplatTotal = 0

let pipelines = PipelineManager.shared.renderPipelinesByType
guard let initializePipeline = pipelines[.gaussianTBDRInitialize],
let drawPipeline = pipelines[.gaussianTBDRDraw],
Expand Down Expand Up @@ -4046,6 +4050,7 @@ public enum RenderPasses {
renderEncoder.pushDebugGroup("Initialize")
renderEncoder.setRenderPipelineState(initializePipelineState)
renderEncoder.dispatchThreadsPerTile(MTLSize(width: 32, height: 32, depth: 1))
profileTotals.dispatchCount += 1
renderEncoder.popDebugGroup()

renderEncoder.pushDebugGroup("Draw Splats")
Expand All @@ -4063,6 +4068,11 @@ public enum RenderPasses {
handleError(.noGaussianComponent, entityId)
continue
}
profileTotals.include(component: gaussianComponent)

let activeSplatCount = min(Int(gaussianComponent.visibleSplatCountForRendering), Int(gaussianComponent.splatCount))
guard activeSplatCount > 0 else { continue }
activeSplatTotal += activeSplatCount

guard gaussianComponent.encodedSplatData != nil else {
handleError(.bufferAllocationFailed, "Encoded Gaussian splat buffer")
Expand Down Expand Up @@ -4135,10 +4145,37 @@ public enum RenderPasses {
)
renderEncoder.setVertexBytes(&renderInfo.viewPort, length: MemoryLayout<simd_float2>.stride, index: Int(gaussianTBDRRenderViewPortIndex.rawValue))

var shMetadata = gaussianComponent.sphericalHarmonicsMetadata ?? GaussianSHMetadata(
degree: 0,
coefficientsPerChannel: 0,
higherOrderCoefficientsPerSplat: 0,
_pad0: 0
)
renderEncoder.setVertexBuffer(
gaussianComponent.sphericalHarmonicsData ?? gaussianComponent.encodedSplatData,
offset: 0,
index: Int(gaussianTBDRRenderSHIndex.rawValue)
)
renderEncoder.setVertexBytes(
&shMetadata,
length: MemoryLayout<GaussianSHMetadata>.stride,
index: Int(gaussianTBDRRenderSHMetadataIndex.rawValue)
)
var localCameraPosition = gaussianLocalCameraPosition(
cameraWorldPosition: effectiveCameraPosition,
modelMatrix: modelMatrix
)
renderEncoder.setVertexBytes(
&localCameraPosition,
length: MemoryLayout<simd_float3>.stride,
index: Int(gaussianTBDRRenderLocalCameraIndex.rawValue)
)

renderEncoder.drawPrimitivesTracked(type: .triangleStrip,
vertexStart: 0,
vertexCount: 4,
instanceCount: Int(gaussianComponent.splatCount))
instanceCount: activeSplatCount)
profileTotals.drawCallCount += 1
}

renderEncoder.popDebugGroup()
Expand All @@ -4149,11 +4186,18 @@ public enum RenderPasses {
var reverseZ = renderInfo.reverseZEnabled
renderEncoder.setFragmentBytes(&reverseZ, length: MemoryLayout<Bool>.stride, index: Int(gaussianTBDRRenderReverseZIndex.rawValue))
renderEncoder.drawPrimitivesTracked(type: .triangle, vertexStart: 0, vertexCount: 3)
profileTotals.drawCallCount += 1
renderEncoder.popDebugGroup()

renderEncoder.updateFence(renderInfo.fence, after: .fragment)
renderEncoder.popDebugGroup()
renderEncoder.endEncoding()
logGaussianProfile(
stage: "Render",
startTime: profileStart,
totals: profileTotals,
extra: String(format: "activeSplats=%d viewport=%.0fx%.0f tile=32x32", activeSplatTotal, renderInfo.viewPort.x, renderInfo.viewPort.y)
)
}

static func executePostProcess(
Expand Down
52 changes: 48 additions & 4 deletions Sources/UntoldEngine/Shaders/BitonicSort.metal
Original file line number Diff line number Diff line change
Expand Up @@ -38,17 +38,62 @@ inline float eye_space_depth(const float4x4 modelView, float3 worldPos) {
return max(-v.z, 0.0f);
}

kernel void gaussianResetVisibleCount(
device atomic_uint *visibleCount [[buffer(gaussianVisibleCountIndex)]],
uint index [[thread_position_in_grid]])
{
if (index == 0u) {
atomic_store_explicit(visibleCount, 0u, memory_order_relaxed);
}
}

kernel void gaussianFrustumCull(
const device EncodedGaussianSplat *splats [[buffer(gaussianEncodedSplatIndex)]],
constant Uniforms &uniforms [[buffer(gaussianUniformIndex)]],
constant uint &numOfSplats [[buffer(gaussianNumberOfSplatsIndex)]],
constant float &clipGuardBand [[buffer(gaussianIndicesIndex)]],
device uint *visibleIndices [[buffer(gaussianVisibleIndicesIndex)]],
device atomic_uint *visibleCount [[buffer(gaussianVisibleCountIndex)]],
uint index [[thread_position_in_grid]])
{
if (index >= numOfSplats) return;

float4 centerClip = uniforms.projectionMatrix *
uniforms.modelViewMatrix *
float4(splats[index].position, 1.0f);

if (centerClip.w <= 0.0f) return;

float limit = max(0.0f, 1.0f + clipGuardBand);
float2 ndc = centerClip.xy / centerClip.w;
if (abs(ndc.x) > limit || abs(ndc.y) > limit) return;
if (centerClip.z < -centerClip.w * clipGuardBand || centerClip.z > centerClip.w * limit) return;

uint writeIndex = atomic_fetch_add_explicit(visibleCount, 1u, memory_order_relaxed);
visibleIndices[writeIndex] = index;
}

kernel void gaussianDepthKeys(device uint64_t *outKeys [[buffer(gaussianIndicesIndex)]],
const device GaussianSplat *splats [[buffer(gaussianSplatIndex)]],
const device EncodedGaussianSplat *splats [[buffer(gaussianEncodedSplatIndex)]],
constant uint &numOfSplats [[buffer(gaussianNumberOfSplatsIndex)]],
constant Uniforms &uniforms [[buffer(gaussianUniformIndex)]],
const device uint *visibleIndices [[buffer(gaussianVisibleIndicesIndex)]],
const device uint *visibleCount [[buffer(gaussianVisibleCountIndex)]],
uint index [[thread_position_in_grid]],
uint gridSize [[threads_per_grid]])
{
if (index >= numOfSplats) return;

if (index >= visibleCount[0]) {
uint64_t invalidPacked = ((uint64_t)0xffffffffu << 32) | (uint64_t)0xffffffffu;
outKeys[index] = invalidPacked;
return;
}

uint splatIndex = visibleIndices[index];

// 1) Eye-space depth
float z = eye_space_depth(uniforms.modelViewMatrix, splats[index].center.xyz);
float z = eye_space_depth(uniforms.modelViewMatrix, splats[splatIndex].position);

// 2) Convert to lexicographically sortable unsigned int
uint keyFrontToBack = float_to_sortable_u32(z);
Expand All @@ -57,7 +102,6 @@ kernel void gaussianDepthKeys(device uint64_t *outKeys [[buffer(gau
//uint keyBackToFront = 0xffffffffu - keyFrontToBack;

// 3) Pack [key | index] for 64-bit radix sort (key in high bits)
uint64_t packed = (uint64_t)keyFrontToBack << 32 | (uint64_t)index;
uint64_t packed = (uint64_t)keyFrontToBack << 32 | (uint64_t)splatIndex;
outKeys[index] = packed;
}

Loading
Loading