chore: sync workspace state

This commit is contained in:
2026-03-29 01:36:53 +08:00
parent eb5de3e3d4
commit e5cb79f3ce
4935 changed files with 35593 additions and 360696 deletions

View File

@@ -0,0 +1,8 @@
fileFormatVersion: 2
guid: 07d8a5410ba18f64e9efe04f3a023cfa
folderAsset: yes
DefaultImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,35 @@
// SPDX-License-Identifier: MIT
using GaussianSplatting.Runtime;
using UnityEditor;
using UnityEditor.EditorTools;
using UnityEngine;
namespace GaussianSplatting.Editor
{
[EditorTool("Gaussian Move Tool", typeof(GaussianSplatRenderer), typeof(GaussianToolContext))]
class GaussianMoveTool : GaussianTool
{
public override void OnToolGUI(EditorWindow window)
{
var gs = GetRenderer();
if (!gs || !CanBeEdited() || !HasSelection())
return;
var tr = gs.transform;
EditorGUI.BeginChangeCheck();
var selCenterLocal = GetSelectionCenterLocal();
var selCenterWorld = tr.TransformPoint(selCenterLocal);
var newPosWorld = Handles.DoPositionHandle(selCenterWorld, Tools.handleRotation);
if (EditorGUI.EndChangeCheck())
{
var newPosLocal = tr.InverseTransformPoint(newPosWorld);
var wasModified = gs.editModified;
gs.EditTranslateSelection(newPosLocal - selCenterLocal);
if (!wasModified)
GaussianSplatRendererEditor.RepaintAll();
Event.current.Use();
}
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 9c9f40b54eb504648b2a0beadabbcc8d
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,70 @@
// SPDX-License-Identifier: MIT
using GaussianSplatting.Runtime;
using UnityEditor;
using UnityEditor.EditorTools;
using UnityEngine;
namespace GaussianSplatting.Editor
{
/* not working correctly yet
[EditorTool("Gaussian Rotate Tool", typeof(GaussianSplatRenderer), typeof(GaussianToolContext))]
class GaussianRotateTool : GaussianTool
{
Quaternion m_CurrentRotation = Quaternion.identity;
Vector3 m_FrozenSelCenterLocal = Vector3.zero;
bool m_FreezePivot = false;
public override void OnActivated()
{
m_FreezePivot = false;
}
public override void OnToolGUI(EditorWindow window)
{
var gs = GetRenderer();
if (!gs || !CanBeEdited() || !HasSelection())
return;
var tr = gs.transform;
var evt = Event.current;
var selCenterLocal = GetSelectionCenterLocal();
if (evt.type == EventType.MouseDown)
{
gs.EditStorePosMouseDown();
gs.EditStoreOtherMouseDown();
m_FrozenSelCenterLocal = selCenterLocal;
m_FreezePivot = true;
}
if (evt.type == EventType.MouseUp)
{
m_CurrentRotation = Quaternion.identity;
m_FreezePivot = false;
}
if (m_FreezePivot)
selCenterLocal = m_FrozenSelCenterLocal;
EditorGUI.BeginChangeCheck();
var selCenterWorld = tr.TransformPoint(selCenterLocal);
var newRotation = Handles.DoRotationHandle(m_CurrentRotation, selCenterWorld);
if (EditorGUI.EndChangeCheck())
{
Matrix4x4 localToWorld = gs.transform.localToWorldMatrix;
Matrix4x4 worldToLocal = gs.transform.worldToLocalMatrix;
var wasModified = gs.editModified;
var rotToApply = newRotation;
gs.EditRotateSelection(selCenterLocal, localToWorld, worldToLocal, rotToApply);
m_CurrentRotation = newRotation;
if (!wasModified)
GaussianSplatRendererEditor.RepaintAll();
if(GUIUtility.hotControl == 0)
{
m_CurrentRotation = Tools.handleRotation;
}
}
}
}
*/
}

View File

@@ -0,0 +1,3 @@
fileFormatVersion: 2
guid: 5128238188a44c86914a22a862195242
timeCreated: 1697805149

View File

@@ -0,0 +1,68 @@
// SPDX-License-Identifier: MIT
using GaussianSplatting.Runtime;
using UnityEditor;
using UnityEditor.EditorTools;
using UnityEngine;
namespace GaussianSplatting.Editor
{
/* // not working correctly yet when the GS itself has scale
[EditorTool("Gaussian Scale Tool", typeof(GaussianSplatRenderer), typeof(GaussianToolContext))]
class GaussianScaleTool : GaussianTool
{
Vector3 m_CurrentScale = Vector3.one;
Vector3 m_FrozenSelCenterLocal = Vector3.zero;
bool m_FreezePivot = false;
public override void OnActivated()
{
m_FreezePivot = false;
}
public override void OnToolGUI(EditorWindow window)
{
var gs = GetRenderer();
if (!gs || !CanBeEdited() || !HasSelection())
return;
var tr = gs.transform;
var evt = Event.current;
var selCenterLocal = GetSelectionCenterLocal();
if (evt.type == EventType.MouseDown)
{
gs.EditStorePosMouseDown();
m_FrozenSelCenterLocal = selCenterLocal;
m_FreezePivot = true;
}
if (evt.type == EventType.MouseUp)
{
m_CurrentScale = Vector3.one;
m_FreezePivot = false;
}
if (m_FreezePivot)
selCenterLocal = m_FrozenSelCenterLocal;
EditorGUI.BeginChangeCheck();
var selCenterWorld = tr.TransformPoint(selCenterLocal);
m_CurrentScale = Handles.DoScaleHandle(m_CurrentScale, selCenterWorld, Tools.handleRotation, HandleUtility.GetHandleSize(selCenterWorld));
if (EditorGUI.EndChangeCheck())
{
Matrix4x4 localToWorld = Matrix4x4.identity;
Matrix4x4 worldToLocal = Matrix4x4.identity;
if (Tools.pivotRotation == PivotRotation.Global)
{
localToWorld = gs.transform.localToWorldMatrix;
worldToLocal = gs.transform.worldToLocalMatrix;
}
var wasModified = gs.editModified;
gs.EditScaleSelection(selCenterLocal, localToWorld, worldToLocal, m_CurrentScale);
if (!wasModified)
GaussianSplatRendererEditor.RepaintAll();
evt.Use();
}
}
}
*/
}

View File

@@ -0,0 +1,3 @@
fileFormatVersion: 2
guid: dbf3d17a31b942b28f5d8c187adb8fdf
timeCreated: 1697732813

File diff suppressed because it is too large Load Diff

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 635bd950b8a74c84f870d5c8f02c3974
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,71 @@
// SPDX-License-Identifier: MIT
using GaussianSplatting.Runtime;
using Unity.Collections.LowLevel.Unsafe;
using UnityEditor;
using UnityEngine;
namespace GaussianSplatting.Editor
{
[CustomEditor(typeof(GaussianSplatAsset))]
[CanEditMultipleObjects]
public class GaussianSplatAssetEditor : UnityEditor.Editor
{
public override void OnInspectorGUI()
{
var gs = target as GaussianSplatAsset;
if (!gs)
return;
using var _ = new EditorGUI.DisabledScope(true);
if (targets.Length == 1)
SingleAssetGUI(gs);
else
{
int totalCount = 0;
foreach (var tgt in targets)
{
var gss = tgt as GaussianSplatAsset;
if (gss)
{
totalCount += gss.splatCount;
}
}
EditorGUILayout.TextField("Total Splats", $"{totalCount:N0}");
}
}
static void SingleAssetGUI(GaussianSplatAsset gs)
{
var splatCount = gs.splatCount;
EditorGUILayout.TextField("Splats", $"{splatCount:N0}");
var prevBackColor = GUI.backgroundColor;
if (gs.formatVersion != GaussianSplatAsset.kCurrentVersion)
GUI.backgroundColor *= Color.red;
EditorGUILayout.IntField("Version", gs.formatVersion);
GUI.backgroundColor = prevBackColor;
long sizePos = gs.posData != null ? gs.posData.dataSize : 0;
long sizeOther = gs.otherData != null ? gs.otherData.dataSize : 0;
long sizeCol = gs.colorData != null ? gs.colorData.dataSize : 0;
long sizeSH = GaussianSplatAsset.CalcSHDataSize(gs.splatCount, gs.shFormat);
long sizeChunk = gs.chunkData != null ? gs.chunkData.dataSize : 0;
EditorGUILayout.TextField("Memory", EditorUtility.FormatBytes(sizePos + sizeOther + sizeSH + sizeCol + sizeChunk));
EditorGUI.indentLevel++;
EditorGUILayout.TextField("Positions", $"{EditorUtility.FormatBytes(sizePos)} ({gs.posFormat})");
EditorGUILayout.TextField("Other", $"{EditorUtility.FormatBytes(sizeOther)} ({gs.scaleFormat})");
EditorGUILayout.TextField("Base color", $"{EditorUtility.FormatBytes(sizeCol)} ({gs.colorFormat})");
EditorGUILayout.TextField("SHs", $"{EditorUtility.FormatBytes(sizeSH)} ({gs.shFormat})");
EditorGUILayout.TextField("Chunks",
$"{EditorUtility.FormatBytes(sizeChunk)} ({UnsafeUtility.SizeOf<GaussianSplatAsset.ChunkInfo>()} B/chunk)");
EditorGUI.indentLevel--;
EditorGUILayout.Vector3Field("Bounds Min", gs.boundsMin);
EditorGUILayout.Vector3Field("Bounds Max", gs.boundsMax);
EditorGUILayout.TextField("Data Hash", gs.dataHash.ToString());
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 75971a29a6deda14c9b1ff5f4ab2f2a0
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,444 @@
// SPDX-License-Identifier: MIT
using System;
using System.Collections.Generic;
using System.IO;
using System.Text;
using GaussianSplatting.Runtime;
using Unity.Collections.LowLevel.Unsafe;
using Unity.Mathematics;
using UnityEditor;
using UnityEditor.EditorTools;
using UnityEngine;
using GaussianSplatRenderer = GaussianSplatting.Runtime.GaussianSplatRenderer;
namespace GaussianSplatting.Editor
{
[CustomEditor(typeof(GaussianSplatRenderer))]
[CanEditMultipleObjects]
public class GaussianSplatRendererEditor : UnityEditor.Editor
{
const string kPrefExportBake = "nesnausk.GaussianSplatting.ExportBakeTransform";
SerializedProperty m_PropAsset;
SerializedProperty m_PropSplatScale;
SerializedProperty m_PropOpacityScale;
SerializedProperty m_PropSHOrder;
SerializedProperty m_PropSHOnly;
SerializedProperty m_PropSortNthFrame;
SerializedProperty m_PropRenderMode;
SerializedProperty m_PropPointDisplaySize;
SerializedProperty m_PropCutouts;
SerializedProperty m_PropShaderSplats;
SerializedProperty m_PropShaderComposite;
SerializedProperty m_PropShaderDebugPoints;
SerializedProperty m_PropShaderDebugBoxes;
SerializedProperty m_PropCSSplatUtilities;
bool m_ResourcesExpanded = false;
int m_CameraIndex = 0;
bool m_ExportBakeTransform;
static int s_EditStatsUpdateCounter = 0;
static HashSet<GaussianSplatRendererEditor> s_AllEditors = new();
public static void BumpGUICounter()
{
++s_EditStatsUpdateCounter;
}
public static void RepaintAll()
{
foreach (var e in s_AllEditors)
e.Repaint();
}
public void OnEnable()
{
m_ExportBakeTransform = EditorPrefs.GetBool(kPrefExportBake, false);
m_PropAsset = serializedObject.FindProperty("m_Asset");
m_PropSplatScale = serializedObject.FindProperty("m_SplatScale");
m_PropOpacityScale = serializedObject.FindProperty("m_OpacityScale");
m_PropSHOrder = serializedObject.FindProperty("m_SHOrder");
m_PropSHOnly = serializedObject.FindProperty("m_SHOnly");
m_PropSortNthFrame = serializedObject.FindProperty("m_SortNthFrame");
m_PropRenderMode = serializedObject.FindProperty("m_RenderMode");
m_PropPointDisplaySize = serializedObject.FindProperty("m_PointDisplaySize");
m_PropCutouts = serializedObject.FindProperty("m_Cutouts");
m_PropShaderSplats = serializedObject.FindProperty("m_ShaderSplats");
m_PropShaderComposite = serializedObject.FindProperty("m_ShaderComposite");
m_PropShaderDebugPoints = serializedObject.FindProperty("m_ShaderDebugPoints");
m_PropShaderDebugBoxes = serializedObject.FindProperty("m_ShaderDebugBoxes");
m_PropCSSplatUtilities = serializedObject.FindProperty("m_CSSplatUtilities");
s_AllEditors.Add(this);
}
public void OnDisable()
{
s_AllEditors.Remove(this);
}
public override void OnInspectorGUI()
{
var gs = target as GaussianSplatRenderer;
if (!gs)
return;
serializedObject.Update();
GUILayout.Label("Data Asset", EditorStyles.boldLabel);
EditorGUILayout.PropertyField(m_PropAsset);
if (!gs.HasValidAsset)
{
var msg = gs.asset != null && gs.asset.formatVersion != GaussianSplatAsset.kCurrentVersion
? "Gaussian Splat asset version is not compatible, please recreate the asset"
: "Gaussian Splat asset is not assigned or is empty";
EditorGUILayout.HelpBox(msg, MessageType.Error);
}
EditorGUILayout.Space();
GUILayout.Label("Render Options", EditorStyles.boldLabel);
EditorGUILayout.PropertyField(m_PropSplatScale);
EditorGUILayout.PropertyField(m_PropOpacityScale);
EditorGUILayout.PropertyField(m_PropSHOrder);
EditorGUILayout.PropertyField(m_PropSHOnly);
EditorGUILayout.PropertyField(m_PropSortNthFrame);
EditorGUILayout.Space();
GUILayout.Label("Debugging Tweaks", EditorStyles.boldLabel);
EditorGUILayout.PropertyField(m_PropRenderMode);
if (m_PropRenderMode.intValue is (int)GaussianSplatRenderer.RenderMode.DebugPoints or (int)GaussianSplatRenderer.RenderMode.DebugPointIndices)
EditorGUILayout.PropertyField(m_PropPointDisplaySize);
EditorGUILayout.Space();
m_ResourcesExpanded = EditorGUILayout.Foldout(m_ResourcesExpanded, "Resources", true, EditorStyles.foldoutHeader);
if (m_ResourcesExpanded)
{
EditorGUILayout.PropertyField(m_PropShaderSplats);
EditorGUILayout.PropertyField(m_PropShaderComposite);
EditorGUILayout.PropertyField(m_PropShaderDebugPoints);
EditorGUILayout.PropertyField(m_PropShaderDebugBoxes);
EditorGUILayout.PropertyField(m_PropCSSplatUtilities);
}
bool validAndEnabled = gs && gs.enabled && gs.gameObject.activeInHierarchy && gs.HasValidAsset;
if (validAndEnabled && !gs.HasValidRenderSetup)
{
EditorGUILayout.HelpBox("Shader resources are not set up", MessageType.Error);
validAndEnabled = false;
}
if (validAndEnabled && targets.Length == 1)
{
EditCameras(gs);
EditGUI(gs);
}
if (validAndEnabled && targets.Length > 1)
{
MultiEditGUI();
}
serializedObject.ApplyModifiedProperties();
}
void EditCameras(GaussianSplatRenderer gs)
{
var asset = gs.asset;
var cameras = asset.cameras;
if (cameras != null && cameras.Length != 0)
{
EditorGUILayout.Space();
GUILayout.Label("Cameras", EditorStyles.boldLabel);
var camIndex = EditorGUILayout.IntSlider("Camera", m_CameraIndex, 0, cameras.Length - 1);
camIndex = math.clamp(camIndex, 0, cameras.Length - 1);
if (camIndex != m_CameraIndex)
{
m_CameraIndex = camIndex;
gs.ActivateCamera(camIndex);
}
}
}
void MultiEditGUI()
{
DrawSeparator();
CountTargetSplats(out var totalSplats, out var totalObjects);
EditorGUILayout.LabelField("Total Objects", $"{totalObjects}");
EditorGUILayout.LabelField("Total Splats", $"{totalSplats:N0}");
if (totalSplats > GaussianSplatAsset.kMaxSplats)
{
EditorGUILayout.HelpBox($"Can't merge, too many splats (max. supported {GaussianSplatAsset.kMaxSplats:N0})", MessageType.Warning);
return;
}
var targetGs = (GaussianSplatRenderer) target;
if (!targetGs || !targetGs.HasValidAsset || !targetGs.isActiveAndEnabled)
{
EditorGUILayout.HelpBox($"Can't merge into {target.name} (no asset or disable)", MessageType.Warning);
return;
}
if (targetGs.asset.chunkData != null)
{
EditorGUILayout.HelpBox($"Can't merge into {target.name} (needs to use Very High quality preset)", MessageType.Warning);
return;
}
if (GUILayout.Button($"Merge into {target.name}"))
{
MergeSplatObjects();
}
}
void CountTargetSplats(out int totalSplats, out int totalObjects)
{
totalObjects = 0;
totalSplats = 0;
foreach (var obj in targets)
{
var gs = obj as GaussianSplatRenderer;
if (!gs || !gs.HasValidAsset || !gs.isActiveAndEnabled)
continue;
++totalObjects;
totalSplats += gs.splatCount;
}
}
void MergeSplatObjects()
{
CountTargetSplats(out var totalSplats, out _);
if (totalSplats > GaussianSplatAsset.kMaxSplats)
return;
var targetGs = (GaussianSplatRenderer) target;
int copyDstOffset = targetGs.splatCount;
targetGs.EditSetSplatCount(totalSplats);
foreach (var obj in targets)
{
var gs = obj as GaussianSplatRenderer;
if (!gs || !gs.HasValidAsset || !gs.isActiveAndEnabled)
continue;
if (gs == targetGs)
continue;
gs.EditCopySplatsInto(targetGs, 0, copyDstOffset, gs.splatCount);
copyDstOffset += gs.splatCount;
gs.gameObject.SetActive(false);
}
Debug.Assert(copyDstOffset == totalSplats, $"Merge count mismatch, {copyDstOffset} vs {totalSplats}");
Selection.activeObject = targetGs;
}
void EditGUI(GaussianSplatRenderer gs)
{
++s_EditStatsUpdateCounter;
DrawSeparator();
bool wasToolActive = ToolManager.activeContextType == typeof(GaussianToolContext);
GUILayout.BeginHorizontal();
bool isToolActive = GUILayout.Toggle(wasToolActive, "Edit", EditorStyles.miniButton);
using (new EditorGUI.DisabledScope(!gs.editModified))
{
if (GUILayout.Button("Reset", GUILayout.ExpandWidth(false)))
{
if (EditorUtility.DisplayDialog("Reset Splat Modifications?",
$"This will reset edits of {gs.name} to match the {gs.asset.name} asset. Continue?",
"Yes, reset", "Cancel"))
{
gs.enabled = false;
gs.enabled = true;
}
}
}
GUILayout.EndHorizontal();
if (!wasToolActive && isToolActive)
{
ToolManager.SetActiveContext<GaussianToolContext>();
if (Tools.current == Tool.View)
Tools.current = Tool.Move;
}
if (wasToolActive && !isToolActive)
{
ToolManager.SetActiveContext<GameObjectToolContext>();
}
if (isToolActive && gs.asset.chunkData != null)
{
EditorGUILayout.HelpBox("Splat move/rotate/scale tools need Very High splat quality preset", MessageType.Warning);
}
EditorGUILayout.Space();
GUILayout.BeginHorizontal();
if (GUILayout.Button("Add Cutout"))
{
GaussianCutout cutout = ObjectFactory.CreateGameObject("GSCutout", typeof(GaussianCutout)).GetComponent<GaussianCutout>();
Transform cutoutTr = cutout.transform;
cutoutTr.SetParent(gs.transform, false);
cutoutTr.localScale = (gs.asset.boundsMax - gs.asset.boundsMin) * 0.25f;
gs.m_Cutouts ??= Array.Empty<GaussianCutout>();
ArrayUtility.Add(ref gs.m_Cutouts, cutout);
gs.UpdateEditCountsAndBounds();
EditorUtility.SetDirty(gs);
Selection.activeGameObject = cutout.gameObject;
}
if (GUILayout.Button("Use All Cutouts"))
{
gs.m_Cutouts = FindObjectsByType<GaussianCutout>(FindObjectsSortMode.InstanceID);
gs.UpdateEditCountsAndBounds();
EditorUtility.SetDirty(gs);
}
if (GUILayout.Button("No Cutouts"))
{
gs.m_Cutouts = Array.Empty<GaussianCutout>();
gs.UpdateEditCountsAndBounds();
EditorUtility.SetDirty(gs);
}
GUILayout.EndHorizontal();
EditorGUILayout.PropertyField(m_PropCutouts);
bool hasCutouts = gs.m_Cutouts != null && gs.m_Cutouts.Length != 0;
bool modifiedOrHasCutouts = gs.editModified || hasCutouts;
var asset = gs.asset;
EditorGUILayout.Space();
EditorGUI.BeginChangeCheck();
m_ExportBakeTransform = EditorGUILayout.Toggle("Export in world space", m_ExportBakeTransform);
if (EditorGUI.EndChangeCheck())
{
EditorPrefs.SetBool(kPrefExportBake, m_ExportBakeTransform);
}
if (GUILayout.Button("Export PLY"))
ExportPlyFile(gs, m_ExportBakeTransform);
if (asset.posFormat > GaussianSplatAsset.VectorFormat.Norm16 ||
asset.scaleFormat > GaussianSplatAsset.VectorFormat.Norm16 ||
asset.colorFormat > GaussianSplatAsset.ColorFormat.Float16x4 ||
asset.shFormat > GaussianSplatAsset.SHFormat.Float16)
{
EditorGUILayout.HelpBox(
"It is recommended to use High or VeryHigh quality preset for editing splats, lower levels are lossy",
MessageType.Warning);
}
bool displayEditStats = isToolActive || modifiedOrHasCutouts;
EditorGUILayout.Space();
EditorGUILayout.LabelField("Splats", $"{gs.splatCount:N0}");
if (displayEditStats)
{
EditorGUILayout.LabelField("Cut", $"{gs.editCutSplats:N0}");
EditorGUILayout.LabelField("Deleted", $"{gs.editDeletedSplats:N0}");
EditorGUILayout.LabelField("Selected", $"{gs.editSelectedSplats:N0}");
if (hasCutouts)
{
if (s_EditStatsUpdateCounter > 10)
{
gs.UpdateEditCountsAndBounds();
s_EditStatsUpdateCounter = 0;
}
}
}
}
static void DrawSeparator()
{
EditorGUILayout.Space(12f, true);
GUILayout.Box(GUIContent.none, "sv_iconselector_sep", GUILayout.Height(2), GUILayout.ExpandWidth(true));
EditorGUILayout.Space();
}
bool HasFrameBounds()
{
return true;
}
Bounds OnGetFrameBounds()
{
var gs = target as GaussianSplatRenderer;
if (!gs || !gs.HasValidRenderSetup)
return new Bounds(Vector3.zero, Vector3.one);
Bounds bounds = default;
bounds.SetMinMax(gs.asset.boundsMin, gs.asset.boundsMax);
if (gs.editSelectedSplats > 0)
{
bounds = gs.editSelectedBounds;
}
bounds.extents *= 0.7f;
return TransformBounds(gs.transform, bounds);
}
public static Bounds TransformBounds(Transform tr, Bounds bounds )
{
var center = tr.TransformPoint(bounds.center);
var ext = bounds.extents;
var axisX = tr.TransformVector(ext.x, 0, 0);
var axisY = tr.TransformVector(0, ext.y, 0);
var axisZ = tr.TransformVector(0, 0, ext.z);
// sum their absolute value to get the world extents
ext.x = Mathf.Abs(axisX.x) + Mathf.Abs(axisY.x) + Mathf.Abs(axisZ.x);
ext.y = Mathf.Abs(axisX.y) + Mathf.Abs(axisY.y) + Mathf.Abs(axisZ.y);
ext.z = Mathf.Abs(axisX.z) + Mathf.Abs(axisY.z) + Mathf.Abs(axisZ.z);
return new Bounds { center = center, extents = ext };
}
static unsafe void ExportPlyFile(GaussianSplatRenderer gs, bool bakeTransform)
{
var path = EditorUtility.SaveFilePanel(
"Export Gaussian Splat PLY file", "", $"{gs.asset.name}-edit.ply", "ply");
if (string.IsNullOrWhiteSpace(path))
return;
int kSplatSize = UnsafeUtility.SizeOf<GaussianSplatAssetCreator.InputSplatData>();
using var gpuData = new GraphicsBuffer(GraphicsBuffer.Target.Structured, gs.splatCount, kSplatSize);
if (!gs.EditExportData(gpuData, bakeTransform))
return;
GaussianSplatAssetCreator.InputSplatData[] data = new GaussianSplatAssetCreator.InputSplatData[gpuData.count];
gpuData.GetData(data);
var gpuDeleted = gs.GpuEditDeleted;
uint[] deleted = new uint[gpuDeleted.count];
gpuDeleted.GetData(deleted);
// count non-deleted splats
int aliveCount = 0;
for (int i = 0; i < data.Length; ++i)
{
int wordIdx = i >> 5;
int bitIdx = i & 31;
bool isDeleted = (deleted[wordIdx] & (1u << bitIdx)) != 0;
bool isCutout = data[i].nor.sqrMagnitude > 0;
if (!isDeleted && !isCutout)
++aliveCount;
}
using FileStream fs = new FileStream(path, FileMode.Create, FileAccess.Write);
// note: this is a long string! but we don't use multiline literal because we want guaranteed LF line ending
var header = $"ply\nformat binary_little_endian 1.0\nelement vertex {aliveCount}\nproperty float x\nproperty float y\nproperty float z\nproperty float nx\nproperty float ny\nproperty float nz\nproperty float f_dc_0\nproperty float f_dc_1\nproperty float f_dc_2\nproperty float f_rest_0\nproperty float f_rest_1\nproperty float f_rest_2\nproperty float f_rest_3\nproperty float f_rest_4\nproperty float f_rest_5\nproperty float f_rest_6\nproperty float f_rest_7\nproperty float f_rest_8\nproperty float f_rest_9\nproperty float f_rest_10\nproperty float f_rest_11\nproperty float f_rest_12\nproperty float f_rest_13\nproperty float f_rest_14\nproperty float f_rest_15\nproperty float f_rest_16\nproperty float f_rest_17\nproperty float f_rest_18\nproperty float f_rest_19\nproperty float f_rest_20\nproperty float f_rest_21\nproperty float f_rest_22\nproperty float f_rest_23\nproperty float f_rest_24\nproperty float f_rest_25\nproperty float f_rest_26\nproperty float f_rest_27\nproperty float f_rest_28\nproperty float f_rest_29\nproperty float f_rest_30\nproperty float f_rest_31\nproperty float f_rest_32\nproperty float f_rest_33\nproperty float f_rest_34\nproperty float f_rest_35\nproperty float f_rest_36\nproperty float f_rest_37\nproperty float f_rest_38\nproperty float f_rest_39\nproperty float f_rest_40\nproperty float f_rest_41\nproperty float f_rest_42\nproperty float f_rest_43\nproperty float f_rest_44\nproperty float opacity\nproperty float scale_0\nproperty float scale_1\nproperty float scale_2\nproperty float rot_0\nproperty float rot_1\nproperty float rot_2\nproperty float rot_3\nend_header\n";
fs.Write(Encoding.UTF8.GetBytes(header));
for (int i = 0; i < data.Length; ++i)
{
int wordIdx = i >> 5;
int bitIdx = i & 31;
bool isDeleted = (deleted[wordIdx] & (1u << bitIdx)) != 0;
bool isCutout = data[i].nor.sqrMagnitude > 0;
if (!isDeleted && !isCutout)
{
var splat = data[i];
byte* ptr = (byte*)&splat;
fs.Write(new ReadOnlySpan<byte>(ptr, kSplatSize));
}
}
Debug.Log($"Exported PLY {path} with {aliveCount:N0} splats");
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: b0ce434aee9ae4ee6b1f5cd10ae7c8cb
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,210 @@
// SPDX-License-Identifier: MIT
using System.IO;
using GaussianSplatting.Runtime;
using Unity.Burst;
using Unity.Collections;
using Unity.Collections.LowLevel.Unsafe;
using Unity.Jobs;
using Unity.Mathematics;
using UnityEditor;
using UnityEngine;
using UnityEngine.Experimental.Rendering;
namespace GaussianSplatting.Editor
{
[BurstCompile]
public static class GaussianSplatValidator
{
struct RefItem
{
public string assetPath;
public int cameraIndex;
public float fov;
}
// currently on RTX 3080Ti: 43.76, 39.36, 43.50 PSNR
[MenuItem("Tools/Gaussian Splats/Debug/Validate Render against SBIR")]
public static void ValidateSBIR()
{
ValidateImpl("SBIR");
}
// currently on RTX 3080Ti: matches
[MenuItem("Tools/Gaussian Splats/Debug/Validate Render against D3D12")]
public static void ValidateD3D12()
{
ValidateImpl("D3D12");
}
static unsafe void ValidateImpl(string refPrefix)
{
var gaussians = Object.FindObjectOfType(typeof(GaussianSplatRenderer)) as GaussianSplatRenderer;
{
if (gaussians == null)
{
Debug.LogError("No GaussianSplatRenderer object found");
return;
}
}
var items = new RefItem[]
{
new() {assetPath = "bicycle", cameraIndex = 0, fov = 39.09651f},
new() {assetPath = "truck", cameraIndex = 30, fov = 50},
new() {assetPath = "garden", cameraIndex = 30, fov = 47},
};
var cam = Camera.main;
var oldAsset = gaussians.asset;
var oldCamPos = cam.transform.localPosition;
var oldCamRot = cam.transform.localRotation;
var oldCamFov = cam.fieldOfView;
for (var index = 0; index < items.Length; index++)
{
var item = items[index];
EditorUtility.DisplayProgressBar("Validating Gaussian splat rendering", item.assetPath, (float)index / items.Length);
var path = $"Assets/GaussianAssets/{item.assetPath}-point_cloud-iteration_30000-point_cloud.asset";
var gs = AssetDatabase.LoadAssetAtPath<GaussianSplatAsset>(path);
if (gs == null)
{
Debug.LogError($"Did not find asset for validation item {item.assetPath} at {path}");
continue;
}
var refImageFile = $"../../docs/RefImages/{refPrefix}_{item.assetPath}{item.cameraIndex}.png"; // use our snapshot by default
if (!File.Exists(refImageFile))
{
Debug.LogError($"Did not find reference image for validation item {item.assetPath} at {refImageFile}");
continue;
}
var compareTexture = new Texture2D(4, 4, GraphicsFormat.R8G8B8A8_SRGB, TextureCreationFlags.None);
byte[] refImageBytes = File.ReadAllBytes(refImageFile);
ImageConversion.LoadImage(compareTexture, refImageBytes, false);
int width = compareTexture.width;
int height = compareTexture.height;
var renderTarget = RenderTexture.GetTemporary(width, height, 24, GraphicsFormat.R8G8B8A8_SRGB);
cam.targetTexture = renderTarget;
cam.fieldOfView = item.fov;
var captureTexture = new Texture2D(width, height, GraphicsFormat.R8G8B8A8_SRGB, TextureCreationFlags.None);
NativeArray<Color32> diffPixels = new(width * height, Allocator.Persistent);
gaussians.m_Asset = gs;
gaussians.Update();
gaussians.ActivateCamera(item.cameraIndex);
cam.Render();
Graphics.SetRenderTarget(renderTarget);
captureTexture.ReadPixels(new Rect(0, 0, width, height), 0, 0);
NativeArray<Color32> refPixels = compareTexture.GetPixelData<Color32>(0);
NativeArray<Color32> gotPixels = captureTexture.GetPixelData<Color32>(0);
float psnr = 0, rmse = 0;
int errorsCount = 0;
DiffImagesJob difJob = new DiffImagesJob();
difJob.difPixels = diffPixels;
difJob.refPixels = refPixels;
difJob.gotPixels = gotPixels;
difJob.psnrPtr = &psnr;
difJob.rmsePtr = &rmse;
difJob.difPixCount = &errorsCount;
difJob.Schedule().Complete();
string pathDif = $"../../Shot-{refPrefix}-{item.assetPath}{item.cameraIndex}-diff.png";
string pathRef = $"../../Shot-{refPrefix}-{item.assetPath}{item.cameraIndex}-ref.png";
string pathGot = $"../../Shot-{refPrefix}-{item.assetPath}{item.cameraIndex}-got.png";
if (errorsCount > 50 || psnr < 90.0f)
{
Debug.LogWarning(
$"{refPrefix} {item.assetPath} cam {item.cameraIndex}: RMSE {rmse:F2} PSNR {psnr:F2} diff pixels {errorsCount:N0}");
NativeArray<byte> pngBytes = ImageConversion.EncodeNativeArrayToPNG(diffPixels,
GraphicsFormat.R8G8B8A8_SRGB, (uint) width, (uint) height);
File.WriteAllBytes(pathDif, pngBytes.ToArray());
pngBytes.Dispose();
pngBytes = ImageConversion.EncodeNativeArrayToPNG(refPixels, GraphicsFormat.R8G8B8A8_SRGB,
(uint) width, (uint) height);
File.WriteAllBytes(pathRef, pngBytes.ToArray());
pngBytes.Dispose();
pngBytes = ImageConversion.EncodeNativeArrayToPNG(gotPixels, GraphicsFormat.R8G8B8A8_SRGB,
(uint) width, (uint) height);
File.WriteAllBytes(pathGot, pngBytes.ToArray());
pngBytes.Dispose();
}
else
{
File.Delete(pathDif);
File.Delete(pathRef);
File.Delete(pathGot);
}
diffPixels.Dispose();
RenderTexture.ReleaseTemporary(renderTarget);
Object.DestroyImmediate(captureTexture);
Object.DestroyImmediate(compareTexture);
}
cam.targetTexture = null;
gaussians.m_Asset = oldAsset;
gaussians.Update();
cam.transform.localPosition = oldCamPos;
cam.transform.localRotation = oldCamRot;
cam.fieldOfView = oldCamFov;
EditorUtility.ClearProgressBar();
}
[BurstCompile]
struct DiffImagesJob : IJob
{
public NativeArray<Color32> refPixels;
public NativeArray<Color32> gotPixels;
public NativeArray<Color32> difPixels;
[NativeDisableUnsafePtrRestriction] public unsafe float* rmsePtr;
[NativeDisableUnsafePtrRestriction] public unsafe float* psnrPtr;
[NativeDisableUnsafePtrRestriction] public unsafe int* difPixCount;
public unsafe void Execute()
{
const int kDiffScale = 5;
const int kDiffThreshold = 3 * kDiffScale;
*difPixCount = 0;
double sumSqDif = 0;
for (int i = 0; i < refPixels.Length; ++i)
{
Color32 cref = refPixels[i];
// note: LoadImage always loads PNGs into ARGB order, so swizzle to normal RGBA
cref = new Color32(cref.g, cref.b, cref.a, 255);
refPixels[i] = cref;
Color32 cgot = gotPixels[i];
cgot.a = 255;
gotPixels[i] = cgot;
Color32 cdif = new Color32(0, 0, 0, 255);
cdif.r = (byte)math.abs(cref.r - cgot.r);
cdif.g = (byte)math.abs(cref.g - cgot.g);
cdif.b = (byte)math.abs(cref.b - cgot.b);
sumSqDif += cdif.r * cdif.r + cdif.g * cdif.g + cdif.b * cdif.b;
cdif.r = (byte)math.min(255, cdif.r * kDiffScale);
cdif.g = (byte)math.min(255, cdif.g * kDiffScale);
cdif.b = (byte)math.min(255, cdif.b * kDiffScale);
difPixels[i] = cdif;
if (cdif.r >= kDiffThreshold || cdif.g >= kDiffThreshold || cdif.b >= kDiffThreshold)
{
(*difPixCount)++;
}
}
double meanSqDif = sumSqDif / (refPixels.Length * 3);
double rmse = math.sqrt(meanSqDif);
double psnr = 20.0 * math.log10(255.0) - 10.0 * math.log10(rmse * rmse);
*rmsePtr = (float) rmse;
*psnrPtr = (float) psnr;
}
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 2f8e75b80eb181a4698f733ba59b694b
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,21 @@
{
"name": "GaussianSplattingEditor",
"rootNamespace": "GaussianSplatting.Editor",
"references": [
"GUID:4b653174f8fcdcd49b4c9a6f1ca8c7c3",
"GUID:2665a8d13d1b3f18800f46e256720795",
"GUID:d8b63aba1907145bea998dd612889d6b",
"GUID:e0cd26848372d4e5c891c569017e11f1"
],
"includePlatforms": [
"Editor"
],
"excludePlatforms": [],
"allowUnsafeCode": true,
"overrideReferences": false,
"precompiledReferences": [],
"autoReferenced": true,
"defineConstraints": [],
"versionDefines": [],
"noEngineReferences": false
}

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: 14414175af4b366469db63f2efee475f
AssemblyDefinitionImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,43 @@
// SPDX-License-Identifier: MIT
using GaussianSplatting.Runtime;
using UnityEditor.EditorTools;
using UnityEngine;
namespace GaussianSplatting.Editor
{
abstract class GaussianTool : EditorTool
{
protected GaussianSplatRenderer GetRenderer()
{
var gs = target as GaussianSplatRenderer;
if (!gs || !gs.HasValidAsset || !gs.HasValidRenderSetup)
return null;
return gs;
}
protected bool CanBeEdited()
{
var gs = GetRenderer();
if (!gs)
return false;
return gs.asset.chunkData == null; // need to be lossless / non-chunked for editing
}
protected bool HasSelection()
{
var gs = GetRenderer();
if (!gs)
return false;
return gs.editSelectedSplats > 0;
}
protected Vector3 GetSelectionCenterLocal()
{
var gs = GetRenderer();
if (!gs || gs.editSelectedSplats == 0)
return Vector3.zero;
return gs.editSelectedBounds.center;
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 6203c808ab9e64a4a8ff0277c5aa7669
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,192 @@
// SPDX-License-Identifier: MIT
using System;
using GaussianSplatting.Runtime;
using UnityEditor;
using UnityEditor.EditorTools;
using UnityEngine;
namespace GaussianSplatting.Editor
{
[EditorToolContext("GaussianSplats", typeof(GaussianSplatRenderer)), Icon(k_IconPath)]
class GaussianToolContext : EditorToolContext
{
const string k_IconPath = "Packages/org.nesnausk.gaussian-splatting/Editor/Icons/GaussianContext.png";
Vector2 m_MouseStartDragPos;
protected override Type GetEditorToolType(Tool tool)
{
if (tool == Tool.Move)
return typeof(GaussianMoveTool);
//if (tool == Tool.Rotate)
// return typeof(GaussianRotateTool); // not correctly working yet
//if (tool == Tool.Scale)
// return typeof(GaussianScaleTool); // not working correctly yet when the GS itself has scale
return null;
}
public override void OnWillBeDeactivated()
{
var gs = target as GaussianSplatRenderer;
if (!gs)
return;
gs.EditDeselectAll();
}
static void HandleKeyboardCommands(Event evt, GaussianSplatRenderer gs)
{
if (evt.type != EventType.ValidateCommand && evt.type != EventType.ExecuteCommand)
return;
bool execute = evt.type == EventType.ExecuteCommand;
switch (evt.commandName)
{
// ugh, EventCommandNames string constants is internal :(
case "SoftDelete":
case "Delete":
if (execute)
{
gs.EditDeleteSelected();
GaussianSplatRendererEditor.RepaintAll();
}
evt.Use();
break;
case "SelectAll":
if (execute)
{
gs.EditSelectAll();
GaussianSplatRendererEditor.RepaintAll();
}
evt.Use();
break;
case "DeselectAll":
if (execute)
{
gs.EditDeselectAll();
GaussianSplatRendererEditor.RepaintAll();
}
evt.Use();
break;
case "InvertSelection":
if (execute)
{
gs.EditInvertSelection();
GaussianSplatRendererEditor.RepaintAll();
}
evt.Use();
break;
}
}
static bool IsViewToolActive()
{
return Tools.viewToolActive || Tools.current == Tool.View || (Event.current != null && Event.current.alt);
}
public override void OnToolGUI(EditorWindow window)
{
if (!(window is SceneView sceneView))
return;
var gs = target as GaussianSplatRenderer;
if (!gs)
return;
GaussianSplatRendererEditor.BumpGUICounter();
int id = GUIUtility.GetControlID(FocusType.Passive);
Event evt = Event.current;
HandleKeyboardCommands(evt, gs);
var evtType = evt.GetTypeForControl(id);
switch (evtType)
{
case EventType.Layout:
// make this be the default tool, so that we get focus when user clicks on nothing else
HandleUtility.AddDefaultControl(id);
break;
case EventType.MouseDown:
if (IsViewToolActive())
break;
if (HandleUtility.nearestControl == id && evt.button == 0)
{
// shift/command adds to selection, ctrl removes from selection: if none of these
// are present, start a new selection
if (!evt.shift && !EditorGUI.actionKey && !evt.control)
gs.EditDeselectAll();
// record selection state at start
gs.EditStoreSelectionMouseDown();
GaussianSplatRendererEditor.RepaintAll();
GUIUtility.hotControl = id;
m_MouseStartDragPos = evt.mousePosition;
evt.Use();
}
break;
case EventType.MouseDrag:
if (GUIUtility.hotControl == id && evt.button == 0)
{
Rect rect = FromToRect(m_MouseStartDragPos, evt.mousePosition);
Vector2 rectMin = HandleUtility.GUIPointToScreenPixelCoordinate(rect.min);
Vector2 rectMax = HandleUtility.GUIPointToScreenPixelCoordinate(rect.max);
gs.EditUpdateSelection(rectMin, rectMax, sceneView.camera, evt.control);
GaussianSplatRendererEditor.RepaintAll();
evt.Use();
}
break;
case EventType.MouseUp:
if (GUIUtility.hotControl == id && evt.button == 0)
{
m_MouseStartDragPos = Vector2.zero;
GUIUtility.hotControl = 0;
evt.Use();
}
break;
case EventType.Repaint:
// draw cutout gizmos
Handles.color = new Color(1,0,1,0.7f);
var prevMatrix = Handles.matrix;
foreach (var cutout in gs.m_Cutouts)
{
if (!cutout)
continue;
Handles.matrix = cutout.transform.localToWorldMatrix;
if (cutout.m_Type == GaussianCutout.Type.Ellipsoid)
{
Handles.DrawWireDisc(Vector3.zero, Vector3.up, 1.0f);
Handles.DrawWireDisc(Vector3.zero, Vector3.right, 1.0f);
Handles.DrawWireDisc(Vector3.zero, Vector3.forward, 1.0f);
}
if (cutout.m_Type == GaussianCutout.Type.Box)
Handles.DrawWireCube(Vector3.zero, Vector3.one * 2);
}
Handles.matrix = prevMatrix;
// draw selection bounding box
if (gs.editSelectedSplats > 0)
{
var selBounds = GaussianSplatRendererEditor.TransformBounds(gs.transform, gs.editSelectedBounds);
Handles.DrawWireCube(selBounds.center, selBounds.size);
}
// draw drag rectangle
if (GUIUtility.hotControl == id && evt.mousePosition != m_MouseStartDragPos)
{
GUIStyle style = "SelectionRect";
Handles.BeginGUI();
style.Draw(FromToRect(m_MouseStartDragPos, evt.mousePosition), false, false, false, false);
Handles.EndGUI();
}
break;
}
}
// build a rect that always has a positive size
static Rect FromToRect(Vector2 from, Vector2 to)
{
if (from.x > to.x)
(from.x, to.x) = (to.x, from.x);
if (from.y > to.y)
(from.y, to.y) = (to.y, from.y);
return new Rect(from.x, from.y, to.x - from.x, to.y - from.y);
}
}
}

View File

@@ -0,0 +1,3 @@
fileFormatVersion: 2
guid: 80d7ecbaa1b24e6399ee95f6fc0b9c90
timeCreated: 1697718362

View File

@@ -0,0 +1,8 @@
fileFormatVersion: 2
guid: 770e497b696b99641aa1bf295d0b3552
folderAsset: yes
DefaultImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

Binary file not shown.

After

Width:  |  Height:  |  Size: 4.4 KiB

View File

@@ -0,0 +1,153 @@
fileFormatVersion: 2
guid: c61202bc8cc557546afa505174da220e
TextureImporter:
internalIDToNameTable: []
externalObjects: {}
serializedVersion: 12
mipmaps:
mipMapMode: 0
enableMipMap: 0
sRGBTexture: 1
linearTexture: 0
fadeOut: 0
borderMipMap: 0
mipMapsPreserveCoverage: 0
alphaTestReferenceValue: 0.5
mipMapFadeDistanceStart: 1
mipMapFadeDistanceEnd: 3
bumpmap:
convertToNormalMap: 0
externalNormalMap: 0
heightScale: 0.25
normalMapFilter: 0
flipGreenChannel: 0
isReadable: 0
streamingMipmaps: 0
streamingMipmapsPriority: 0
vTOnly: 0
ignoreMipmapLimit: 0
grayScaleToAlpha: 0
generateCubemap: 6
cubemapConvolution: 0
seamlessCubemap: 0
textureFormat: 1
maxTextureSize: 2048
textureSettings:
serializedVersion: 2
filterMode: 1
aniso: 1
mipBias: 0
wrapU: 1
wrapV: 1
wrapW: 0
nPOTScale: 0
lightmap: 0
compressionQuality: 50
spriteMode: 0
spriteExtrude: 1
spriteMeshType: 1
alignment: 0
spritePivot: {x: 0.5, y: 0.5}
spritePixelsToUnits: 100
spriteBorder: {x: 0, y: 0, z: 0, w: 0}
spriteGenerateFallbackPhysicsShape: 1
alphaUsage: 1
alphaIsTransparency: 1
spriteTessellationDetail: -1
textureType: 2
textureShape: 1
singleChannelComponent: 0
flipbookRows: 1
flipbookColumns: 1
maxTextureSizeSet: 0
compressionQualitySet: 0
textureFormatSet: 0
ignorePngGamma: 0
applyGammaDecoding: 0
swizzle: 50462976
cookieLightType: 1
platformSettings:
- serializedVersion: 3
buildTarget: DefaultTexturePlatform
maxTextureSize: 128
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Standalone
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Nintendo Switch
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: LinuxHeadlessSimulation
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Server
maxTextureSize: 64
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
spriteSheet:
serializedVersion: 2
sprites: []
outline: []
physicsShape: []
bones: []
spriteID:
internalID: 0
vertices: []
indices:
edges: []
weights: []
secondaryTextures: []
nameFileIdTable: {}
mipmapLimitGroupName:
pSDRemoveMatte: 0
userData:
assetBundleName:
assetBundleVariant:

Binary file not shown.

After

Width:  |  Height:  |  Size: 4.8 KiB

View File

@@ -0,0 +1,153 @@
fileFormatVersion: 2
guid: 15d0de03329e14440b034e884fe10379
TextureImporter:
internalIDToNameTable: []
externalObjects: {}
serializedVersion: 12
mipmaps:
mipMapMode: 0
enableMipMap: 0
sRGBTexture: 1
linearTexture: 0
fadeOut: 0
borderMipMap: 0
mipMapsPreserveCoverage: 0
alphaTestReferenceValue: 0.5
mipMapFadeDistanceStart: 1
mipMapFadeDistanceEnd: 3
bumpmap:
convertToNormalMap: 0
externalNormalMap: 0
heightScale: 0.25
normalMapFilter: 0
flipGreenChannel: 0
isReadable: 0
streamingMipmaps: 0
streamingMipmapsPriority: 0
vTOnly: 0
ignoreMipmapLimit: 0
grayScaleToAlpha: 0
generateCubemap: 6
cubemapConvolution: 0
seamlessCubemap: 0
textureFormat: 1
maxTextureSize: 2048
textureSettings:
serializedVersion: 2
filterMode: 1
aniso: 1
mipBias: 0
wrapU: 1
wrapV: 1
wrapW: 0
nPOTScale: 0
lightmap: 0
compressionQuality: 50
spriteMode: 0
spriteExtrude: 1
spriteMeshType: 1
alignment: 0
spritePivot: {x: 0.5, y: 0.5}
spritePixelsToUnits: 100
spriteBorder: {x: 0, y: 0, z: 0, w: 0}
spriteGenerateFallbackPhysicsShape: 1
alphaUsage: 1
alphaIsTransparency: 1
spriteTessellationDetail: -1
textureType: 2
textureShape: 1
singleChannelComponent: 0
flipbookRows: 1
flipbookColumns: 1
maxTextureSizeSet: 0
compressionQualitySet: 0
textureFormatSet: 0
ignorePngGamma: 0
applyGammaDecoding: 0
swizzle: 50462976
cookieLightType: 1
platformSettings:
- serializedVersion: 3
buildTarget: DefaultTexturePlatform
maxTextureSize: 128
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Standalone
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Nintendo Switch
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: LinuxHeadlessSimulation
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Server
maxTextureSize: 64
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
spriteSheet:
serializedVersion: 2
sprites: []
outline: []
physicsShape: []
bones: []
spriteID:
internalID: 0
vertices: []
indices:
edges: []
weights: []
secondaryTextures: []
nameFileIdTable: {}
mipmapLimitGroupName:
pSDRemoveMatte: 0
userData:
assetBundleName:
assetBundleVariant:

Binary file not shown.

After

Width:  |  Height:  |  Size: 4.4 KiB

View File

@@ -0,0 +1,231 @@
fileFormatVersion: 2
guid: 81df9c0903abfa345a9022d090982f5d
TextureImporter:
internalIDToNameTable: []
externalObjects: {}
serializedVersion: 12
mipmaps:
mipMapMode: 0
enableMipMap: 0
sRGBTexture: 1
linearTexture: 0
fadeOut: 0
borderMipMap: 0
mipMapsPreserveCoverage: 0
alphaTestReferenceValue: 0.5
mipMapFadeDistanceStart: 1
mipMapFadeDistanceEnd: 3
bumpmap:
convertToNormalMap: 0
externalNormalMap: 0
heightScale: 0.25
normalMapFilter: 0
flipGreenChannel: 0
isReadable: 0
streamingMipmaps: 0
streamingMipmapsPriority: 0
vTOnly: 0
ignoreMipmapLimit: 0
grayScaleToAlpha: 0
generateCubemap: 6
cubemapConvolution: 0
seamlessCubemap: 0
textureFormat: 1
maxTextureSize: 2048
textureSettings:
serializedVersion: 2
filterMode: 1
aniso: 1
mipBias: 0
wrapU: 1
wrapV: 1
wrapW: 0
nPOTScale: 0
lightmap: 0
compressionQuality: 50
spriteMode: 0
spriteExtrude: 1
spriteMeshType: 1
alignment: 0
spritePivot: {x: 0.5, y: 0.5}
spritePixelsToUnits: 100
spriteBorder: {x: 0, y: 0, z: 0, w: 0}
spriteGenerateFallbackPhysicsShape: 1
alphaUsage: 1
alphaIsTransparency: 1
spriteTessellationDetail: -1
textureType: 2
textureShape: 1
singleChannelComponent: 0
flipbookRows: 1
flipbookColumns: 1
maxTextureSizeSet: 0
compressionQualitySet: 0
textureFormatSet: 0
ignorePngGamma: 0
applyGammaDecoding: 0
swizzle: 50462976
cookieLightType: 1
platformSettings:
- serializedVersion: 3
buildTarget: DefaultTexturePlatform
maxTextureSize: 128
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Standalone
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Android
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: iPhone
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Windows Store Apps
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: tvOS
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Lumin
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: CloudRendering
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Nintendo Switch
maxTextureSize: 64
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: LinuxHeadlessSimulation
maxTextureSize: 64
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Server
maxTextureSize: 64
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
spriteSheet:
serializedVersion: 2
sprites: []
outline: []
physicsShape: []
bones: []
spriteID:
internalID: 0
vertices: []
indices:
edges: []
weights: []
secondaryTextures: []
nameFileIdTable: {}
mipmapLimitGroupName:
pSDRemoveMatte: 0
userData:
assetBundleName:
assetBundleVariant:

Binary file not shown.

After

Width:  |  Height:  |  Size: 4.8 KiB

View File

@@ -0,0 +1,153 @@
fileFormatVersion: 2
guid: e37880566baf3964e9b75e45adb36f3f
TextureImporter:
internalIDToNameTable: []
externalObjects: {}
serializedVersion: 12
mipmaps:
mipMapMode: 0
enableMipMap: 0
sRGBTexture: 1
linearTexture: 0
fadeOut: 0
borderMipMap: 0
mipMapsPreserveCoverage: 0
alphaTestReferenceValue: 0.5
mipMapFadeDistanceStart: 1
mipMapFadeDistanceEnd: 3
bumpmap:
convertToNormalMap: 0
externalNormalMap: 0
heightScale: 0.25
normalMapFilter: 0
flipGreenChannel: 0
isReadable: 0
streamingMipmaps: 0
streamingMipmapsPriority: 0
vTOnly: 0
ignoreMipmapLimit: 0
grayScaleToAlpha: 0
generateCubemap: 6
cubemapConvolution: 0
seamlessCubemap: 0
textureFormat: 1
maxTextureSize: 2048
textureSettings:
serializedVersion: 2
filterMode: 1
aniso: 1
mipBias: 0
wrapU: 1
wrapV: 1
wrapW: 0
nPOTScale: 0
lightmap: 0
compressionQuality: 50
spriteMode: 0
spriteExtrude: 1
spriteMeshType: 1
alignment: 0
spritePivot: {x: 0.5, y: 0.5}
spritePixelsToUnits: 100
spriteBorder: {x: 0, y: 0, z: 0, w: 0}
spriteGenerateFallbackPhysicsShape: 1
alphaUsage: 1
alphaIsTransparency: 1
spriteTessellationDetail: -1
textureType: 2
textureShape: 1
singleChannelComponent: 0
flipbookRows: 1
flipbookColumns: 1
maxTextureSizeSet: 0
compressionQualitySet: 0
textureFormatSet: 0
ignorePngGamma: 0
applyGammaDecoding: 0
swizzle: 50462976
cookieLightType: 1
platformSettings:
- serializedVersion: 3
buildTarget: DefaultTexturePlatform
maxTextureSize: 128
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Standalone
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Nintendo Switch
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: LinuxHeadlessSimulation
maxTextureSize: 2048
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 1
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
- serializedVersion: 3
buildTarget: Server
maxTextureSize: 64
resizeAlgorithm: 0
textureFormat: -1
textureCompression: 0
compressionQuality: 50
crunchedCompression: 0
allowsAlphaSplitting: 0
overridden: 0
ignorePlatformSupport: 0
androidETC2FallbackOverride: 0
forceMaximumCompressionQuality_BC6H_BC7: 0
spriteSheet:
serializedVersion: 2
sprites: []
outline: []
physicsShape: []
bones: []
spriteID:
internalID: 0
vertices: []
indices:
edges: []
weights: []
secondaryTextures: []
nameFileIdTable: {}
mipmapLimitGroupName:
pSDRemoveMatte: 0
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,8 @@
fileFormatVersion: 2
guid: f812890ad0ea4c747bdc67b6d2c1c627
folderAsset: yes
DefaultImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,26 @@
// SPDX-License-Identifier: MIT
using UnityEditor;
using UnityEngine;
namespace GaussianSplatting.Editor.Utils
{
public class CaptureScreenshot : MonoBehaviour
{
[MenuItem("Tools/Gaussian Splats/Debug/Capture Screenshot %g")]
public static void CaptureShot()
{
int counter = 0;
string path;
while(true)
{
path = $"Shot-{counter:0000}.png";
if (!System.IO.File.Exists(path))
break;
++counter;
}
ScreenCapture.CaptureScreenshot(path);
Debug.Log($"Captured {path}");
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 6c80a2b8daebbc1449b79e5ec436f39d
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,274 @@
// SPDX-License-Identifier: MIT
using System;
using System.Collections.Generic;
using System.IO;
using System.Linq;
using UnityEditor;
using UnityEditor.Experimental;
using UnityEngine;
namespace GaussianSplatting.Editor.Utils
{
public class FilePickerControl
{
const string kLastPathPref = "nesnausk.utils.FilePickerLastPath";
static Texture2D s_FolderIcon => EditorGUIUtility.FindTexture(EditorResources.emptyFolderIconName);
static Texture2D s_FileIcon => EditorGUIUtility.FindTexture(EditorResources.folderIconName);
static GUIStyle s_StyleTextFieldText;
static GUIStyle s_StyleTextFieldDropdown;
static readonly int kPathFieldControlID = "FilePickerPathField".GetHashCode();
const int kIconSize = 15;
const int kRecentPathsCount = 20;
public static string PathToDisplayString(string path)
{
if (string.IsNullOrWhiteSpace(path))
return "<none>";
path = path.Replace('\\', '/');
string[] parts = path.Split('/');
// check if filename is not some super generic one
var baseName = Path.GetFileNameWithoutExtension(parts[^1]).ToLowerInvariant();
if (baseName != "point_cloud" && baseName != "splat" && baseName != "input")
return parts[^1];
// otherwise if filename is just some generic "point cloud" type, then take some folder names above it into account
if (parts.Length >= 4)
path = string.Join('/', parts.TakeLast(4));
path = path.Replace('/', '-');
return path;
}
class PreviousPaths
{
public PreviousPaths(List<string> paths)
{
this.paths = paths;
UpdateContent();
}
public void UpdateContent()
{
this.content = paths.Select(p => new GUIContent(PathToDisplayString(p))).ToArray();
}
public List<string> paths;
public GUIContent[] content;
}
Dictionary<string, PreviousPaths> m_PreviousPaths = new();
void PopulatePreviousPaths(string nameKey)
{
if (m_PreviousPaths.ContainsKey(nameKey))
return;
List<string> prevPaths = new();
for (int i = 0; i < kRecentPathsCount; ++i)
{
string path = EditorPrefs.GetString($"{kLastPathPref}-{nameKey}-{i}");
if (!string.IsNullOrWhiteSpace(path))
prevPaths.Add(path);
}
m_PreviousPaths.Add(nameKey, new PreviousPaths(prevPaths));
}
void UpdatePreviousPaths(string nameKey, string path)
{
if (!m_PreviousPaths.ContainsKey(nameKey))
{
m_PreviousPaths.Add(nameKey, new PreviousPaths(new List<string>()));
}
var prevPaths = m_PreviousPaths[nameKey];
prevPaths.paths.Remove(path);
prevPaths.paths.Insert(0, path);
while (prevPaths.paths.Count > kRecentPathsCount)
prevPaths.paths.RemoveAt(prevPaths.paths.Count - 1);
prevPaths.UpdateContent();
for (int i = 0; i < prevPaths.paths.Count; ++i)
{
EditorPrefs.SetString($"{kLastPathPref}-{nameKey}-{i}", prevPaths.paths[i]);
}
}
static bool CheckPath(string path, bool isFolder)
{
if (string.IsNullOrWhiteSpace(path))
return false;
if (isFolder)
{
if (!Directory.Exists(path))
return false;
}
else
{
if (!File.Exists(path))
return false;
}
return true;
}
static string PathAbsToStorage(string path)
{
path = path.Replace('\\', '/');
var dataPath = Application.dataPath;
if (path.StartsWith(dataPath, StringComparison.Ordinal))
{
path = Path.GetRelativePath($"{dataPath}/..", path);
path = path.Replace('\\', '/');
}
return path;
}
bool CheckAndSetNewPath(ref string path, string nameKey, bool isFolder)
{
path = PathAbsToStorage(path);
if (CheckPath(path, isFolder))
{
EditorPrefs.SetString($"{kLastPathPref}-{nameKey}", path);
UpdatePreviousPaths(nameKey, path);
GUI.changed = true;
Event.current.Use();
return true;
}
return false;
}
string PreviousPathsDropdown(Rect position, string value, string nameKey, bool isFolder)
{
PopulatePreviousPaths(nameKey);
if (string.IsNullOrWhiteSpace(value))
value = EditorPrefs.GetString($"{kLastPathPref}-{nameKey}");
m_PreviousPaths.TryGetValue(nameKey, out var prevPaths);
EditorGUI.BeginDisabledGroup(prevPaths == null || prevPaths.paths.Count == 0);
EditorGUI.BeginChangeCheck();
int oldIndent = EditorGUI.indentLevel;
EditorGUI.indentLevel = 0;
int parameterIndex = EditorGUI.Popup(position, GUIContent.none, -1, prevPaths.content, s_StyleTextFieldDropdown);
if (EditorGUI.EndChangeCheck() && parameterIndex < prevPaths.paths.Count)
{
string newValue = prevPaths.paths[parameterIndex];
if (CheckAndSetNewPath(ref newValue, nameKey, isFolder))
value = newValue;
}
EditorGUI.indentLevel = oldIndent;
EditorGUI.EndDisabledGroup();
return value;
}
// null extension picks folders
public string PathFieldGUI(Rect position, GUIContent label, string value, string extension, string nameKey)
{
s_StyleTextFieldText ??= new GUIStyle("TextFieldDropDownText");
s_StyleTextFieldDropdown ??= new GUIStyle("TextFieldDropdown");
bool isFolder = extension == null;
int controlId = GUIUtility.GetControlID(kPathFieldControlID, FocusType.Keyboard, position);
Rect fullRect = EditorGUI.PrefixLabel(position, controlId, label);
Rect textRect = new Rect(fullRect.x, fullRect.y, fullRect.width - s_StyleTextFieldDropdown.fixedWidth, fullRect.height);
Rect dropdownRect = new Rect(textRect.xMax, fullRect.y, s_StyleTextFieldDropdown.fixedWidth, fullRect.height);
Rect iconRect = new Rect(textRect.xMax - kIconSize, textRect.y, kIconSize, textRect.height);
value = PreviousPathsDropdown(dropdownRect, value, nameKey, isFolder);
string displayText = PathToDisplayString(value);
Event evt = Event.current;
switch (evt.type)
{
case EventType.KeyDown:
if (GUIUtility.keyboardControl == controlId)
{
if (evt.keyCode is KeyCode.Backspace or KeyCode.Delete)
{
value = null;
EditorPrefs.SetString($"{kLastPathPref}-{nameKey}", "");
GUI.changed = true;
evt.Use();
}
}
break;
case EventType.Repaint:
s_StyleTextFieldText.Draw(textRect, new GUIContent(displayText), controlId, DragAndDrop.activeControlID == controlId);
GUI.DrawTexture(iconRect, isFolder ? s_FolderIcon : s_FileIcon, ScaleMode.ScaleToFit);
break;
case EventType.MouseDown:
if (evt.button != 0 || !GUI.enabled)
break;
if (textRect.Contains(evt.mousePosition))
{
if (iconRect.Contains(evt.mousePosition))
{
if (string.IsNullOrWhiteSpace(value))
value = EditorPrefs.GetString($"{kLastPathPref}-{nameKey}");
string newPath;
string openToPath = string.Empty;
if (isFolder)
{
if (Directory.Exists(value))
openToPath = value;
newPath = EditorUtility.OpenFolderPanel("Select folder", openToPath, "");
}
else
{
if (File.Exists(value))
openToPath = Path.GetDirectoryName(value);
newPath = EditorUtility.OpenFilePanel("Select file", openToPath, extension);
}
if (CheckAndSetNewPath(ref newPath, nameKey, isFolder))
{
value = newPath;
GUI.changed = true;
evt.Use();
}
}
else if (File.Exists(value) || Directory.Exists(value))
{
EditorUtility.RevealInFinder(value);
}
GUIUtility.keyboardControl = controlId;
}
break;
case EventType.DragUpdated:
case EventType.DragPerform:
if (textRect.Contains(evt.mousePosition) && GUI.enabled)
{
if (DragAndDrop.paths.Length > 0)
{
DragAndDrop.visualMode = DragAndDropVisualMode.Generic;
string path = DragAndDrop.paths[0];
path = PathAbsToStorage(path);
if (CheckPath(path, isFolder))
{
if (evt.type == EventType.DragPerform)
{
UpdatePreviousPaths(nameKey, path);
value = path;
GUI.changed = true;
DragAndDrop.AcceptDrag();
DragAndDrop.activeControlID = 0;
}
else
DragAndDrop.activeControlID = controlId;
}
else
DragAndDrop.visualMode = DragAndDropVisualMode.Rejected;
evt.Use();
}
}
break;
case EventType.DragExited:
if (GUI.enabled)
{
HandleUtility.Repaint();
}
break;
}
return value;
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 69e6c946494a9b2479ce96542339029c
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,595 @@
// SPDX-License-Identifier: MIT
using System;
using Unity.Burst;
using Unity.Burst.Intrinsics;
using Unity.Collections;
using Unity.Collections.LowLevel.Unsafe;
using Unity.Jobs;
using Unity.Mathematics;
using Unity.Profiling;
using Unity.Profiling.LowLevel;
namespace GaussianSplatting.Editor.Utils
{
// Implementation of "Mini Batch" k-means clustering ("Web-Scale K-Means Clustering", Sculley 2010)
// using k-means++ for cluster initialization.
[BurstCompile]
public struct KMeansClustering
{
static ProfilerMarker s_ProfCalculate = new(ProfilerCategory.Render, "KMeans.Calculate", MarkerFlags.SampleGPU);
static ProfilerMarker s_ProfPlusPlus = new(ProfilerCategory.Render, "KMeans.InitialPlusPlus", MarkerFlags.SampleGPU);
static ProfilerMarker s_ProfInitialDistanceSum = new(ProfilerCategory.Render, "KMeans.Initialize.DistanceSum", MarkerFlags.SampleGPU);
static ProfilerMarker s_ProfInitialPickPoint = new(ProfilerCategory.Render, "KMeans.Initialize.PickPoint", MarkerFlags.SampleGPU);
static ProfilerMarker s_ProfInitialDistanceUpdate = new(ProfilerCategory.Render, "KMeans.Initialize.DistanceUpdate", MarkerFlags.SampleGPU);
static ProfilerMarker s_ProfAssignClusters = new(ProfilerCategory.Render, "KMeans.AssignClusters", MarkerFlags.SampleGPU);
static ProfilerMarker s_ProfUpdateMeans = new(ProfilerCategory.Render, "KMeans.UpdateMeans", MarkerFlags.SampleGPU);
public static bool Calculate(int dim, NativeArray<float> inputData, int batchSize, float passesOverData, Func<float,bool> progress, NativeArray<float> outClusterMeans, NativeArray<int> outDataLabels)
{
// Parameter checks
if (dim < 1)
throw new InvalidOperationException($"KMeans: dimensionality has to be >= 1, was {dim}");
if (batchSize < 1)
throw new InvalidOperationException($"KMeans: batch size has to be >= 1, was {batchSize}");
if (passesOverData < 0.0001f)
throw new InvalidOperationException($"KMeans: passes over data must be positive, was {passesOverData}");
if (inputData.Length % dim != 0)
throw new InvalidOperationException($"KMeans: input length must be multiple of dim={dim}, was {inputData.Length}");
if (outClusterMeans.Length % dim != 0)
throw new InvalidOperationException($"KMeans: output means length must be multiple of dim={dim}, was {outClusterMeans.Length}");
int dataSize = inputData.Length / dim;
int k = outClusterMeans.Length / dim;
if (k < 1)
throw new InvalidOperationException($"KMeans: cluster count length must be at least 1, was {k}");
if (dataSize < k)
throw new InvalidOperationException($"KMeans: input length ({inputData.Length}) must at least as long as clusters ({outClusterMeans.Length})");
if (dataSize != outDataLabels.Length)
throw new InvalidOperationException($"KMeans: output labels length must be {dataSize}, was {outDataLabels.Length}");
using var prof = s_ProfCalculate.Auto();
batchSize = math.min(dataSize, batchSize);
uint rngState = 1;
// Do initial cluster placement
int initBatchSize = 10 * k;
const int kInitAttempts = 3;
if (!InitializeCentroids(dim, inputData, initBatchSize, ref rngState, kInitAttempts, outClusterMeans, progress))
return false;
NativeArray<float> counts = new(k, Allocator.TempJob);
NativeArray<float> batchPoints = new(batchSize * dim, Allocator.TempJob);
NativeArray<int> batchClusters = new(batchSize, Allocator.TempJob);
bool cancelled = false;
for (float calcDone = 0.0f, calcLimit = dataSize * passesOverData; calcDone < calcLimit; calcDone += batchSize)
{
if (progress != null && !progress(0.3f + calcDone / calcLimit * 0.4f))
{
cancelled = true;
break;
}
// generate a batch of random input points
MakeRandomBatch(dim, inputData, ref rngState, batchPoints);
// find which of the current centroids each batch point is closest to
{
using var profPart = s_ProfAssignClusters.Auto();
AssignClustersJob job = new AssignClustersJob
{
dim = dim,
data = batchPoints,
means = outClusterMeans,
indexOffset = 0,
clusters = batchClusters,
};
job.Schedule(batchSize, 1).Complete();
}
// update the centroids
{
using var profPart = s_ProfUpdateMeans.Auto();
UpdateCentroidsJob job = new UpdateCentroidsJob
{
m_Clusters = outClusterMeans,
m_Dim = dim,
m_Counts = counts,
m_BatchSize = batchSize,
m_BatchClusters = batchClusters,
m_BatchPoints = batchPoints
};
job.Schedule().Complete();
}
}
// finally find out closest clusters for all input points
{
using var profPart = s_ProfAssignClusters.Auto();
const int kAssignBatchCount = 256 * 1024;
AssignClustersJob job = new AssignClustersJob
{
dim = dim,
data = inputData,
means = outClusterMeans,
indexOffset = 0,
clusters = outDataLabels,
};
for (int i = 0; i < dataSize; i += kAssignBatchCount)
{
if (progress != null && !progress(0.7f + (float) i / dataSize * 0.3f))
{
cancelled = true;
break;
}
job.indexOffset = i;
job.Schedule(math.min(kAssignBatchCount, dataSize - i), 512).Complete();
}
}
counts.Dispose();
batchPoints.Dispose();
batchClusters.Dispose();
return !cancelled;
}
static unsafe float DistanceSquared(int dim, NativeArray<float> a, int aIndex, NativeArray<float> b, int bIndex)
{
aIndex *= dim;
bIndex *= dim;
float d = 0;
if (X86.Avx.IsAvxSupported)
{
// 8x wide with AVX
int i = 0;
float* aptr = (float*) a.GetUnsafeReadOnlyPtr() + aIndex;
float* bptr = (float*) b.GetUnsafeReadOnlyPtr() + bIndex;
for (; i + 7 < dim; i += 8)
{
v256 va = X86.Avx.mm256_loadu_ps(aptr);
v256 vb = X86.Avx.mm256_loadu_ps(bptr);
v256 vd = X86.Avx.mm256_sub_ps(va, vb);
vd = X86.Avx.mm256_mul_ps(vd, vd);
vd = X86.Avx.mm256_hadd_ps(vd, vd);
d += vd.Float0 + vd.Float1 + vd.Float4 + vd.Float5;
aptr += 8;
bptr += 8;
}
// remainder
for (; i < dim; ++i)
{
float delta = *aptr - *bptr;
d += delta * delta;
aptr++;
bptr++;
}
}
else if (Arm.Neon.IsNeonSupported)
{
// 4x wide with NEON
int i = 0;
float* aptr = (float*) a.GetUnsafeReadOnlyPtr() + aIndex;
float* bptr = (float*) b.GetUnsafeReadOnlyPtr() + bIndex;
for (; i + 3 < dim; i += 4)
{
v128 va = Arm.Neon.vld1q_f32(aptr);
v128 vb = Arm.Neon.vld1q_f32(bptr);
v128 vd = Arm.Neon.vsubq_f32(va, vb);
vd = Arm.Neon.vmulq_f32(vd, vd);
d += Arm.Neon.vaddvq_f32(vd);
aptr += 4;
bptr += 4;
}
// remainder
for (; i < dim; ++i)
{
float delta = *aptr - *bptr;
d += delta * delta;
aptr++;
bptr++;
}
}
else
{
for (var i = 0; i < dim; ++i)
{
float delta = a[aIndex + i] - b[bIndex + i];
d += delta * delta;
}
}
return d;
}
static unsafe void CopyElem(int dim, NativeArray<float> src, int srcIndex, NativeArray<float> dst, int dstIndex)
{
UnsafeUtility.MemCpy((float*) dst.GetUnsafePtr() + dstIndex * dim,
(float*) src.GetUnsafeReadOnlyPtr() + srcIndex * dim, dim * 4);
}
[BurstCompile]
struct ClosestDistanceInitialJob : IJobParallelFor
{
public int dim;
[ReadOnly] public NativeArray<float> data;
[ReadOnly] public NativeArray<float> means;
public NativeArray<float> minDistSq;
public int pointIndex;
public void Execute(int index)
{
if (index == pointIndex)
return;
minDistSq[index] = DistanceSquared(dim, data, index, means, 0);
}
}
[BurstCompile]
struct ClosestDistanceUpdateJob : IJobParallelFor
{
public int dim;
[ReadOnly] public NativeArray<float> data;
[ReadOnly] public NativeArray<float> means;
[ReadOnly] public NativeBitArray taken;
public NativeArray<float> minDistSq;
public int meanIndex;
public void Execute(int index)
{
if (taken.IsSet(index))
return;
float distSq = DistanceSquared(dim, data, index, means, meanIndex);
minDistSq[index] = math.min(minDistSq[index], distSq);
}
}
[BurstCompile]
struct CalcDistSqJob : IJobParallelFor
{
public const int kBatchSize = 1024;
public int dataSize;
[ReadOnly] public NativeBitArray taken;
[ReadOnly] public NativeArray<float> minDistSq;
public NativeArray<float> partialSums;
public void Execute(int batchIndex)
{
int iStart = math.min(batchIndex * kBatchSize, dataSize);
int iEnd = math.min((batchIndex + 1) * kBatchSize, dataSize);
float sum = 0;
for (int i = iStart; i < iEnd; ++i)
{
if (taken.IsSet(i))
continue;
sum += minDistSq[i];
}
partialSums[batchIndex] = sum;
}
}
[BurstCompile]
static int PickPointIndex(int dataSize, ref NativeArray<float> partialSums, ref NativeBitArray taken, ref NativeArray<float> minDistSq, float rval)
{
// Skip batches until we hit the ones that might have value to pick from: binary search for the batch
int indexL = 0;
int indexR = partialSums.Length;
while (indexL < indexR)
{
int indexM = (indexL + indexR) / 2;
if (partialSums[indexM] < rval)
indexL = indexM + 1;
else
indexR = indexM;
}
float acc = 0.0f;
if (indexL > 0)
{
acc = partialSums[indexL-1];
}
// Now search for the needed point
int pointIndex = -1;
for (int i = indexL * CalcDistSqJob.kBatchSize; i < dataSize; ++i)
{
if (taken.IsSet(i))
continue;
acc += minDistSq[i];
if (acc >= rval)
{
pointIndex = i;
break;
}
}
// If we have not found a point, pick the last available one
if (pointIndex < 0)
{
for (int i = dataSize - 1; i >= 0; --i)
{
if (taken.IsSet(i))
continue;
pointIndex = i;
break;
}
}
if (pointIndex < 0)
pointIndex = 0;
return pointIndex;
}
static void KMeansPlusPlus(int dim, int k, NativeArray<float> data, NativeArray<float> means, NativeArray<float> minDistSq, ref uint rngState)
{
using var prof = s_ProfPlusPlus.Auto();
int dataSize = data.Length / dim;
NativeBitArray taken = new NativeBitArray(dataSize, Allocator.TempJob);
// Select first mean randomly
int pointIndex = (int)(pcg_random(ref rngState) % dataSize);
taken.Set(pointIndex, true);
CopyElem(dim, data, pointIndex, means, 0);
// For each point: closest squared distance to the picked point
{
ClosestDistanceInitialJob job = new ClosestDistanceInitialJob
{
dim = dim,
data = data,
means = means,
minDistSq = minDistSq,
pointIndex = pointIndex
};
job.Schedule(dataSize, 1024).Complete();
}
int sumBatches = (dataSize + CalcDistSqJob.kBatchSize - 1) / CalcDistSqJob.kBatchSize;
NativeArray<float> partialSums = new(sumBatches, Allocator.TempJob);
int resultCount = 1;
while (resultCount < k)
{
// Find total sum of distances of not yet taken points
float distSqTotal = 0;
{
using var profPart = s_ProfInitialDistanceSum.Auto();
CalcDistSqJob job = new CalcDistSqJob
{
dataSize = dataSize,
taken = taken,
minDistSq = minDistSq,
partialSums = partialSums
};
job.Schedule(sumBatches, 1).Complete();
for (int i = 0; i < sumBatches; ++i)
{
distSqTotal += partialSums[i];
partialSums[i] = distSqTotal;
}
}
// Pick a non-taken point, with a probability proportional
// to distance: points furthest from any cluster are picked more.
{
using var profPart = s_ProfInitialPickPoint.Auto();
float rval = pcg_hash_float(rngState + (uint)resultCount, distSqTotal);
pointIndex = PickPointIndex(dataSize, ref partialSums, ref taken, ref minDistSq, rval);
}
// Take this point as a new cluster mean
taken.Set(pointIndex, true);
CopyElem(dim, data, pointIndex, means, resultCount);
++resultCount;
if (resultCount < k)
{
// Update distances of the points: since it tracks closest one,
// calculate distance to the new cluster and update if smaller.
using var profPart = s_ProfInitialDistanceUpdate.Auto();
ClosestDistanceUpdateJob job = new ClosestDistanceUpdateJob
{
dim = dim,
data = data,
means = means,
minDistSq = minDistSq,
taken = taken,
meanIndex = resultCount - 1
};
job.Schedule(dataSize, 256).Complete();
}
}
taken.Dispose();
partialSums.Dispose();
}
// For each data point, find cluster index that is closest to it
[BurstCompile]
struct AssignClustersJob : IJobParallelFor
{
public int indexOffset;
public int dim;
[ReadOnly] public NativeArray<float> data;
[ReadOnly] public NativeArray<float> means;
[NativeDisableParallelForRestriction] public NativeArray<int> clusters;
[NativeDisableContainerSafetyRestriction] [NativeDisableParallelForRestriction] public NativeArray<float> distances;
public void Execute(int index)
{
index += indexOffset;
int meansCount = means.Length / dim;
float minDist = float.MaxValue;
int minIndex = 0;
for (int i = 0; i < meansCount; ++i)
{
float dist = DistanceSquared(dim, data, index, means, i);
if (dist < minDist)
{
minIndex = i;
minDist = dist;
}
}
clusters[index] = minIndex;
if (distances.IsCreated)
distances[index] = minDist;
}
}
static void MakeRandomBatch(int dim, NativeArray<float> inputData, ref uint rngState, NativeArray<float> outBatch)
{
var job = new MakeBatchJob
{
m_Dim = dim,
m_InputData = inputData,
m_Seed = pcg_random(ref rngState),
m_OutBatch = outBatch
};
job.Schedule().Complete();
}
[BurstCompile]
struct MakeBatchJob : IJob
{
public int m_Dim;
public NativeArray<float> m_InputData;
public NativeArray<float> m_OutBatch;
public uint m_Seed;
public void Execute()
{
uint dataSize = (uint)(m_InputData.Length / m_Dim);
int batchSize = m_OutBatch.Length / m_Dim;
NativeHashSet<int> picked = new(batchSize, Allocator.Temp);
while (picked.Count < batchSize)
{
int index = (int)(pcg_hash(m_Seed++) % dataSize);
if (!picked.Contains(index))
{
CopyElem(m_Dim, m_InputData, index, m_OutBatch, picked.Count);
picked.Add(index);
}
}
picked.Dispose();
}
}
[BurstCompile]
struct UpdateCentroidsJob : IJob
{
public int m_Dim;
public int m_BatchSize;
[ReadOnly] public NativeArray<int> m_BatchClusters;
public NativeArray<float> m_Counts;
[ReadOnly] public NativeArray<float> m_BatchPoints;
public NativeArray<float> m_Clusters;
public void Execute()
{
for (int i = 0; i < m_BatchSize; ++i)
{
int clusterIndex = m_BatchClusters[i];
m_Counts[clusterIndex]++;
float alpha = 1.0f / m_Counts[clusterIndex];
for (int j = 0; j < m_Dim; ++j)
{
m_Clusters[clusterIndex * m_Dim + j] = math.lerp(m_Clusters[clusterIndex * m_Dim + j],
m_BatchPoints[i * m_Dim + j], alpha);
}
}
}
}
static bool InitializeCentroids(int dim, NativeArray<float> inputData, int initBatchSize, ref uint rngState, int initAttempts, NativeArray<float> outClusters, Func<float,bool> progress)
{
using var prof = s_ProfPlusPlus.Auto();
int k = outClusters.Length / dim;
int dataSize = inputData.Length / dim;
initBatchSize = math.min(initBatchSize, dataSize);
NativeArray<float> centroidBatch = new(initBatchSize * dim, Allocator.TempJob);
NativeArray<float> validationBatch = new(initBatchSize * dim, Allocator.TempJob);
MakeRandomBatch(dim, inputData, ref rngState, centroidBatch);
MakeRandomBatch(dim, inputData, ref rngState, validationBatch);
NativeArray<int> tmpIndices = new(initBatchSize, Allocator.TempJob);
NativeArray<float> tmpDistances = new(initBatchSize, Allocator.TempJob);
NativeArray<float> curCentroids = new(k * dim, Allocator.TempJob);
float minDistSum = float.MaxValue;
bool cancelled = false;
for (int ia = 0; ia < initAttempts; ++ia)
{
if (progress != null && !progress((float) ia / initAttempts * 0.3f))
{
cancelled = true;
break;
}
KMeansPlusPlus(dim, k, centroidBatch, curCentroids, tmpDistances, ref rngState);
{
using var profPart = s_ProfAssignClusters.Auto();
AssignClustersJob job = new AssignClustersJob
{
dim = dim,
data = validationBatch,
means = curCentroids,
indexOffset = 0,
clusters = tmpIndices,
distances = tmpDistances
};
job.Schedule(initBatchSize, 1).Complete();
}
float distSum = 0;
foreach (var d in tmpDistances)
distSum += d;
// is this centroid better?
if (distSum < minDistSum)
{
minDistSum = distSum;
outClusters.CopyFrom(curCentroids);
}
}
centroidBatch.Dispose();
validationBatch.Dispose();
tmpDistances.Dispose();
tmpIndices.Dispose();
curCentroids.Dispose();
return !cancelled;
}
// https://www.reedbeta.com/blog/hash-functions-for-gpu-rendering/
static uint pcg_hash(uint input)
{
uint state = input * 747796405u + 2891336453u;
uint word = ((state >> (int)((state >> 28) + 4u)) ^ state) * 277803737u;
return (word >> 22) ^ word;
}
static float pcg_hash_float(uint input, float upTo)
{
uint val = pcg_hash(input);
float f = math.asfloat(0x3f800000 | (val >> 9)) - 1.0f;
return f * upTo;
}
static uint pcg_random(ref uint rng_state)
{
uint state = rng_state;
rng_state = rng_state * 747796405u + 2891336453u;
uint word = ((state >> (int)((state >> 28) + 4u)) ^ state) * 277803737u;
return (word >> 22) ^ word;
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 9cecadf9c980a4ad9a30d0e1ae09d16a
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,107 @@
// SPDX-License-Identifier: MIT
using System;
using System.Collections.Generic;
using System.IO;
using System.Linq;
using System.Text;
using Unity.Collections;
namespace GaussianSplatting.Editor.Utils
{
public static class PLYFileReader
{
public static void ReadFileHeader(string filePath, out int vertexCount, out int vertexStride, out List<string> attrNames)
{
vertexCount = 0;
vertexStride = 0;
attrNames = new List<string>();
if (!File.Exists(filePath))
return;
using var fs = new FileStream(filePath, FileMode.Open, FileAccess.Read);
ReadHeaderImpl(filePath, out vertexCount, out vertexStride, out attrNames, fs);
}
static void ReadHeaderImpl(string filePath, out int vertexCount, out int vertexStride, out List<string> attrNames, FileStream fs)
{
// C# arrays and NativeArrays make it hard to have a "byte" array larger than 2GB :/
if (fs.Length >= 2 * 1024 * 1024 * 1024L)
throw new IOException($"PLY {filePath} read error: currently files larger than 2GB are not supported");
// read header
vertexCount = 0;
vertexStride = 0;
attrNames = new List<string>();
const int kMaxHeaderLines = 9000;
for (int lineIdx = 0; lineIdx < kMaxHeaderLines; ++lineIdx)
{
var line = ReadLine(fs);
if (line == "end_header" || line.Length == 0)
break;
var tokens = line.Split(' ');
if (tokens.Length == 3 && tokens[0] == "element" && tokens[1] == "vertex")
vertexCount = int.Parse(tokens[2]);
if (tokens.Length == 3 && tokens[0] == "property")
{
ElementType type = tokens[1] switch
{
"float" => ElementType.Float,
"double" => ElementType.Double,
"uchar" => ElementType.UChar,
_ => ElementType.None
};
vertexStride += TypeToSize(type);
attrNames.Add(tokens[2]);
}
}
//Debug.Log($"PLY {filePath} vtx {vertexCount} stride {vertexStride} attrs #{attrNames.Count} {string.Join(',', attrNames)}");
}
public static void ReadFile(string filePath, out int vertexCount, out int vertexStride, out List<string> attrNames, out NativeArray<byte> vertices)
{
using var fs = new FileStream(filePath, FileMode.Open, FileAccess.Read);
ReadHeaderImpl(filePath, out vertexCount, out vertexStride, out attrNames, fs);
vertices = new NativeArray<byte>(vertexCount * vertexStride, Allocator.Persistent);
var readBytes = fs.Read(vertices);
if (readBytes != vertices.Length)
throw new IOException($"PLY {filePath} read error, expected {vertices.Length} data bytes got {readBytes}");
}
enum ElementType
{
None,
Float,
Double,
UChar
}
static int TypeToSize(ElementType t)
{
return t switch
{
ElementType.None => 0,
ElementType.Float => 4,
ElementType.Double => 8,
ElementType.UChar => 1,
_ => throw new ArgumentOutOfRangeException(nameof(t), t, null)
};
}
static string ReadLine(FileStream fs)
{
var byteBuffer = new List<byte>();
while (true)
{
int b = fs.ReadByte();
if (b == -1 || b == '\n')
break;
byteBuffer.Add((byte)b);
}
// if line had CRLF line endings, remove the CR part
if (byteBuffer.Count > 0 && byteBuffer.Last() == '\r')
byteBuffer.RemoveAt(byteBuffer.Count-1);
return Encoding.UTF8.GetString(byteBuffer.ToArray());
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 27964c85711004ddca73909489af2e2e
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,403 @@
/*
Embedded TinyJSON from https://github.com/pbhogan/TinyJSON (version 71fed96, 2019 May 10) directly here, with
custom namespace wrapped around it so it does not clash with any other embedded TinyJSON. Original license:
The MIT License (MIT)
Copyright (c) 2018 Alex Parker
Permission is hereby granted, free of charge, to any person obtaining a copy of
this software and associated documentation files (the "Software"), to deal in
the Software without restriction, including without limitation the rights to
use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of
the Software, and to permit persons to whom the Software is furnished to do so,
subject to the following conditions:
The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN
CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
using System;
using System.Collections;
using System.Collections.Generic;
using System.Reflection;
using System.Runtime.Serialization;
using System.Text;
namespace GaussianSplatting.Editor.Utils
{
// Really simple JSON parser in ~300 lines
// - Attempts to parse JSON files with minimal GC allocation
// - Nice and simple "[1,2,3]".FromJson<List<int>>() API
// - Classes and structs can be parsed too!
// class Foo { public int Value; }
// "{\"Value\":10}".FromJson<Foo>()
// - Can parse JSON without type information into Dictionary<string,object> and List<object> e.g.
// "[1,2,3]".FromJson<object>().GetType() == typeof(List<object>)
// "{\"Value\":10}".FromJson<object>().GetType() == typeof(Dictionary<string,object>)
// - No JIT Emit support to support AOT compilation on iOS
// - Attempts are made to NOT throw an exception if the JSON is corrupted or invalid: returns null instead.
// - Only public fields and property setters on classes/structs will be written to
//
// Limitations:
// - No JIT Emit support to parse structures quickly
// - Limited to parsing <2GB JSON files (due to int.MaxValue)
// - Parsing of abstract classes or interfaces is NOT supported and will throw an exception.
public static class JSONParser
{
[ThreadStatic] static Stack<List<string>> splitArrayPool;
[ThreadStatic] static StringBuilder stringBuilder;
[ThreadStatic] static Dictionary<Type, Dictionary<string, FieldInfo>> fieldInfoCache;
[ThreadStatic] static Dictionary<Type, Dictionary<string, PropertyInfo>> propertyInfoCache;
public static T FromJson<T>(this string json)
{
// Initialize, if needed, the ThreadStatic variables
if (propertyInfoCache == null) propertyInfoCache = new Dictionary<Type, Dictionary<string, PropertyInfo>>();
if (fieldInfoCache == null) fieldInfoCache = new Dictionary<Type, Dictionary<string, FieldInfo>>();
if (stringBuilder == null) stringBuilder = new StringBuilder();
if (splitArrayPool == null) splitArrayPool = new Stack<List<string>>();
//Remove all whitespace not within strings to make parsing simpler
stringBuilder.Length = 0;
for (int i = 0; i < json.Length; i++)
{
char c = json[i];
if (c == '"')
{
i = AppendUntilStringEnd(true, i, json);
continue;
}
if (char.IsWhiteSpace(c))
continue;
stringBuilder.Append(c);
}
//Parse the thing!
return (T)ParseValue(typeof(T), stringBuilder.ToString());
}
static int AppendUntilStringEnd(bool appendEscapeCharacter, int startIdx, string json)
{
stringBuilder.Append(json[startIdx]);
for (int i = startIdx + 1; i < json.Length; i++)
{
if (json[i] == '\\')
{
if (appendEscapeCharacter)
stringBuilder.Append(json[i]);
stringBuilder.Append(json[i + 1]);
i++;//Skip next character as it is escaped
}
else if (json[i] == '"')
{
stringBuilder.Append(json[i]);
return i;
}
else
stringBuilder.Append(json[i]);
}
return json.Length - 1;
}
//Splits { <value>:<value>, <value>:<value> } and [ <value>, <value> ] into a list of <value> strings
static List<string> Split(string json)
{
List<string> splitArray = splitArrayPool.Count > 0 ? splitArrayPool.Pop() : new List<string>();
splitArray.Clear();
if (json.Length == 2)
return splitArray;
int parseDepth = 0;
stringBuilder.Length = 0;
for (int i = 1; i < json.Length - 1; i++)
{
switch (json[i])
{
case '[':
case '{':
parseDepth++;
break;
case ']':
case '}':
parseDepth--;
break;
case '"':
i = AppendUntilStringEnd(true, i, json);
continue;
case ',':
case ':':
if (parseDepth == 0)
{
splitArray.Add(stringBuilder.ToString());
stringBuilder.Length = 0;
continue;
}
break;
}
stringBuilder.Append(json[i]);
}
splitArray.Add(stringBuilder.ToString());
return splitArray;
}
internal static object ParseValue(Type type, string json)
{
if (type == typeof(string))
{
if (json.Length <= 2)
return string.Empty;
StringBuilder parseStringBuilder = new StringBuilder(json.Length);
for (int i = 1; i < json.Length - 1; ++i)
{
if (json[i] == '\\' && i + 1 < json.Length - 1)
{
int j = "\"\\nrtbf/".IndexOf(json[i + 1]);
if (j >= 0)
{
parseStringBuilder.Append("\"\\\n\r\t\b\f/"[j]);
++i;
continue;
}
if (json[i + 1] == 'u' && i + 5 < json.Length - 1)
{
UInt32 c = 0;
if (UInt32.TryParse(json.Substring(i + 2, 4), System.Globalization.NumberStyles.AllowHexSpecifier, null, out c))
{
parseStringBuilder.Append((char)c);
i += 5;
continue;
}
}
}
parseStringBuilder.Append(json[i]);
}
return parseStringBuilder.ToString();
}
if (type.IsPrimitive)
{
var result = Convert.ChangeType(json, type, System.Globalization.CultureInfo.InvariantCulture);
return result;
}
if (type == typeof(decimal))
{
decimal result;
decimal.TryParse(json, System.Globalization.NumberStyles.Float, System.Globalization.CultureInfo.InvariantCulture, out result);
return result;
}
if (type == typeof(DateTime))
{
DateTime result;
DateTime.TryParse(json.Replace("\"",""), System.Globalization.CultureInfo.InvariantCulture, System.Globalization.DateTimeStyles.None, out result);
return result;
}
if (json == "null")
{
return null;
}
if (type.IsEnum)
{
if (json[0] == '"')
json = json.Substring(1, json.Length - 2);
try
{
return Enum.Parse(type, json, false);
}
catch
{
return 0;
}
}
if (type.IsArray)
{
Type arrayType = type.GetElementType();
if (json[0] != '[' || json[json.Length - 1] != ']')
return null;
List<string> elems = Split(json);
Array newArray = Array.CreateInstance(arrayType, elems.Count);
for (int i = 0; i < elems.Count; i++)
newArray.SetValue(ParseValue(arrayType, elems[i]), i);
splitArrayPool.Push(elems);
return newArray;
}
if (type.IsGenericType && type.GetGenericTypeDefinition() == typeof(List<>))
{
Type listType = type.GetGenericArguments()[0];
if (json[0] != '[' || json[json.Length - 1] != ']')
return null;
List<string> elems = Split(json);
var list = (IList)type.GetConstructor(new Type[] { typeof(int) }).Invoke(new object[] { elems.Count });
for (int i = 0; i < elems.Count; i++)
list.Add(ParseValue(listType, elems[i]));
splitArrayPool.Push(elems);
return list;
}
if (type.IsGenericType && type.GetGenericTypeDefinition() == typeof(Dictionary<,>))
{
Type keyType, valueType;
{
Type[] args = type.GetGenericArguments();
keyType = args[0];
valueType = args[1];
}
//Refuse to parse dictionary keys that aren't of type string
if (keyType != typeof(string))
return null;
//Must be a valid dictionary element
if (json[0] != '{' || json[json.Length - 1] != '}')
return null;
//The list is split into key/value pairs only, this means the split must be divisible by 2 to be valid JSON
List<string> elems = Split(json);
if (elems.Count % 2 != 0)
return null;
var dictionary = (IDictionary)type.GetConstructor(new Type[] { typeof(int) }).Invoke(new object[] { elems.Count / 2 });
for (int i = 0; i < elems.Count; i += 2)
{
if (elems[i].Length <= 2)
continue;
string keyValue = elems[i].Substring(1, elems[i].Length - 2);
object val = ParseValue(valueType, elems[i + 1]);
dictionary[keyValue] = val;
}
return dictionary;
}
if (type == typeof(object))
{
return ParseAnonymousValue(json);
}
if (json[0] == '{' && json[json.Length - 1] == '}')
{
return ParseObject(type, json);
}
return null;
}
static object ParseAnonymousValue(string json)
{
if (json.Length == 0)
return null;
if (json[0] == '{' && json[json.Length - 1] == '}')
{
List<string> elems = Split(json);
if (elems.Count % 2 != 0)
return null;
var dict = new Dictionary<string, object>(elems.Count / 2);
for (int i = 0; i < elems.Count; i += 2)
dict[elems[i].Substring(1, elems[i].Length - 2)] = ParseAnonymousValue(elems[i + 1]);
return dict;
}
if (json[0] == '[' && json[json.Length - 1] == ']')
{
List<string> items = Split(json);
var finalList = new List<object>(items.Count);
for (int i = 0; i < items.Count; i++)
finalList.Add(ParseAnonymousValue(items[i]));
return finalList;
}
if (json[0] == '"' && json[json.Length - 1] == '"')
{
string str = json.Substring(1, json.Length - 2);
return str.Replace("\\", string.Empty);
}
if (char.IsDigit(json[0]) || json[0] == '-')
{
if (json.Contains("."))
{
double result;
double.TryParse(json, System.Globalization.NumberStyles.Float, System.Globalization.CultureInfo.InvariantCulture, out result);
return result;
}
else
{
int result;
int.TryParse(json, out result);
return result;
}
}
if (json == "true")
return true;
if (json == "false")
return false;
// handles json == "null" as well as invalid JSON
return null;
}
static Dictionary<string, T> CreateMemberNameDictionary<T>(T[] members) where T : MemberInfo
{
Dictionary<string, T> nameToMember = new Dictionary<string, T>(StringComparer.OrdinalIgnoreCase);
for (int i = 0; i < members.Length; i++)
{
T member = members[i];
if (member.IsDefined(typeof(IgnoreDataMemberAttribute), true))
continue;
string name = member.Name;
if (member.IsDefined(typeof(DataMemberAttribute), true))
{
DataMemberAttribute dataMemberAttribute = (DataMemberAttribute)Attribute.GetCustomAttribute(member, typeof(DataMemberAttribute), true);
if (!string.IsNullOrEmpty(dataMemberAttribute.Name))
name = dataMemberAttribute.Name;
}
nameToMember.Add(name, member);
}
return nameToMember;
}
static object ParseObject(Type type, string json)
{
object instance = FormatterServices.GetUninitializedObject(type);
//The list is split into key/value pairs only, this means the split must be divisible by 2 to be valid JSON
List<string> elems = Split(json);
if (elems.Count % 2 != 0)
return instance;
Dictionary<string, FieldInfo> nameToField;
Dictionary<string, PropertyInfo> nameToProperty;
if (!fieldInfoCache.TryGetValue(type, out nameToField))
{
nameToField = CreateMemberNameDictionary(type.GetFields(BindingFlags.Instance | BindingFlags.Public | BindingFlags.FlattenHierarchy));
fieldInfoCache.Add(type, nameToField);
}
if (!propertyInfoCache.TryGetValue(type, out nameToProperty))
{
nameToProperty = CreateMemberNameDictionary(type.GetProperties(BindingFlags.Instance | BindingFlags.Public | BindingFlags.FlattenHierarchy));
propertyInfoCache.Add(type, nameToProperty);
}
for (int i = 0; i < elems.Count; i += 2)
{
if (elems[i].Length <= 2)
continue;
string key = elems[i].Substring(1, elems[i].Length - 2);
string value = elems[i + 1];
FieldInfo fieldInfo;
PropertyInfo propertyInfo;
if (nameToField.TryGetValue(key, out fieldInfo))
fieldInfo.SetValue(instance, ParseValue(fieldInfo.FieldType, value));
else if (nameToProperty.TryGetValue(key, out propertyInfo))
propertyInfo.SetValue(instance, ParseValue(propertyInfo.PropertyType, value), null);
}
return instance;
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: e9ea5041388393f459c378c31e4d7b1f
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

21
MVS/3DGS-Unity/LICENSE.md Normal file
View File

@@ -0,0 +1,21 @@
MIT License
Copyright (c) 2023 Aras Pranckevičius
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: da286c32b8dba1744aecca8cb1ab4ad6
TextScriptImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,8 @@
fileFormatVersion: 2
guid: 78bfe028c2744c741bd4f94574de884a
folderAsset: yes
DefaultImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,30 @@
%YAML 1.1
%TAG !u! tag:unity3d.com,2011:
--- !u!21 &2100000
Material:
serializedVersion: 8
m_ObjectHideFlags: 0
m_CorrespondingSourceObject: {fileID: 0}
m_PrefabInstance: {fileID: 0}
m_PrefabAsset: {fileID: 0}
m_Name: BlackSkybox
m_Shader: {fileID: 4800000, guid: a4867e5be68354ccda78062a92c74391, type: 3}
m_Parent: {fileID: 0}
m_ModifiedSerializedProperties: 0
m_ValidKeywords: []
m_InvalidKeywords: []
m_LightmapFlags: 4
m_EnableInstancingVariants: 0
m_DoubleSidedGI: 0
m_CustomRenderQueue: -1
stringTagMap: {}
disabledShaderPasses: []
m_LockedProperties:
m_SavedProperties:
serializedVersion: 3
m_TexEnvs: []
m_Ints: []
m_Floats: []
m_Colors:
- _Color: {r: 0, g: 0, b: 0, a: 0}
m_BuildTextureStacks: []

View File

@@ -0,0 +1,8 @@
fileFormatVersion: 2
guid: e9c9951d4a35e4a54812fa0280fa548c
NativeFormatImporter:
externalObjects: {}
mainObjectFileID: 2100000
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,8 @@
fileFormatVersion: 2
guid: 71627bcf67390da439d82a2a05a57bb4
folderAsset: yes
DefaultImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,77 @@
// SPDX-License-Identifier: MIT
using System.Linq;
using UnityEditor;
using UnityEngine;
namespace GaussianSplatting.Runtime
{
public class GaussianCutout : MonoBehaviour
{
public enum Type
{
Ellipsoid,
Box
}
public Type m_Type = Type.Ellipsoid;
public bool m_Invert = false;
public struct ShaderData // match GaussianCutoutShaderData in CS
{
public Matrix4x4 matrix;
public uint typeAndFlags;
}
public static ShaderData GetShaderData(GaussianCutout self, Matrix4x4 rendererMatrix)
{
ShaderData sd = default;
if (self && self.isActiveAndEnabled)
{
var tr = self.transform;
sd.matrix = tr.worldToLocalMatrix * rendererMatrix;
sd.typeAndFlags = ((uint)self.m_Type) | (self.m_Invert ? 0x100u : 0u);
}
else
{
sd.typeAndFlags = ~0u;
}
return sd;
}
#if UNITY_EDITOR
public void OnDrawGizmos()
{
Gizmos.matrix = transform.localToWorldMatrix;
var color = Color.magenta;
color.a = 0.2f;
if (Selection.Contains(gameObject))
color.a = 0.9f;
else
{
// mid amount of alpha if a GS object that contains us as a cutout is selected
var activeGo = Selection.activeGameObject;
if (activeGo != null)
{
var activeSplat = activeGo.GetComponent<GaussianSplatRenderer>();
if (activeSplat != null)
{
if (activeSplat.m_Cutouts != null && activeSplat.m_Cutouts.Contains(this))
color.a = 0.5f;
}
}
}
Gizmos.color = color;
if (m_Type == Type.Ellipsoid)
{
Gizmos.DrawWireSphere(Vector3.zero, 1.0f);
}
if (m_Type == Type.Box)
{
Gizmos.DrawWireCube(Vector3.zero, Vector3.one * 2);
}
}
#endif // #if UNITY_EDITOR
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 2c57a1c501bd05549ae157cc474bd4c4
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,247 @@
// SPDX-License-Identifier: MIT
using System;
using Unity.Collections.LowLevel.Unsafe;
using Unity.Mathematics;
using UnityEngine;
using UnityEngine.Experimental.Rendering;
namespace GaussianSplatting.Runtime
{
public class GaussianSplatAsset : ScriptableObject
{
public const int kCurrentVersion = 2023_10_20;
public const int kChunkSize = 256;
public const int kTextureWidth = 2048; // allows up to 32M splats on desktop GPU (2k width x 16k height)
public const int kMaxSplats = 8_600_000; // mostly due to 2GB GPU buffer size limit when exporting a splat (2GB / 248B is just over 8.6M)
[SerializeField] int m_FormatVersion;
[SerializeField] int m_SplatCount;
[SerializeField] Vector3 m_BoundsMin;
[SerializeField] Vector3 m_BoundsMax;
[SerializeField] Hash128 m_DataHash;
public int formatVersion => m_FormatVersion;
public int splatCount => m_SplatCount;
public Vector3 boundsMin => m_BoundsMin;
public Vector3 boundsMax => m_BoundsMax;
public Hash128 dataHash => m_DataHash;
// Match VECTOR_FMT_* in HLSL
public enum VectorFormat
{
Float32, // 12 bytes: 32F.32F.32F
Norm16, // 6 bytes: 16.16.16
Norm11, // 4 bytes: 11.10.11
Norm6 // 2 bytes: 6.5.5
}
public static int GetVectorSize(VectorFormat fmt)
{
return fmt switch
{
VectorFormat.Float32 => 12,
VectorFormat.Norm16 => 6,
VectorFormat.Norm11 => 4,
VectorFormat.Norm6 => 2,
_ => throw new ArgumentOutOfRangeException(nameof(fmt), fmt, null)
};
}
public enum ColorFormat
{
Float32x4,
Float16x4,
Norm8x4,
BC7,
}
public static int GetColorSize(ColorFormat fmt)
{
return fmt switch
{
ColorFormat.Float32x4 => 16,
ColorFormat.Float16x4 => 8,
ColorFormat.Norm8x4 => 4,
ColorFormat.BC7 => 1,
_ => throw new ArgumentOutOfRangeException(nameof(fmt), fmt, null)
};
}
public enum SHFormat
{
Float32,
Float16,
Norm11,
Norm6,
Cluster64k,
Cluster32k,
Cluster16k,
Cluster8k,
Cluster4k,
}
public struct SHTableItemFloat32
{
public float3 sh1, sh2, sh3, sh4, sh5, sh6, sh7, sh8, sh9, shA, shB, shC, shD, shE, shF;
public float3 shPadding; // pad to multiple of 16 bytes
}
public struct SHTableItemFloat16
{
public half3 sh1, sh2, sh3, sh4, sh5, sh6, sh7, sh8, sh9, shA, shB, shC, shD, shE, shF;
public half3 shPadding; // pad to multiple of 16 bytes
}
public struct SHTableItemNorm11
{
public uint sh1, sh2, sh3, sh4, sh5, sh6, sh7, sh8, sh9, shA, shB, shC, shD, shE, shF;
}
public struct SHTableItemNorm6
{
public ushort sh1, sh2, sh3, sh4, sh5, sh6, sh7, sh8, sh9, shA, shB, shC, shD, shE, shF;
public ushort shPadding; // pad to multiple of 4 bytes
}
public void Initialize(int splats, VectorFormat formatPos, VectorFormat formatScale, ColorFormat formatColor, SHFormat formatSh, Vector3 bMin, Vector3 bMax, CameraInfo[] cameraInfos)
{
m_SplatCount = splats;
m_FormatVersion = kCurrentVersion;
m_PosFormat = formatPos;
m_ScaleFormat = formatScale;
m_ColorFormat = formatColor;
m_SHFormat = formatSh;
m_Cameras = cameraInfos;
m_BoundsMin = bMin;
m_BoundsMax = bMax;
}
public void SetDataHash(Hash128 hash)
{
m_DataHash = hash;
}
public void SetAssetFiles(TextAsset dataChunk, TextAsset dataPos, TextAsset dataOther, TextAsset dataColor, TextAsset dataSh)
{
m_ChunkData = dataChunk;
m_PosData = dataPos;
m_OtherData = dataOther;
m_ColorData = dataColor;
m_SHData = dataSh;
}
public static int GetOtherSizeNoSHIndex(VectorFormat scaleFormat)
{
return 4 + GetVectorSize(scaleFormat);
}
public static int GetSHCount(SHFormat fmt, int splatCount)
{
return fmt switch
{
SHFormat.Float32 => splatCount,
SHFormat.Float16 => splatCount,
SHFormat.Norm11 => splatCount,
SHFormat.Norm6 => splatCount,
SHFormat.Cluster64k => 64 * 1024,
SHFormat.Cluster32k => 32 * 1024,
SHFormat.Cluster16k => 16 * 1024,
SHFormat.Cluster8k => 8 * 1024,
SHFormat.Cluster4k => 4 * 1024,
_ => throw new ArgumentOutOfRangeException(nameof(fmt), fmt, null)
};
}
public static (int,int) CalcTextureSize(int splatCount)
{
int width = kTextureWidth;
int height = math.max(1, (splatCount + width - 1) / width);
// our swizzle tiles are 16x16, so make texture multiple of that height
int blockHeight = 16;
height = (height + blockHeight - 1) / blockHeight * blockHeight;
return (width, height);
}
public static GraphicsFormat ColorFormatToGraphics(ColorFormat format)
{
return format switch
{
ColorFormat.Float32x4 => GraphicsFormat.R32G32B32A32_SFloat,
ColorFormat.Float16x4 => GraphicsFormat.R16G16B16A16_SFloat,
ColorFormat.Norm8x4 => GraphicsFormat.R8G8B8A8_UNorm,
ColorFormat.BC7 => GraphicsFormat.RGBA_BC7_UNorm,
_ => throw new ArgumentOutOfRangeException(nameof(format), format, null)
};
}
public static long CalcPosDataSize(int splatCount, VectorFormat formatPos)
{
return splatCount * GetVectorSize(formatPos);
}
public static long CalcOtherDataSize(int splatCount, VectorFormat formatScale)
{
return splatCount * GetOtherSizeNoSHIndex(formatScale);
}
public static long CalcColorDataSize(int splatCount, ColorFormat formatColor)
{
var (width, height) = CalcTextureSize(splatCount);
return width * height * GetColorSize(formatColor);
}
public static long CalcSHDataSize(int splatCount, SHFormat formatSh)
{
int shCount = GetSHCount(formatSh, splatCount);
return formatSh switch
{
SHFormat.Float32 => shCount * UnsafeUtility.SizeOf<SHTableItemFloat32>(),
SHFormat.Float16 => shCount * UnsafeUtility.SizeOf<SHTableItemFloat16>(),
SHFormat.Norm11 => shCount * UnsafeUtility.SizeOf<SHTableItemNorm11>(),
SHFormat.Norm6 => shCount * UnsafeUtility.SizeOf<SHTableItemNorm6>(),
_ => shCount * UnsafeUtility.SizeOf<SHTableItemFloat16>() + splatCount * 2
};
}
public static long CalcChunkDataSize(int splatCount)
{
int chunkCount = (splatCount + kChunkSize - 1) / kChunkSize;
return chunkCount * UnsafeUtility.SizeOf<ChunkInfo>();
}
[SerializeField] VectorFormat m_PosFormat = VectorFormat.Norm11;
[SerializeField] VectorFormat m_ScaleFormat = VectorFormat.Norm11;
[SerializeField] SHFormat m_SHFormat = SHFormat.Norm11;
[SerializeField] ColorFormat m_ColorFormat;
[SerializeField] TextAsset m_PosData;
[SerializeField] TextAsset m_ColorData;
[SerializeField] TextAsset m_OtherData;
[SerializeField] TextAsset m_SHData;
// Chunk data is optional (if data formats are fully lossless then there's no chunking)
[SerializeField] TextAsset m_ChunkData;
[SerializeField] CameraInfo[] m_Cameras;
public VectorFormat posFormat => m_PosFormat;
public VectorFormat scaleFormat => m_ScaleFormat;
public SHFormat shFormat => m_SHFormat;
public ColorFormat colorFormat => m_ColorFormat;
public TextAsset posData => m_PosData;
public TextAsset colorData => m_ColorData;
public TextAsset otherData => m_OtherData;
public TextAsset shData => m_SHData;
public TextAsset chunkData => m_ChunkData;
public CameraInfo[] cameras => m_Cameras;
public struct ChunkInfo
{
public uint colR, colG, colB, colA;
public float2 posX, posY, posZ;
public uint sclX, sclY, sclZ;
public uint shR, shG, shB;
}
[Serializable]
public struct CameraInfo
{
public Vector3 pos;
public Vector3 axisX, axisY, axisZ;
public float fov;
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 33b71fae31e6c7d438e8566dc713e666
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,59 @@
// SPDX-License-Identifier: MIT
#if GS_ENABLE_HDRP
using UnityEngine;
using UnityEngine.Rendering.HighDefinition;
using UnityEngine.Rendering;
using UnityEngine.Experimental.Rendering;
namespace GaussianSplatting.Runtime
{
// Note: I have no idea what is the proper usage of CustomPass.
// Code below "seems to work" but I'm just fumbling along, without understanding any of it.
class GaussianSplatHDRPPass : CustomPass
{
RTHandle m_RenderTarget;
// It can be used to configure render targets and their clear state. Also to create temporary render target textures.
// When empty this render pass will render to the active camera render target.
// You should never call CommandBuffer.SetRenderTarget. Instead call <c>ConfigureTarget</c> and <c>ConfigureClear</c>.
// The render pipeline will ensure target setup and clearing happens in an performance manner.
protected override void Setup(ScriptableRenderContext renderContext, CommandBuffer cmd)
{
m_RenderTarget = RTHandles.Alloc(Vector2.one,
colorFormat: GraphicsFormat.R16G16B16A16_SFloat, useDynamicScale: true,
depthBufferBits: DepthBits.None, msaaSamples: MSAASamples.None,
filterMode: FilterMode.Point, wrapMode: TextureWrapMode.Clamp, name: "_GaussianSplatRT");
}
protected override void Execute(CustomPassContext ctx)
{
var cam = ctx.hdCamera.camera;
var system = GaussianSplatRenderSystem.instance;
if (!system.GatherSplatsForCamera(cam))
return;
ctx.cmd.SetGlobalTexture(m_RenderTarget.name, m_RenderTarget.nameID);
CoreUtils.SetRenderTarget(ctx.cmd, m_RenderTarget, ctx.cameraDepthBuffer, ClearFlag.Color,
new Color(0, 0, 0, 0));
// add sorting, view calc and drawing commands for each splat object
Material matComposite =
GaussianSplatRenderSystem.instance.SortAndRenderSplats(ctx.hdCamera.camera, ctx.cmd);
// compose
ctx.cmd.BeginSample(GaussianSplatRenderSystem.s_ProfCompose);
CoreUtils.SetRenderTarget(ctx.cmd, ctx.cameraColorBuffer, ClearFlag.None);
CoreUtils.DrawFullScreen(ctx.cmd, matComposite, ctx.propertyBlock, shaderPassId: 0);
ctx.cmd.EndSample(GaussianSplatRenderSystem.s_ProfCompose);
}
protected override void Cleanup()
{
m_RenderTarget.Release();
}
}
}
#endif

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: f40f16e78da87c646826cc5335ccb1f8
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

File diff suppressed because it is too large Load Diff

View File

@@ -0,0 +1,17 @@
fileFormatVersion: 2
guid: c12d929a7f62c48adbe9ff45c9a33ff0
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences:
- m_ShaderSplats: {fileID: 4800000, guid: ed800126ae8844a67aad1974ddddd59c, type: 3}
- m_ShaderComposite: {fileID: 4800000, guid: 7e184af7d01193a408eb916d8acafff9, type: 3}
- m_ShaderDebugPoints: {fileID: 4800000, guid: b44409fc67214394f8f47e4e2648425e, type: 3}
- m_ShaderDebugBoxes: {fileID: 4800000, guid: 4006f2680fd7c8b4cbcb881454c782be, type: 3}
- m_CSSplatUtilities: {fileID: 7200000, guid: ec84f78b836bd4f96a105d6b804f08bd, type: 3}
- m_CSFfxSort: {fileID: 7200000, guid: dec36776b6c843544a1f6f9b436a4474, type: 3}
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,158 @@
// SPDX-License-Identifier: MIT
#if GS_ENABLE_URP
using UnityEngine;
using UnityEngine.Experimental.Rendering;
using UnityEngine.Rendering;
using UnityEngine.Rendering.Universal;
#if UNITY_6000_0_OR_NEWER
using UnityEngine.Rendering.RenderGraphModule;
#endif
namespace GaussianSplatting.Runtime
{
// Note: I have no idea what is the purpose of ScriptableRendererFeature vs ScriptableRenderPass, which one of those
// is supposed to do resource management vs logic, etc. etc. Code below "seems to work" but I'm just fumbling along,
// without understanding any of it.
//
// ReSharper disable once InconsistentNaming
class GaussianSplatURPFeature : ScriptableRendererFeature
{
class GSRenderPass : ScriptableRenderPass
{
const string GaussianSplatRTName = "_GaussianSplatRT";
#if !UNITY_6000_0_OR_NEWER
RTHandle m_RenderTarget;
internal ScriptableRenderer m_Renderer;
internal CommandBuffer m_Cmb;
#endif
public void Dispose()
{
#if !UNITY_6000_0_OR_NEWER
m_RenderTarget?.Release();
#endif
}
#if UNITY_6000_0_OR_NEWER
const string ProfilerTag = "GaussianSplatRenderGraph";
static readonly ProfilingSampler s_profilingSampler = new(ProfilerTag);
static readonly int s_gaussianSplatRT = Shader.PropertyToID(GaussianSplatRTName);
class PassData
{
internal UniversalCameraData CameraData;
internal TextureHandle SourceTexture;
internal TextureHandle SourceDepth;
internal TextureHandle GaussianSplatRT;
}
public override void RecordRenderGraph(RenderGraph renderGraph, ContextContainer frameData)
{
using var builder = renderGraph.AddUnsafePass(ProfilerTag, out PassData passData);
var cameraData = frameData.Get<UniversalCameraData>();
var resourceData = frameData.Get<UniversalResourceData>();
RenderTextureDescriptor rtDesc = cameraData.cameraTargetDescriptor;
rtDesc.depthBufferBits = 0;
rtDesc.msaaSamples = 1;
rtDesc.graphicsFormat = GraphicsFormat.R16G16B16A16_SFloat;
var textureHandle = UniversalRenderer.CreateRenderGraphTexture(renderGraph, rtDesc, GaussianSplatRTName, true);
passData.CameraData = cameraData;
passData.SourceTexture = resourceData.activeColorTexture;
passData.SourceDepth = resourceData.activeDepthTexture;
passData.GaussianSplatRT = textureHandle;
builder.UseTexture(resourceData.activeColorTexture, AccessFlags.ReadWrite);
builder.UseTexture(resourceData.activeDepthTexture);
builder.UseTexture(textureHandle, AccessFlags.Write);
builder.AllowPassCulling(false);
builder.SetRenderFunc(static (PassData data, UnsafeGraphContext context) =>
{
var commandBuffer = CommandBufferHelpers.GetNativeCommandBuffer(context.cmd);
using var _ = new ProfilingScope(commandBuffer, s_profilingSampler);
commandBuffer.SetGlobalTexture(s_gaussianSplatRT, data.GaussianSplatRT);
CoreUtils.SetRenderTarget(commandBuffer, data.GaussianSplatRT, data.SourceDepth, ClearFlag.Color, Color.clear);
Material matComposite = GaussianSplatRenderSystem.instance.SortAndRenderSplats(data.CameraData.camera, commandBuffer);
commandBuffer.BeginSample(GaussianSplatRenderSystem.s_ProfCompose);
Blitter.BlitCameraTexture(commandBuffer, data.GaussianSplatRT, data.SourceTexture, matComposite, 0);
commandBuffer.EndSample(GaussianSplatRenderSystem.s_ProfCompose);
});
}
#else
public override void OnCameraSetup(CommandBuffer cmd, ref RenderingData renderingData)
{
RenderTextureDescriptor rtDesc = renderingData.cameraData.cameraTargetDescriptor;
rtDesc.depthBufferBits = 0;
rtDesc.msaaSamples = 1;
rtDesc.graphicsFormat = GraphicsFormat.R16G16B16A16_SFloat;
RenderingUtils.ReAllocateIfNeeded(ref m_RenderTarget, rtDesc, FilterMode.Point, TextureWrapMode.Clamp, name: GaussianSplatRTName);
cmd.SetGlobalTexture(m_RenderTarget.name, m_RenderTarget.nameID);
ConfigureTarget(m_RenderTarget, m_Renderer.cameraDepthTargetHandle);
ConfigureClear(ClearFlag.Color, new Color(0,0,0,0));
}
public override void Execute(ScriptableRenderContext context, ref RenderingData renderingData)
{
if (m_Cmb == null)
return;
// add sorting, view calc and drawing commands for each splat object
Material matComposite = GaussianSplatRenderSystem.instance.SortAndRenderSplats(renderingData.cameraData.camera, m_Cmb);
// compose
m_Cmb.BeginSample(GaussianSplatRenderSystem.s_ProfCompose);
Blitter.BlitCameraTexture(m_Cmb, m_RenderTarget, m_Renderer.cameraColorTargetHandle, RenderBufferLoadAction.Load, RenderBufferStoreAction.Store, matComposite, 0);
m_Cmb.EndSample(GaussianSplatRenderSystem.s_ProfCompose);
context.ExecuteCommandBuffer(m_Cmb);
}
#endif
}
GSRenderPass m_Pass;
bool m_HasCamera;
public override void Create()
{
m_Pass = new GSRenderPass
{
renderPassEvent = RenderPassEvent.BeforeRenderingTransparents
};
}
public override void OnCameraPreCull(ScriptableRenderer renderer, in CameraData cameraData)
{
m_HasCamera = false;
var system = GaussianSplatRenderSystem.instance;
if (!system.GatherSplatsForCamera(cameraData.camera))
return;
#if !UNITY_6000_0_OR_NEWER
CommandBuffer cmb = system.InitialClearCmdBuffer(cameraData.camera);
m_Pass.m_Cmb = cmb;
#endif
m_HasCamera = true;
}
public override void AddRenderPasses(ScriptableRenderer renderer, ref RenderingData renderingData)
{
if (!m_HasCamera)
return;
#if !UNITY_6000_0_OR_NEWER
m_Pass.m_Renderer = renderer;
#endif
renderer.EnqueuePass(m_Pass);
}
protected override void Dispose(bool disposing)
{
m_Pass?.Dispose();
m_Pass = null;
}
}
}
#endif // #if GS_ENABLE_URP

View File

@@ -0,0 +1,17 @@
fileFormatVersion: 2
guid: 01052c1a4da12064d8681d7c1800f94a
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences:
- m_ShaderSplats: {fileID: 4800000, guid: ed800126ae8844a67aad1974ddddd59c, type: 3}
- m_ShaderComposite: {fileID: 4800000, guid: 7e184af7d01193a408eb916d8acafff9, type: 3}
- m_ShaderDebugPoints: {fileID: 4800000, guid: b44409fc67214394f8f47e4e2648425e, type: 3}
- m_ShaderDebugBoxes: {fileID: 4800000, guid: 4006f2680fd7c8b4cbcb881454c782be, type: 3}
- m_CSSplatUtilities: {fileID: 7200000, guid: ec84f78b836bd4f96a105d6b804f08bd, type: 3}
- m_CSFfxSort: {fileID: 7200000, guid: dec36776b6c843544a1f6f9b436a4474, type: 3}
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,31 @@
{
"name": "GaussianSplatting",
"rootNamespace": "GaussianSplatting.Runtime",
"references": [
"GUID:d8b63aba1907145bea998dd612889d6b",
"GUID:8992d429105beaf428dfc91fb5b9f531",
"GUID:15fc0a57446b3144c949da3e2b9737a9",
"GUID:df380645f10b7bc4b97d4f5eb6303d95",
"GUID:457756d89b35d2941b3e7b37b4ece6f1"
],
"includePlatforms": [],
"excludePlatforms": [],
"allowUnsafeCode": true,
"overrideReferences": false,
"precompiledReferences": [],
"autoReferenced": true,
"defineConstraints": [],
"versionDefines": [
{
"name": "com.unity.render-pipelines.universal",
"expression": "",
"define": "GS_ENABLE_URP"
},
{
"name": "com.unity.render-pipelines.high-definition",
"expression": "",
"define": "GS_ENABLE_HDRP"
}
],
"noEngineReferences": false
}

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: 4b653174f8fcdcd49b4c9a6f1ca8c7c3
AssemblyDefinitionImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,107 @@
// SPDX-License-Identifier: MIT
using Unity.Mathematics;
namespace GaussianSplatting.Runtime
{
public static class GaussianUtils
{
public static float Sigmoid(float v)
{
return math.rcp(1.0f + math.exp(-v));
}
public static float3 SH0ToColor(float3 dc0)
{
const float kSH_C0 = 0.2820948f;
return dc0 * kSH_C0 + 0.5f;
}
public static float3 LinearScale(float3 logScale)
{
return math.abs(math.exp(logScale));
}
public static float SquareCentered01(float x)
{
x -= 0.5f;
x *= x * math.sign(x);
return x * 2.0f + 0.5f;
}
public static float InvSquareCentered01(float x)
{
x -= 0.5f;
x *= 0.5f;
x = math.sqrt(math.abs(x)) * math.sign(x);
return x + 0.5f;
}
public static float4 NormalizeSwizzleRotation(float4 wxyz)
{
return math.normalize(wxyz).yzwx;
}
// Returns three smallest quaternion components in xyz (normalized to 0..1 range), and index/3 of the largest one in w
public static float4 PackSmallest3Rotation(float4 q)
{
// find biggest component
float4 absQ = math.abs(q);
int index = 0;
float maxV = absQ.x;
if (absQ.y > maxV)
{
index = 1;
maxV = absQ.y;
}
if (absQ.z > maxV)
{
index = 2;
maxV = absQ.z;
}
if (absQ.w > maxV)
{
index = 3;
maxV = absQ.w;
}
if (index == 0) q = q.yzwx;
if (index == 1) q = q.xzwy;
if (index == 2) q = q.xywz;
float3 three = q.xyz * (q.w >= 0 ? 1 : -1); // -1/sqrt2..+1/sqrt2 range
three = (three * math.SQRT2) * 0.5f + 0.5f; // 0..1 range
return new float4(three, index / 3.0f);
}
// Based on https://fgiesen.wordpress.com/2009/12/13/decoding-morton-codes/
// Insert two 0 bits after each of the 21 low bits of x
static ulong MortonPart1By2(ulong x)
{
x &= 0x1fffff;
x = (x ^ (x << 32)) & 0x1f00000000ffffUL;
x = (x ^ (x << 16)) & 0x1f0000ff0000ffUL;
x = (x ^ (x << 8)) & 0x100f00f00f00f00fUL;
x = (x ^ (x << 4)) & 0x10c30c30c30c30c3UL;
x = (x ^ (x << 2)) & 0x1249249249249249UL;
return x;
}
// Encode three 21-bit integers into 3D Morton order
public static ulong MortonEncode3(uint3 v)
{
return (MortonPart1By2(v.z) << 2) | (MortonPart1By2(v.y) << 1) | MortonPart1By2(v.x);
}
// See GaussianSplatting.hlsl
public static uint2 DecodeMorton2D_16x16(uint t)
{
t = (t & 0xFF) | ((t & 0xFE) << 7); // -EAFBGCHEAFBGCHD
t &= 0x5555; // -E-F-G-H-A-B-C-D
t = (t ^ (t >> 1)) & 0x3333; // --EF--GH--AB--CD
t = (t ^ (t >> 2)) & 0x0f0f; // ----EFGH----ABCD
return new uint2(t & 0xF, t >> 8); // --------EFGHABCD
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: ff862528cafe3e243aa42978a6d284d8
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,200 @@
using UnityEngine;
using UnityEngine.Assertions;
using UnityEngine.Rendering;
namespace GaussianSplatting.Runtime
{
// GPU (uint key, uint payload) 8 bit-LSD radix sort, using reduce-then-scan
// Copyright Thomas Smith 2024, MIT license
// https://github.com/b0nes164/GPUSorting
public class GpuSorting
{
//The size of a threadblock partition in the sort
const uint DEVICE_RADIX_SORT_PARTITION_SIZE = 3840;
//The size of our radix in bits
const uint DEVICE_RADIX_SORT_BITS = 8;
//Number of digits in our radix, 1 << DEVICE_RADIX_SORT_BITS
const uint DEVICE_RADIX_SORT_RADIX = 256;
//Number of sorting passes required to sort a 32bit key, KEY_BITS / DEVICE_RADIX_SORT_BITS
const uint DEVICE_RADIX_SORT_PASSES = 4;
//Keywords to enable for the shader
private LocalKeyword m_keyUintKeyword;
private LocalKeyword m_payloadUintKeyword;
private LocalKeyword m_ascendKeyword;
private LocalKeyword m_sortPairKeyword;
private LocalKeyword m_vulkanKeyword;
public struct Args
{
public uint count;
public GraphicsBuffer inputKeys;
public GraphicsBuffer inputValues;
public SupportResources resources;
internal int workGroupCount;
}
public struct SupportResources
{
public GraphicsBuffer altBuffer;
public GraphicsBuffer altPayloadBuffer;
public GraphicsBuffer passHistBuffer;
public GraphicsBuffer globalHistBuffer;
public static SupportResources Load(uint count)
{
//This is threadBlocks * DEVICE_RADIX_SORT_RADIX
uint scratchBufferSize = DivRoundUp(count, DEVICE_RADIX_SORT_PARTITION_SIZE) * DEVICE_RADIX_SORT_RADIX;
uint reducedScratchBufferSize = DEVICE_RADIX_SORT_RADIX * DEVICE_RADIX_SORT_PASSES;
var target = GraphicsBuffer.Target.Structured;
var resources = new SupportResources
{
altBuffer = new GraphicsBuffer(target, (int)count, 4) { name = "DeviceRadixAlt" },
altPayloadBuffer = new GraphicsBuffer(target, (int)count, 4) { name = "DeviceRadixAltPayload" },
passHistBuffer = new GraphicsBuffer(target, (int)scratchBufferSize, 4) { name = "DeviceRadixPassHistogram" },
globalHistBuffer = new GraphicsBuffer(target, (int)reducedScratchBufferSize, 4) { name = "DeviceRadixGlobalHistogram" },
};
return resources;
}
public void Dispose()
{
altBuffer?.Dispose();
altPayloadBuffer?.Dispose();
passHistBuffer?.Dispose();
globalHistBuffer?.Dispose();
altBuffer = null;
altPayloadBuffer = null;
passHistBuffer = null;
globalHistBuffer = null;
}
}
readonly ComputeShader m_CS;
readonly int m_kernelInitDeviceRadixSort = -1;
readonly int m_kernelUpsweep = -1;
readonly int m_kernelScan = -1;
readonly int m_kernelDownsweep = -1;
readonly bool m_Valid;
public bool Valid => m_Valid;
public GpuSorting(ComputeShader cs)
{
m_CS = cs;
if (cs)
{
m_kernelInitDeviceRadixSort = cs.FindKernel("InitDeviceRadixSort");
m_kernelUpsweep = cs.FindKernel("Upsweep");
m_kernelScan = cs.FindKernel("Scan");
m_kernelDownsweep = cs.FindKernel("Downsweep");
}
m_Valid = m_kernelInitDeviceRadixSort >= 0 &&
m_kernelUpsweep >= 0 &&
m_kernelScan >= 0 &&
m_kernelDownsweep >= 0;
if (m_Valid)
{
if (!cs.IsSupported(m_kernelInitDeviceRadixSort) ||
!cs.IsSupported(m_kernelUpsweep) ||
!cs.IsSupported(m_kernelScan) ||
!cs.IsSupported(m_kernelDownsweep))
{
m_Valid = false;
}
}
m_keyUintKeyword = new LocalKeyword(cs, "KEY_UINT");
m_payloadUintKeyword = new LocalKeyword(cs, "PAYLOAD_UINT");
m_ascendKeyword = new LocalKeyword(cs, "SHOULD_ASCEND");
m_sortPairKeyword = new LocalKeyword(cs, "SORT_PAIRS");
m_vulkanKeyword = new LocalKeyword(cs, "VULKAN");
cs.EnableKeyword(m_keyUintKeyword);
cs.EnableKeyword(m_payloadUintKeyword);
cs.EnableKeyword(m_ascendKeyword);
cs.EnableKeyword(m_sortPairKeyword);
if (SystemInfo.graphicsDeviceType == UnityEngine.Rendering.GraphicsDeviceType.Vulkan)
cs.EnableKeyword(m_vulkanKeyword);
else
cs.DisableKeyword(m_vulkanKeyword);
}
static uint DivRoundUp(uint x, uint y) => (x + y - 1) / y;
//Can we remove the last 4 padding without breaking?
struct SortConstants
{
public uint numKeys; // The number of keys to sort
public uint radixShift; // The radix shift value for the current pass
public uint threadBlocks; // threadBlocks
public uint padding0; // Padding - unused
}
public void Dispatch(CommandBuffer cmd, Args args)
{
Assert.IsTrue(Valid);
GraphicsBuffer srcKeyBuffer = args.inputKeys;
GraphicsBuffer srcPayloadBuffer = args.inputValues;
GraphicsBuffer dstKeyBuffer = args.resources.altBuffer;
GraphicsBuffer dstPayloadBuffer = args.resources.altPayloadBuffer;
SortConstants constants = default;
constants.numKeys = args.count;
constants.threadBlocks = DivRoundUp(args.count, DEVICE_RADIX_SORT_PARTITION_SIZE);
// Setup overall constants
cmd.SetComputeIntParam(m_CS, "e_numKeys", (int)constants.numKeys);
cmd.SetComputeIntParam(m_CS, "e_threadBlocks", (int)constants.threadBlocks);
//Set statically located buffers
//Upsweep
cmd.SetComputeBufferParam(m_CS, m_kernelUpsweep, "b_passHist", args.resources.passHistBuffer);
cmd.SetComputeBufferParam(m_CS, m_kernelUpsweep, "b_globalHist", args.resources.globalHistBuffer);
//Scan
cmd.SetComputeBufferParam(m_CS, m_kernelScan, "b_passHist", args.resources.passHistBuffer);
//Downsweep
cmd.SetComputeBufferParam(m_CS, m_kernelDownsweep, "b_passHist", args.resources.passHistBuffer);
cmd.SetComputeBufferParam(m_CS, m_kernelDownsweep, "b_globalHist", args.resources.globalHistBuffer);
//Clear the global histogram
cmd.SetComputeBufferParam(m_CS, m_kernelInitDeviceRadixSort, "b_globalHist", args.resources.globalHistBuffer);
cmd.DispatchCompute(m_CS, m_kernelInitDeviceRadixSort, 1, 1, 1);
// Execute the sort algorithm in 8-bit increments
for (constants.radixShift = 0; constants.radixShift < 32; constants.radixShift += DEVICE_RADIX_SORT_BITS)
{
cmd.SetComputeIntParam(m_CS, "e_radixShift", (int)constants.radixShift);
//Upsweep
cmd.SetComputeBufferParam(m_CS, m_kernelUpsweep, "b_sort", srcKeyBuffer);
cmd.DispatchCompute(m_CS, m_kernelUpsweep, (int)constants.threadBlocks, 1, 1);
// Scan
cmd.DispatchCompute(m_CS, m_kernelScan, (int)DEVICE_RADIX_SORT_RADIX, 1, 1);
// Downsweep
cmd.SetComputeBufferParam(m_CS, m_kernelDownsweep, "b_sort", srcKeyBuffer);
cmd.SetComputeBufferParam(m_CS, m_kernelDownsweep, "b_sortPayload", srcPayloadBuffer);
cmd.SetComputeBufferParam(m_CS, m_kernelDownsweep, "b_alt", dstKeyBuffer);
cmd.SetComputeBufferParam(m_CS, m_kernelDownsweep, "b_altPayload", dstPayloadBuffer);
cmd.DispatchCompute(m_CS, m_kernelDownsweep, (int)constants.threadBlocks, 1, 1);
// Swap
(srcKeyBuffer, dstKeyBuffer) = (dstKeyBuffer, srcKeyBuffer);
(srcPayloadBuffer, dstPayloadBuffer) = (dstPayloadBuffer, srcPayloadBuffer);
}
}
}
}

View File

@@ -0,0 +1,11 @@
fileFormatVersion: 2
guid: 65a55f12dc9f42e4196260841dd87c15
MonoImporter:
externalObjects: {}
serializedVersion: 2
defaultReferences: []
executionOrder: 0
icon: {instanceID: 0}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,8 @@
fileFormatVersion: 2
guid: 2cbb533de67b91d45afad2ab53f7f03c
folderAsset: yes
DefaultImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,44 @@
// SPDX-License-Identifier: MIT
Shader "Unlit/BlackSkybox"
{
Properties
{
_Color ("Color", Color) = (0,0,0,0)
}
SubShader
{
Pass
{
CGPROGRAM
#pragma vertex vert
#pragma fragment frag
#include "UnityCG.cginc"
struct appdata
{
float4 vertex : POSITION;
};
struct v2f
{
float4 vertex : SV_POSITION;
};
v2f vert (appdata v)
{
v2f o;
o.vertex = UnityObjectToClipPos(v.vertex);
return o;
}
half4 _Color;
half4 frag (v2f i) : SV_Target
{
return _Color;
}
ENDCG
}
}
}

View File

@@ -0,0 +1,9 @@
fileFormatVersion: 2
guid: a4867e5be68354ccda78062a92c74391
ShaderImporter:
externalObjects: {}
defaultTextures: []
nonModifiableTextures: []
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,531 @@
/******************************************************************************
* DeviceRadixSort
* Device Level 8-bit LSD Radix Sort using reduce then scan
*
* SPDX-License-Identifier: MIT
* Copyright Thomas Smith 5/17/2024
* https://github.com/b0nes164/GPUSorting
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
******************************************************************************/
#include "SortCommon.hlsl"
#define US_DIM 128U //The number of threads in a Upsweep threadblock
#define SCAN_DIM 128U //The number of threads in a Scan threadblock
RWStructuredBuffer<uint> b_globalHist; //buffer holding device level offsets for each binning pass
RWStructuredBuffer<uint> b_passHist; //buffer used to store reduced sums of partition tiles
groupshared uint g_us[RADIX * 2]; //Shared memory for upsweep
groupshared uint g_scan[SCAN_DIM]; //Shared memory for the scan
//*****************************************************************************
//INIT KERNEL
//*****************************************************************************
//Clear the global histogram, as we will be adding to it atomically
[numthreads(1024, 1, 1)]
void InitDeviceRadixSort(int3 id : SV_DispatchThreadID)
{
b_globalHist[id.x] = 0;
}
//*****************************************************************************
//UPSWEEP KERNEL
//*****************************************************************************
//histogram, 64 threads to a histogram
inline void HistogramDigitCounts(uint gtid, uint gid)
{
const uint histOffset = gtid / 64 * RADIX;
const uint partitionEnd = gid == e_threadBlocks - 1 ?
e_numKeys : (gid + 1) * PART_SIZE;
for (uint i = gtid + gid * PART_SIZE; i < partitionEnd; i += US_DIM)
{
#if defined(KEY_UINT)
InterlockedAdd(g_us[ExtractDigit(b_sort[i]) + histOffset], 1);
#elif defined(KEY_INT)
InterlockedAdd(g_us[ExtractDigit(IntToUint(b_sort[i])) + histOffset], 1);
#elif defined(KEY_FLOAT)
InterlockedAdd(g_us[ExtractDigit(FloatToUint(b_sort[i])) + histOffset], 1);
#endif
}
}
//reduce and pass to tile histogram
inline void ReduceWriteDigitCounts(uint gtid, uint gid)
{
for (uint i = gtid; i < RADIX; i += US_DIM)
{
g_us[i] += g_us[i + RADIX];
b_passHist[i * e_threadBlocks + gid] = g_us[i];
g_us[i] += WavePrefixSum(g_us[i]);
}
}
//Exclusive scan over digit counts, then atomically add to global hist
inline void GlobalHistExclusiveScanWGE16(uint gtid, uint waveSize)
{
GroupMemoryBarrierWithGroupSync();
if (gtid < (RADIX / waveSize))
{
g_us[(gtid + 1) * waveSize - 1] +=
WavePrefixSum(g_us[(gtid + 1) * waveSize - 1]);
}
GroupMemoryBarrierWithGroupSync();
//atomically add to global histogram
const uint globalHistOffset = GlobalHistOffset();
const uint laneMask = waveSize - 1;
const uint circularLaneShift = WaveGetLaneIndex() + 1 & laneMask;
for (uint i = gtid; i < RADIX; i += US_DIM)
{
const uint index = circularLaneShift + (i & ~laneMask);
uint t = WaveGetLaneIndex() != laneMask ? g_us[i] : 0;
if (i >= waveSize)
t += WaveReadLaneAt(g_us[i - 1], 0);
InterlockedAdd(b_globalHist[index + globalHistOffset], t);
}
}
inline void GlobalHistExclusiveScanWLT16(uint gtid, uint waveSize)
{
const uint globalHistOffset = GlobalHistOffset();
if (gtid < waveSize)
{
const uint circularLaneShift = WaveGetLaneIndex() + 1 &
waveSize - 1;
InterlockedAdd(b_globalHist[circularLaneShift + globalHistOffset],
circularLaneShift ? g_us[gtid] : 0);
}
GroupMemoryBarrierWithGroupSync();
const uint laneLog = countbits(waveSize - 1);
uint offset = laneLog;
uint j = waveSize;
for (; j < (RADIX >> 1); j <<= laneLog)
{
if (gtid < (RADIX >> offset))
{
g_us[((gtid + 1) << offset) - 1] +=
WavePrefixSum(g_us[((gtid + 1) << offset) - 1]);
}
GroupMemoryBarrierWithGroupSync();
for (uint i = gtid + j; i < RADIX; i += US_DIM)
{
if ((i & ((j << laneLog) - 1)) >= j)
{
if (i < (j << laneLog))
{
InterlockedAdd(b_globalHist[i + globalHistOffset],
WaveReadLaneAt(g_us[((i >> offset) << offset) - 1], 0) +
((i & (j - 1)) ? g_us[i - 1] : 0));
}
else
{
if ((i + 1) & (j - 1))
{
g_us[i] +=
WaveReadLaneAt(g_us[((i >> offset) << offset) - 1], 0);
}
}
}
}
offset += laneLog;
}
GroupMemoryBarrierWithGroupSync();
//If RADIX is not a power of lanecount
for (uint i = gtid + j; i < RADIX; i += US_DIM)
{
InterlockedAdd(b_globalHist[i + globalHistOffset],
WaveReadLaneAt(g_us[((i >> offset) << offset) - 1], 0) +
((i & (j - 1)) ? g_us[i - 1] : 0));
}
}
[numthreads(US_DIM, 1, 1)]
void Upsweep(uint3 gtid : SV_GroupThreadID, uint3 gid : SV_GroupID)
{
//get the wave size
const uint waveSize = getWaveSize();
//clear shared memory
const uint histsEnd = RADIX * 2;
for (uint i = gtid.x; i < histsEnd; i += US_DIM)
g_us[i] = 0;
GroupMemoryBarrierWithGroupSync();
HistogramDigitCounts(gtid.x, gid.x);
GroupMemoryBarrierWithGroupSync();
ReduceWriteDigitCounts(gtid.x, gid.x);
if (waveSize >= 16)
GlobalHistExclusiveScanWGE16(gtid.x, waveSize);
if (waveSize < 16)
GlobalHistExclusiveScanWLT16(gtid.x, waveSize);
}
//*****************************************************************************
//SCAN KERNEL
//*****************************************************************************
inline void ExclusiveThreadBlockScanFullWGE16(
uint gtid,
uint laneMask,
uint circularLaneShift,
uint partEnd,
uint deviceOffset,
uint waveSize,
inout uint reduction)
{
for (uint i = gtid; i < partEnd; i += SCAN_DIM)
{
g_scan[gtid] = b_passHist[i + deviceOffset];
g_scan[gtid] += WavePrefixSum(g_scan[gtid]);
GroupMemoryBarrierWithGroupSync();
if (gtid < SCAN_DIM / waveSize)
{
g_scan[(gtid + 1) * waveSize - 1] +=
WavePrefixSum(g_scan[(gtid + 1) * waveSize - 1]);
}
GroupMemoryBarrierWithGroupSync();
uint t = (WaveGetLaneIndex() != laneMask ? g_scan[gtid] : 0) + reduction;
if (gtid >= waveSize)
t += WaveReadLaneAt(g_scan[gtid - 1], 0);
b_passHist[circularLaneShift + (i & ~laneMask) + deviceOffset] = t;
reduction += g_scan[SCAN_DIM - 1];
GroupMemoryBarrierWithGroupSync();
}
}
inline void ExclusiveThreadBlockScanPartialWGE16(
uint gtid,
uint laneMask,
uint circularLaneShift,
uint partEnd,
uint deviceOffset,
uint waveSize,
uint reduction)
{
uint i = gtid + partEnd;
if (i < e_threadBlocks)
g_scan[gtid] = b_passHist[deviceOffset + i];
g_scan[gtid] += WavePrefixSum(g_scan[gtid]);
GroupMemoryBarrierWithGroupSync();
if (gtid < SCAN_DIM / waveSize)
{
g_scan[(gtid + 1) * waveSize - 1] +=
WavePrefixSum(g_scan[(gtid + 1) * waveSize - 1]);
}
GroupMemoryBarrierWithGroupSync();
const uint index = circularLaneShift + (i & ~laneMask);
if (index < e_threadBlocks)
{
uint t = (WaveGetLaneIndex() != laneMask ? g_scan[gtid] : 0) + reduction;
if (gtid >= waveSize)
t += g_scan[(gtid & ~laneMask) - 1];
b_passHist[index + deviceOffset] = t;
}
}
inline void ExclusiveThreadBlockScanWGE16(uint gtid, uint gid, uint waveSize)
{
uint reduction = 0;
const uint laneMask = waveSize - 1;
const uint circularLaneShift = WaveGetLaneIndex() + 1 & laneMask;
const uint partionsEnd = e_threadBlocks / SCAN_DIM * SCAN_DIM;
const uint deviceOffset = gid * e_threadBlocks;
ExclusiveThreadBlockScanFullWGE16(
gtid,
laneMask,
circularLaneShift,
partionsEnd,
deviceOffset,
waveSize,
reduction);
ExclusiveThreadBlockScanPartialWGE16(
gtid,
laneMask,
circularLaneShift,
partionsEnd,
deviceOffset,
waveSize,
reduction);
}
inline void ExclusiveThreadBlockScanFullWLT16(
uint gtid,
uint partitions,
uint deviceOffset,
uint laneLog,
uint circularLaneShift,
uint waveSize,
inout uint reduction)
{
for (uint k = 0; k < partitions; ++k)
{
g_scan[gtid] = b_passHist[gtid + k * SCAN_DIM + deviceOffset];
g_scan[gtid] += WavePrefixSum(g_scan[gtid]);
GroupMemoryBarrierWithGroupSync();
if (gtid < waveSize)
{
b_passHist[circularLaneShift + k * SCAN_DIM + deviceOffset] =
(circularLaneShift ? g_scan[gtid] : 0) + reduction;
}
uint offset = laneLog;
uint j = waveSize;
for (; j < (SCAN_DIM >> 1); j <<= laneLog)
{
if (gtid < (SCAN_DIM >> offset))
{
g_scan[((gtid + 1) << offset) - 1] +=
WavePrefixSum(g_scan[((gtid + 1) << offset) - 1]);
}
GroupMemoryBarrierWithGroupSync();
if ((gtid & ((j << laneLog) - 1)) >= j)
{
if (gtid < (j << laneLog))
{
b_passHist[gtid + k * SCAN_DIM + deviceOffset] =
WaveReadLaneAt(g_scan[((gtid >> offset) << offset) - 1], 0) +
((gtid & (j - 1)) ? g_scan[gtid - 1] : 0) + reduction;
}
else
{
if ((gtid + 1) & (j - 1))
{
g_scan[gtid] +=
WaveReadLaneAt(g_scan[((gtid >> offset) << offset) - 1], 0);
}
}
}
offset += laneLog;
}
GroupMemoryBarrierWithGroupSync();
//If SCAN_DIM is not a power of lanecount
for (uint i = gtid + j; i < SCAN_DIM; i += SCAN_DIM)
{
b_passHist[i + k * SCAN_DIM + deviceOffset] =
WaveReadLaneAt(g_scan[((i >> offset) << offset) - 1], 0) +
((i & (j - 1)) ? g_scan[i - 1] : 0) + reduction;
}
reduction += WaveReadLaneAt(g_scan[SCAN_DIM - 1], 0) +
WaveReadLaneAt(g_scan[(((SCAN_DIM - 1) >> offset) << offset) - 1], 0);
GroupMemoryBarrierWithGroupSync();
}
}
inline void ExclusiveThreadBlockScanParitalWLT16(
uint gtid,
uint partitions,
uint deviceOffset,
uint laneLog,
uint circularLaneShift,
uint waveSize,
uint reduction)
{
const uint finalPartSize = e_threadBlocks - partitions * SCAN_DIM;
if (gtid < finalPartSize)
{
g_scan[gtid] = b_passHist[gtid + partitions * SCAN_DIM + deviceOffset];
g_scan[gtid] += WavePrefixSum(g_scan[gtid]);
}
GroupMemoryBarrierWithGroupSync();
if (gtid < waveSize && circularLaneShift < finalPartSize)
{
b_passHist[circularLaneShift + partitions * SCAN_DIM + deviceOffset] =
(circularLaneShift ? g_scan[gtid] : 0) + reduction;
}
uint offset = laneLog;
for (uint j = waveSize; j < finalPartSize; j <<= laneLog)
{
if (gtid < (finalPartSize >> offset))
{
g_scan[((gtid + 1) << offset) - 1] +=
WavePrefixSum(g_scan[((gtid + 1) << offset) - 1]);
}
GroupMemoryBarrierWithGroupSync();
if ((gtid & ((j << laneLog) - 1)) >= j && gtid < finalPartSize)
{
if (gtid < (j << laneLog))
{
b_passHist[gtid + partitions * SCAN_DIM + deviceOffset] =
WaveReadLaneAt(g_scan[((gtid >> offset) << offset) - 1], 0) +
((gtid & (j - 1)) ? g_scan[gtid - 1] : 0) + reduction;
}
else
{
if ((gtid + 1) & (j - 1))
{
g_scan[gtid] +=
WaveReadLaneAt(g_scan[((gtid >> offset) << offset) - 1], 0);
}
}
}
offset += laneLog;
}
}
inline void ExclusiveThreadBlockScanWLT16(uint gtid, uint gid, uint waveSize)
{
uint reduction = 0;
const uint partitions = e_threadBlocks / SCAN_DIM;
const uint deviceOffset = gid * e_threadBlocks;
const uint laneLog = countbits(waveSize - 1);
const uint circularLaneShift = WaveGetLaneIndex() + 1 & waveSize - 1;
ExclusiveThreadBlockScanFullWLT16(
gtid,
partitions,
deviceOffset,
laneLog,
circularLaneShift,
waveSize,
reduction);
ExclusiveThreadBlockScanParitalWLT16(
gtid,
partitions,
deviceOffset,
laneLog,
circularLaneShift,
waveSize,
reduction);
}
//Scan does not need flattening of gids
[numthreads(SCAN_DIM, 1, 1)]
void Scan(uint3 gtid : SV_GroupThreadID, uint3 gid : SV_GroupID)
{
const uint waveSize = getWaveSize();
if (waveSize >= 16)
ExclusiveThreadBlockScanWGE16(gtid.x, gid.x, waveSize);
if (waveSize < 16)
ExclusiveThreadBlockScanWLT16(gtid.x, gid.x, waveSize);
}
//*****************************************************************************
//DOWNSWEEP KERNEL
//*****************************************************************************
inline void LoadThreadBlockReductions(uint gtid, uint gid, uint exclusiveHistReduction)
{
if (gtid < RADIX)
{
g_d[gtid + PART_SIZE] = b_globalHist[gtid + GlobalHistOffset()] +
b_passHist[gtid * e_threadBlocks + gid] - exclusiveHistReduction;
}
}
[numthreads(D_DIM, 1, 1)]
void Downsweep(uint3 gtid : SV_GroupThreadID, uint3 gid : SV_GroupID)
{
KeyStruct keys;
OffsetStruct offsets;
const uint waveSize = getWaveSize();
ClearWaveHists(gtid.x, waveSize);
GroupMemoryBarrierWithGroupSync();
if (gid.x < e_threadBlocks - 1)
{
if (waveSize >= 16)
keys = LoadKeysWGE16(gtid.x, waveSize, gid.x);
if (waveSize < 16)
keys = LoadKeysWLT16(gtid.x, waveSize, gid.x, SerialIterations(waveSize));
}
if (gid.x == e_threadBlocks - 1)
{
if (waveSize >= 16)
keys = LoadKeysPartialWGE16(gtid.x, waveSize, gid.x);
if (waveSize < 16)
keys = LoadKeysPartialWLT16(gtid.x, waveSize, gid.x, SerialIterations(waveSize));
}
uint exclusiveHistReduction;
if (waveSize >= 16)
{
offsets = RankKeysWGE16(waveSize, getWaveIndex(gtid.x, waveSize) * RADIX, keys);
GroupMemoryBarrierWithGroupSync();
uint histReduction;
if (gtid.x < RADIX)
{
histReduction = WaveHistInclusiveScanCircularShiftWGE16(gtid.x, waveSize);
histReduction += WavePrefixSum(histReduction); //take advantage of barrier to begin scan
}
GroupMemoryBarrierWithGroupSync();
WaveHistReductionExclusiveScanWGE16(gtid.x, waveSize, histReduction);
GroupMemoryBarrierWithGroupSync();
UpdateOffsetsWGE16(gtid.x, waveSize, offsets, keys);
if (gtid.x < RADIX)
exclusiveHistReduction = g_d[gtid.x]; //take advantage of barrier to grab value
GroupMemoryBarrierWithGroupSync();
}
if (waveSize < 16)
{
offsets = RankKeysWLT16(waveSize, getWaveIndex(gtid.x, waveSize), keys, SerialIterations(waveSize));
if (gtid.x < HALF_RADIX)
{
uint histReduction = WaveHistInclusiveScanCircularShiftWLT16(gtid.x);
g_d[gtid.x] = histReduction + (histReduction << 16); //take advantage of barrier to begin scan
}
WaveHistReductionExclusiveScanWLT16(gtid.x);
GroupMemoryBarrierWithGroupSync();
UpdateOffsetsWLT16(gtid.x, waveSize, SerialIterations(waveSize), offsets, keys);
if (gtid.x < RADIX) //take advantage of barrier to grab value
exclusiveHistReduction = g_d[gtid.x >> 1] >> ((gtid.x & 1) ? 16 : 0) & 0xffff;
GroupMemoryBarrierWithGroupSync();
}
ScatterKeysShared(offsets, keys);
LoadThreadBlockReductions(gtid.x, gid.x, exclusiveHistReduction);
GroupMemoryBarrierWithGroupSync();
if (gid.x < e_threadBlocks - 1)
ScatterDevice(gtid.x, waveSize, gid.x, offsets);
if (gid.x == e_threadBlocks - 1)
ScatterDevicePartial(gtid.x, waveSize, gid.x, offsets);
}

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: 02209b8d952e7fc418492b88139826fd
ShaderIncludeImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,45 @@
// SPDX-License-Identifier: MIT
Shader "Hidden/Gaussian Splatting/Composite"
{
SubShader
{
Pass
{
ZWrite Off
ZTest Always
Cull Off
Blend SrcAlpha OneMinusSrcAlpha
CGPROGRAM
#pragma vertex vert
#pragma fragment frag
#pragma require compute
#pragma use_dxc
#include "UnityCG.cginc"
struct v2f
{
float4 vertex : SV_POSITION;
};
v2f vert (uint vtxID : SV_VertexID)
{
v2f o;
float2 quadPos = float2(vtxID&1, (vtxID>>1)&1) * 4.0 - 1.0;
o.vertex = float4(quadPos, 1, 1);
return o;
}
Texture2D _GaussianSplatRT;
half4 frag (v2f i) : SV_Target
{
half4 col = _GaussianSplatRT.Load(int3(i.vertex.xy, 0));
col.rgb = GammaToLinearSpace(col.rgb);
col.a = saturate(col.a * 1.5);
return col;
}
ENDCG
}
}
}

View File

@@ -0,0 +1,9 @@
fileFormatVersion: 2
guid: 7e184af7d01193a408eb916d8acafff9
ShaderImporter:
externalObjects: {}
defaultTextures: []
nonModifiableTextures: []
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,102 @@
// SPDX-License-Identifier: MIT
Shader "Gaussian Splatting/Debug/Render Boxes"
{
SubShader
{
Tags { "RenderType"="Transparent" "Queue"="Transparent" }
Pass
{
ZWrite Off
Blend OneMinusDstAlpha One
Cull Front
CGPROGRAM
#pragma vertex vert
#pragma fragment frag
#pragma require compute
#pragma use_dxc
#include "UnityCG.cginc"
#include "GaussianSplatting.hlsl"
StructuredBuffer<uint> _OrderBuffer;
bool _DisplayChunks;
struct v2f
{
half4 col : COLOR0;
float4 vertex : SV_POSITION;
};
float _SplatScale;
float _SplatOpacityScale;
// based on https://iquilezles.org/articles/palettes/
// cosine based palette, 4 vec3 params
half3 palette(float t, half3 a, half3 b, half3 c, half3 d)
{
return a + b*cos(6.28318*(c*t+d));
}
v2f vert (uint vtxID : SV_VertexID, uint instID : SV_InstanceID)
{
v2f o;
bool chunks = _DisplayChunks;
uint idx = vtxID;
float3 localPos = float3(idx&1, (idx>>1)&1, (idx>>2)&1) * 2.0 - 1.0;
float3 centerWorldPos = 0;
if (!chunks)
{
// display splat boxes
instID = _OrderBuffer[instID];
SplatData splat = LoadSplatData(instID);
float4 boxRot = splat.rot;
float3 boxSize = splat.scale;
boxSize *= _SplatScale;
float3x3 splatRotScaleMat = CalcMatrixFromRotationScale(boxRot, boxSize);
splatRotScaleMat = mul((float3x3)unity_ObjectToWorld, splatRotScaleMat);
centerWorldPos = splat.pos;
centerWorldPos = mul(unity_ObjectToWorld, float4(centerWorldPos,1)).xyz;
o.col.rgb = saturate(splat.sh.col);
o.col.a = saturate(splat.opacity * _SplatOpacityScale);
localPos = mul(splatRotScaleMat, localPos) * 2;
}
else
{
// display chunk boxes
localPos = localPos * 0.5 + 0.5;
SplatChunkInfo chunk = _SplatChunks[instID];
float3 posMin = float3(chunk.posX.x, chunk.posY.x, chunk.posZ.x);
float3 posMax = float3(chunk.posX.y, chunk.posY.y, chunk.posZ.y);
localPos = lerp(posMin, posMax, localPos);
localPos = mul(unity_ObjectToWorld, float4(localPos,1)).xyz;
o.col.rgb = palette((float)instID / (float)_SplatChunkCount, half3(0.5,0.5,0.5), half3(0.5,0.5,0.5), half3(1,1,1), half3(0.0, 0.33, 0.67));
o.col.a = 0.1;
}
float3 worldPos = centerWorldPos + localPos;
o.vertex = UnityWorldToClipPos(worldPos);
return o;
}
half4 frag (v2f i) : SV_Target
{
half4 res = half4(i.col.rgb * i.col.a, i.col.a);
return res;
}
ENDCG
}
}
}

View File

@@ -0,0 +1,9 @@
fileFormatVersion: 2
guid: 4006f2680fd7c8b4cbcb881454c782be
ShaderImporter:
externalObjects: {}
defaultTextures: []
nonModifiableTextures: []
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,65 @@
// SPDX-License-Identifier: MIT
Shader "Gaussian Splatting/Debug/Render Points"
{
SubShader
{
Tags { "RenderType"="Transparent" "Queue"="Transparent" }
Pass
{
ZWrite On
Cull Off
CGPROGRAM
#pragma vertex vert
#pragma fragment frag
#pragma require compute
#pragma use_dxc
#include "GaussianSplatting.hlsl"
struct v2f
{
half3 color : TEXCOORD0;
float4 vertex : SV_POSITION;
};
float _SplatSize;
bool _DisplayIndex;
int _SplatCount;
v2f vert (uint vtxID : SV_VertexID, uint instID : SV_InstanceID)
{
v2f o;
uint splatIndex = instID;
SplatData splat = LoadSplatData(splatIndex);
float3 centerWorldPos = splat.pos;
centerWorldPos = mul(unity_ObjectToWorld, float4(centerWorldPos,1)).xyz;
float4 centerClipPos = mul(UNITY_MATRIX_VP, float4(centerWorldPos, 1));
o.vertex = centerClipPos;
uint idx = vtxID;
float2 quadPos = float2(idx&1, (idx>>1)&1) * 2.0 - 1.0;
o.vertex.xy += (quadPos * _SplatSize / _ScreenParams.xy) * o.vertex.w;
o.color.rgb = saturate(splat.sh.col);
if (_DisplayIndex)
{
o.color.r = frac((float)splatIndex / (float)_SplatCount * 100);
o.color.g = frac((float)splatIndex / (float)_SplatCount * 10);
o.color.b = (float)splatIndex / (float)_SplatCount;
}
return o;
}
half4 frag (v2f i) : SV_Target
{
return half4(i.color.rgb, 1);
}
ENDCG
}
}
}

View File

@@ -0,0 +1,9 @@
fileFormatVersion: 2
guid: b44409fc67214394f8f47e4e2648425e
ShaderImporter:
externalObjects: {}
defaultTextures: []
nonModifiableTextures: []
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,617 @@
// SPDX-License-Identifier: MIT
#ifndef GAUSSIAN_SPLATTING_HLSL
#define GAUSSIAN_SPLATTING_HLSL
float InvSquareCentered01(float x)
{
x -= 0.5;
x *= 0.5;
x = sqrt(abs(x)) * sign(x);
return x + 0.5;
}
float3 QuatRotateVector(float3 v, float4 r)
{
float3 t = 2 * cross(r.xyz, v);
return v + r.w * t + cross(r.xyz, t);
}
float4 QuatMul(float4 a, float4 b)
{
return float4(a.wwww * b + (a.xyzx * b.wwwx + a.yzxy * b.zxyy) * float4(1,1,1,-1) - a.zxyz * b.yzxz);
}
float4 QuatInverse(float4 q)
{
return rcp(dot(q, q)) * q * float4(-1,-1,-1,1);
}
float3x3 CalcMatrixFromRotationScale(float4 rot, float3 scale)
{
float3x3 ms = float3x3(
scale.x, 0, 0,
0, scale.y, 0,
0, 0, scale.z
);
float x = rot.x;
float y = rot.y;
float z = rot.z;
float w = rot.w;
float3x3 mr = float3x3(
1-2*(y*y + z*z), 2*(x*y - w*z), 2*(x*z + w*y),
2*(x*y + w*z), 1-2*(x*x + z*z), 2*(y*z - w*x),
2*(x*z - w*y), 2*(y*z + w*x), 1-2*(x*x + y*y)
);
return mul(mr, ms);
}
void CalcCovariance3D(float3x3 rotMat, out float3 sigma0, out float3 sigma1)
{
float3x3 sig = mul(rotMat, transpose(rotMat));
sigma0 = float3(sig._m00, sig._m01, sig._m02);
sigma1 = float3(sig._m11, sig._m12, sig._m22);
}
// from "EWA Splatting" (Zwicker et al 2002) eq. 31
float3 CalcCovariance2D(float3 worldPos, float3 cov3d0, float3 cov3d1, float4x4 matrixV, float4x4 matrixP, float4 screenParams)
{
float4x4 viewMatrix = matrixV;
float3 viewPos = mul(viewMatrix, float4(worldPos, 1)).xyz;
// this is needed in order for splats that are visible in view but clipped "quite a lot" to work
float aspect = matrixP._m00 / matrixP._m11;
float tanFovX = rcp(matrixP._m00);
float tanFovY = rcp(matrixP._m11 * aspect);
float limX = 1.3 * tanFovX;
float limY = 1.3 * tanFovY;
viewPos.x = clamp(viewPos.x / viewPos.z, -limX, limX) * viewPos.z;
viewPos.y = clamp(viewPos.y / viewPos.z, -limY, limY) * viewPos.z;
float focal = screenParams.x * matrixP._m00 / 2;
float3x3 J = float3x3(
focal / viewPos.z, 0, -(focal * viewPos.x) / (viewPos.z * viewPos.z),
0, focal / viewPos.z, -(focal * viewPos.y) / (viewPos.z * viewPos.z),
0, 0, 0
);
float3x3 W = (float3x3)viewMatrix;
float3x3 T = mul(J, W);
float3x3 V = float3x3(
cov3d0.x, cov3d0.y, cov3d0.z,
cov3d0.y, cov3d1.x, cov3d1.y,
cov3d0.z, cov3d1.y, cov3d1.z
);
float3x3 cov = mul(T, mul(V, transpose(T)));
// Low pass filter to make each splat at least 1px size.
cov._m00 += 0.3;
cov._m11 += 0.3;
return float3(cov._m00, cov._m01, cov._m11);
}
float3 CalcConic(float3 cov2d)
{
float det = cov2d.x * cov2d.z - cov2d.y * cov2d.y;
return float3(cov2d.z, -cov2d.y, cov2d.x) * rcp(det);
}
float2 CalcScreenSpaceDelta(float2 svPositionXY, float2 centerXY, float4 projectionParams)
{
float2 d = svPositionXY - centerXY;
d.y *= projectionParams.x;
return d;
}
float CalcPowerFromConic(float3 conic, float2 d)
{
return -0.5 * (conic.x * d.x*d.x + conic.z * d.y*d.y) + conic.y * d.x*d.y;
}
// Morton interleaving 16x16 group i.e. by 4 bits of coordinates, based on this thread:
// https://twitter.com/rygorous/status/986715358852608000
// which is simplified version of https://fgiesen.wordpress.com/2009/12/13/decoding-morton-codes/
uint EncodeMorton2D_16x16(uint2 c)
{
uint t = ((c.y & 0xF) << 8) | (c.x & 0xF); // ----EFGH----ABCD
t = (t ^ (t << 2)) & 0x3333; // --EF--GH--AB--CD
t = (t ^ (t << 1)) & 0x5555; // -E-F-G-H-A-B-C-D
return (t | (t >> 7)) & 0xFF; // --------EAFBGCHD
}
uint2 DecodeMorton2D_16x16(uint t) // --------EAFBGCHD
{
t = (t & 0xFF) | ((t & 0xFE) << 7); // -EAFBGCHEAFBGCHD
t &= 0x5555; // -E-F-G-H-A-B-C-D
t = (t ^ (t >> 1)) & 0x3333; // --EF--GH--AB--CD
t = (t ^ (t >> 2)) & 0x0f0f; // ----EFGH----ABCD
return uint2(t & 0xF, t >> 8); // --------EFGHABCD
}
static const float SH_C1 = 0.4886025;
static const float SH_C2[] = { 1.0925484, -1.0925484, 0.3153916, -1.0925484, 0.5462742 };
static const float SH_C3[] = { -0.5900436, 2.8906114, -0.4570458, 0.3731763, -0.4570458, 1.4453057, -0.5900436 };
struct SplatSHData
{
half3 col, sh1, sh2, sh3, sh4, sh5, sh6, sh7, sh8, sh9, sh10, sh11, sh12, sh13, sh14, sh15;
};
half3 ShadeSH(SplatSHData splat, half3 dir, int shOrder, bool onlySH)
{
dir *= -1;
half x = dir.x, y = dir.y, z = dir.z;
// ambient band
half3 res = splat.col; // col = sh0 * SH_C0 + 0.5 is already precomputed
if (onlySH)
res = 0.5;
// 1st degree
if (shOrder >= 1)
{
res += SH_C1 * (-splat.sh1 * y + splat.sh2 * z - splat.sh3 * x);
// 2nd degree
if (shOrder >= 2)
{
half xx = x * x, yy = y * y, zz = z * z;
half xy = x * y, yz = y * z, xz = x * z;
res +=
(SH_C2[0] * xy) * splat.sh4 +
(SH_C2[1] * yz) * splat.sh5 +
(SH_C2[2] * (2 * zz - xx - yy)) * splat.sh6 +
(SH_C2[3] * xz) * splat.sh7 +
(SH_C2[4] * (xx - yy)) * splat.sh8;
// 3rd degree
if (shOrder >= 3)
{
res +=
(SH_C3[0] * y * (3 * xx - yy)) * splat.sh9 +
(SH_C3[1] * xy * z) * splat.sh10 +
(SH_C3[2] * y * (4 * zz - xx - yy)) * splat.sh11 +
(SH_C3[3] * z * (2 * zz - 3 * xx - 3 * yy)) * splat.sh12 +
(SH_C3[4] * x * (4 * zz - xx - yy)) * splat.sh13 +
(SH_C3[5] * z * (xx - yy)) * splat.sh14 +
(SH_C3[6] * x * (xx - 3 * yy)) * splat.sh15;
}
}
}
return max(res, 0);
}
static const uint kTexWidth = 2048;
uint3 SplatIndexToPixelIndex(uint idx)
{
uint3 res;
uint2 xy = DecodeMorton2D_16x16(idx);
uint width = kTexWidth / 16;
idx >>= 8;
res.x = (idx % width) * 16 + xy.x;
res.y = (idx / width) * 16 + xy.y;
res.z = 0;
return res;
}
struct SplatChunkInfo
{
uint colR, colG, colB, colA;
float2 posX, posY, posZ;
uint sclX, sclY, sclZ;
uint shR, shG, shB;
};
StructuredBuffer<SplatChunkInfo> _SplatChunks;
uint _SplatChunkCount;
static const uint kChunkSize = 256;
struct SplatData
{
float3 pos;
float4 rot;
float3 scale;
half opacity;
SplatSHData sh;
};
// Decode quaternion from a "smallest 3" e.g. 10.10.10.2 format
float4 DecodeRotation(float4 pq)
{
uint idx = (uint)round(pq.w * 3.0); // note: need to round or index might come out wrong in some formats (e.g. fp16.fp16.fp16.fp16)
float4 q;
q.xyz = pq.xyz * sqrt(2.0) - (1.0 / sqrt(2.0));
q.w = sqrt(1.0 - saturate(dot(q.xyz, q.xyz)));
if (idx == 0) q = q.wxyz;
if (idx == 1) q = q.xwyz;
if (idx == 2) q = q.xywz;
return q;
}
float4 PackSmallest3Rotation(float4 q)
{
// find biggest component
float4 absQ = abs(q);
int index = 0;
float maxV = absQ.x;
if (absQ.y > maxV)
{
index = 1;
maxV = absQ.y;
}
if (absQ.z > maxV)
{
index = 2;
maxV = absQ.z;
}
if (absQ.w > maxV)
{
index = 3;
maxV = absQ.w;
}
if (index == 0) q = q.yzwx;
if (index == 1) q = q.xzwy;
if (index == 2) q = q.xywz;
float3 three = q.xyz * (q.w >= 0 ? 1 : -1); // -1/sqrt2..+1/sqrt2 range
three = (three * sqrt(2.0)) * 0.5 + 0.5; // 0..1 range
return float4(three, index / 3.0);
}
half3 DecodePacked_6_5_5(uint enc)
{
return half3(
(enc & 63) / 63.0,
((enc >> 6) & 31) / 31.0,
((enc >> 11) & 31) / 31.0);
}
half3 DecodePacked_5_6_5(uint enc)
{
return half3(
(enc & 31) / 31.0,
((enc >> 5) & 63) / 63.0,
((enc >> 11) & 31) / 31.0);
}
half3 DecodePacked_11_10_11(uint enc)
{
return half3(
(enc & 2047) / 2047.0,
((enc >> 11) & 1023) / 1023.0,
((enc >> 21) & 2047) / 2047.0);
}
float3 DecodePacked_16_16_16(uint2 enc)
{
return float3(
(enc.x & 65535) / 65535.0,
((enc.x >> 16) & 65535) / 65535.0,
(enc.y & 65535) / 65535.0);
}
float4 DecodePacked_10_10_10_2(uint enc)
{
return float4(
(enc & 1023) / 1023.0,
((enc >> 10) & 1023) / 1023.0,
((enc >> 20) & 1023) / 1023.0,
((enc >> 30) & 3) / 3.0);
}
uint EncodeQuatToNorm10(float4 v) // 32 bits: 10.10.10.2
{
return (uint) (v.x * 1023.5f) | ((uint) (v.y * 1023.5f) << 10) | ((uint) (v.z * 1023.5f) << 20) | ((uint) (v.w * 3.5f) << 30);
}
#ifdef SHADER_STAGE_COMPUTE
#define SplatBufferDataType RWByteAddressBuffer
#else
#define SplatBufferDataType ByteAddressBuffer
#endif
SplatBufferDataType _SplatPos;
SplatBufferDataType _SplatOther;
SplatBufferDataType _SplatSH;
Texture2D _SplatColor;
uint _SplatFormat;
// Match GaussianSplatAsset.VectorFormat
#define VECTOR_FMT_32F 0
#define VECTOR_FMT_16 1
#define VECTOR_FMT_11 2
#define VECTOR_FMT_6 3
uint LoadUShort(SplatBufferDataType dataBuffer, uint addrU)
{
uint addrA = addrU & ~0x3;
uint val = dataBuffer.Load(addrA);
if (addrU != addrA)
val >>= 16;
return val & 0xFFFF;
}
uint LoadUInt(SplatBufferDataType dataBuffer, uint addrU)
{
uint addrA = addrU & ~0x3;
uint val = dataBuffer.Load(addrA);
if (addrU != addrA)
{
uint val1 = dataBuffer.Load(addrA + 4);
val = (val >> 16) | ((val1 & 0xFFFF) << 16);
}
return val;
}
float3 LoadAndDecodeVector(SplatBufferDataType dataBuffer, uint addrU, uint fmt)
{
uint addrA = addrU & ~0x3;
uint val0 = dataBuffer.Load(addrA);
float3 res = 0;
if (fmt == VECTOR_FMT_32F)
{
uint val1 = dataBuffer.Load(addrA + 4);
uint val2 = dataBuffer.Load(addrA + 8);
if (addrU != addrA)
{
uint val3 = dataBuffer.Load(addrA + 12);
val0 = (val0 >> 16) | ((val1 & 0xFFFF) << 16);
val1 = (val1 >> 16) | ((val2 & 0xFFFF) << 16);
val2 = (val2 >> 16) | ((val3 & 0xFFFF) << 16);
}
res = float3(asfloat(val0), asfloat(val1), asfloat(val2));
}
else if (fmt == VECTOR_FMT_16)
{
uint val1 = dataBuffer.Load(addrA + 4);
if (addrU != addrA)
{
val0 = (val0 >> 16) | ((val1 & 0xFFFF) << 16);
val1 >>= 16;
}
res = DecodePacked_16_16_16(uint2(val0, val1));
}
else if (fmt == VECTOR_FMT_11)
{
uint val1 = dataBuffer.Load(addrA + 4);
if (addrU != addrA)
{
val0 = (val0 >> 16) | ((val1 & 0xFFFF) << 16);
}
res = DecodePacked_11_10_11(val0);
}
else if (fmt == VECTOR_FMT_6)
{
if (addrU != addrA)
val0 >>= 16;
res = DecodePacked_6_5_5(val0);
}
return res;
}
float3 LoadSplatPosValue(uint index)
{
uint fmt = _SplatFormat & 0xFF;
uint stride = 0;
if (fmt == VECTOR_FMT_32F)
stride = 12;
else if (fmt == VECTOR_FMT_16)
stride = 6;
else if (fmt == VECTOR_FMT_11)
stride = 4;
else if (fmt == VECTOR_FMT_6)
stride = 2;
return LoadAndDecodeVector(_SplatPos, index * stride, fmt);
}
float3 LoadSplatPos(uint idx)
{
float3 pos = LoadSplatPosValue(idx);
uint chunkIdx = idx / kChunkSize;
if (chunkIdx < _SplatChunkCount)
{
SplatChunkInfo chunk = _SplatChunks[chunkIdx];
float3 posMin = float3(chunk.posX.x, chunk.posY.x, chunk.posZ.x);
float3 posMax = float3(chunk.posX.y, chunk.posY.y, chunk.posZ.y);
pos = lerp(posMin, posMax, pos);
}
return pos;
}
half4 LoadSplatColTex(uint3 coord)
{
return _SplatColor.Load(coord);
}
SplatData LoadSplatData(uint idx)
{
SplatData s = (SplatData)0;
// figure out raw data offsets / locations
uint3 coord = SplatIndexToPixelIndex(idx);
uint scaleFmt = (_SplatFormat >> 8) & 0xFF;
uint shFormat = (_SplatFormat >> 16) & 0xFF;
uint otherStride = 4; // rotation is 10.10.10.2
if (scaleFmt == VECTOR_FMT_32F)
otherStride += 12;
else if (scaleFmt == VECTOR_FMT_16)
otherStride += 6;
else if (scaleFmt == VECTOR_FMT_11)
otherStride += 4;
else if (scaleFmt == VECTOR_FMT_6)
otherStride += 2;
if (shFormat > VECTOR_FMT_6)
otherStride += 2;
uint otherAddr = idx * otherStride;
uint shStride = 0;
if (shFormat == VECTOR_FMT_32F)
shStride = 192; // 15*3 fp32, rounded up to multiple of 16
else if (shFormat == VECTOR_FMT_16 || shFormat > VECTOR_FMT_6)
shStride = 96; // 15*3 fp16, rounded up to multiple of 16
else if (shFormat == VECTOR_FMT_11)
shStride = 60; // 15x uint
else if (shFormat == VECTOR_FMT_6)
shStride = 32; // 15x ushort, rounded up to multiple of 4
// load raw splat data, which might be chunk-relative
s.pos = LoadSplatPosValue(idx);
s.rot = DecodeRotation(DecodePacked_10_10_10_2(LoadUInt(_SplatOther, otherAddr)));
s.scale = LoadAndDecodeVector(_SplatOther, otherAddr + 4, scaleFmt);
half4 col = LoadSplatColTex(coord);
uint shIndex = idx;
if (shFormat > VECTOR_FMT_6)
shIndex = LoadUShort(_SplatOther, otherAddr + otherStride - 2);
uint shOffset = shIndex * shStride;
uint4 shRaw0 = _SplatSH.Load4(shOffset);
uint4 shRaw1 = _SplatSH.Load4(shOffset + 16);
if (shFormat == VECTOR_FMT_32F)
{
uint4 shRaw2 = _SplatSH.Load4(shOffset + 32);
uint4 shRaw3 = _SplatSH.Load4(shOffset + 48);
uint4 shRaw4 = _SplatSH.Load4(shOffset + 64);
uint4 shRaw5 = _SplatSH.Load4(shOffset + 80);
uint4 shRaw6 = _SplatSH.Load4(shOffset + 96);
uint4 shRaw7 = _SplatSH.Load4(shOffset + 112);
uint4 shRaw8 = _SplatSH.Load4(shOffset + 128);
uint4 shRaw9 = _SplatSH.Load4(shOffset + 144);
uint4 shRawA = _SplatSH.Load4(shOffset + 160);
uint shRawB = _SplatSH.Load(shOffset + 176);
s.sh.sh1.r = asfloat(shRaw0.x); s.sh.sh1.g = asfloat(shRaw0.y); s.sh.sh1.b = asfloat(shRaw0.z);
s.sh.sh2.r = asfloat(shRaw0.w); s.sh.sh2.g = asfloat(shRaw1.x); s.sh.sh2.b = asfloat(shRaw1.y);
s.sh.sh3.r = asfloat(shRaw1.z); s.sh.sh3.g = asfloat(shRaw1.w); s.sh.sh3.b = asfloat(shRaw2.x);
s.sh.sh4.r = asfloat(shRaw2.y); s.sh.sh4.g = asfloat(shRaw2.z); s.sh.sh4.b = asfloat(shRaw2.w);
s.sh.sh5.r = asfloat(shRaw3.x); s.sh.sh5.g = asfloat(shRaw3.y); s.sh.sh5.b = asfloat(shRaw3.z);
s.sh.sh6.r = asfloat(shRaw3.w); s.sh.sh6.g = asfloat(shRaw4.x); s.sh.sh6.b = asfloat(shRaw4.y);
s.sh.sh7.r = asfloat(shRaw4.z); s.sh.sh7.g = asfloat(shRaw4.w); s.sh.sh7.b = asfloat(shRaw5.x);
s.sh.sh8.r = asfloat(shRaw5.y); s.sh.sh8.g = asfloat(shRaw5.z); s.sh.sh8.b = asfloat(shRaw5.w);
s.sh.sh9.r = asfloat(shRaw6.x); s.sh.sh9.g = asfloat(shRaw6.y); s.sh.sh9.b = asfloat(shRaw6.z);
s.sh.sh10.r = asfloat(shRaw6.w); s.sh.sh10.g = asfloat(shRaw7.x); s.sh.sh10.b = asfloat(shRaw7.y);
s.sh.sh11.r = asfloat(shRaw7.z); s.sh.sh11.g = asfloat(shRaw7.w); s.sh.sh11.b = asfloat(shRaw8.x);
s.sh.sh12.r = asfloat(shRaw8.y); s.sh.sh12.g = asfloat(shRaw8.z); s.sh.sh12.b = asfloat(shRaw8.w);
s.sh.sh13.r = asfloat(shRaw9.x); s.sh.sh13.g = asfloat(shRaw9.y); s.sh.sh13.b = asfloat(shRaw9.z);
s.sh.sh14.r = asfloat(shRaw9.w); s.sh.sh14.g = asfloat(shRawA.x); s.sh.sh14.b = asfloat(shRawA.y);
s.sh.sh15.r = asfloat(shRawA.z); s.sh.sh15.g = asfloat(shRawA.w); s.sh.sh15.b = asfloat(shRawB);
}
else if (shFormat == VECTOR_FMT_16 || shFormat > VECTOR_FMT_6)
{
uint4 shRaw2 = _SplatSH.Load4(shOffset + 32);
uint4 shRaw3 = _SplatSH.Load4(shOffset + 48);
uint4 shRaw4 = _SplatSH.Load4(shOffset + 64);
uint3 shRaw5 = _SplatSH.Load3(shOffset + 80);
s.sh.sh1.r = f16tof32(shRaw0.x ); s.sh.sh1.g = f16tof32(shRaw0.x >> 16); s.sh.sh1.b = f16tof32(shRaw0.y );
s.sh.sh2.r = f16tof32(shRaw0.y >> 16); s.sh.sh2.g = f16tof32(shRaw0.z ); s.sh.sh2.b = f16tof32(shRaw0.z >> 16);
s.sh.sh3.r = f16tof32(shRaw0.w ); s.sh.sh3.g = f16tof32(shRaw0.w >> 16); s.sh.sh3.b = f16tof32(shRaw1.x );
s.sh.sh4.r = f16tof32(shRaw1.x >> 16); s.sh.sh4.g = f16tof32(shRaw1.y ); s.sh.sh4.b = f16tof32(shRaw1.y >> 16);
s.sh.sh5.r = f16tof32(shRaw1.z ); s.sh.sh5.g = f16tof32(shRaw1.z >> 16); s.sh.sh5.b = f16tof32(shRaw1.w );
s.sh.sh6.r = f16tof32(shRaw1.w >> 16); s.sh.sh6.g = f16tof32(shRaw2.x ); s.sh.sh6.b = f16tof32(shRaw2.x >> 16);
s.sh.sh7.r = f16tof32(shRaw2.y ); s.sh.sh7.g = f16tof32(shRaw2.y >> 16); s.sh.sh7.b = f16tof32(shRaw2.z );
s.sh.sh8.r = f16tof32(shRaw2.z >> 16); s.sh.sh8.g = f16tof32(shRaw2.w ); s.sh.sh8.b = f16tof32(shRaw2.w >> 16);
s.sh.sh9.r = f16tof32(shRaw3.x ); s.sh.sh9.g = f16tof32(shRaw3.x >> 16); s.sh.sh9.b = f16tof32(shRaw3.y );
s.sh.sh10.r = f16tof32(shRaw3.y >> 16); s.sh.sh10.g = f16tof32(shRaw3.z ); s.sh.sh10.b = f16tof32(shRaw3.z >> 16);
s.sh.sh11.r = f16tof32(shRaw3.w ); s.sh.sh11.g = f16tof32(shRaw3.w >> 16); s.sh.sh11.b = f16tof32(shRaw4.x );
s.sh.sh12.r = f16tof32(shRaw4.x >> 16); s.sh.sh12.g = f16tof32(shRaw4.y ); s.sh.sh12.b = f16tof32(shRaw4.y >> 16);
s.sh.sh13.r = f16tof32(shRaw4.z ); s.sh.sh13.g = f16tof32(shRaw4.z >> 16); s.sh.sh13.b = f16tof32(shRaw4.w );
s.sh.sh14.r = f16tof32(shRaw4.w >> 16); s.sh.sh14.g = f16tof32(shRaw5.x ); s.sh.sh14.b = f16tof32(shRaw5.x >> 16);
s.sh.sh15.r = f16tof32(shRaw5.y ); s.sh.sh15.g = f16tof32(shRaw5.y >> 16); s.sh.sh15.b = f16tof32(shRaw5.z );
}
else if (shFormat == VECTOR_FMT_11)
{
uint4 shRaw2 = _SplatSH.Load4(shOffset + 32);
uint3 shRaw3 = _SplatSH.Load3(shOffset + 48);
s.sh.sh1 = DecodePacked_11_10_11(shRaw0.x);
s.sh.sh2 = DecodePacked_11_10_11(shRaw0.y);
s.sh.sh3 = DecodePacked_11_10_11(shRaw0.z);
s.sh.sh4 = DecodePacked_11_10_11(shRaw0.w);
s.sh.sh5 = DecodePacked_11_10_11(shRaw1.x);
s.sh.sh6 = DecodePacked_11_10_11(shRaw1.y);
s.sh.sh7 = DecodePacked_11_10_11(shRaw1.z);
s.sh.sh8 = DecodePacked_11_10_11(shRaw1.w);
s.sh.sh9 = DecodePacked_11_10_11(shRaw2.x);
s.sh.sh10 = DecodePacked_11_10_11(shRaw2.y);
s.sh.sh11 = DecodePacked_11_10_11(shRaw2.z);
s.sh.sh12 = DecodePacked_11_10_11(shRaw2.w);
s.sh.sh13 = DecodePacked_11_10_11(shRaw3.x);
s.sh.sh14 = DecodePacked_11_10_11(shRaw3.y);
s.sh.sh15 = DecodePacked_11_10_11(shRaw3.z);
}
else if (shFormat == VECTOR_FMT_6)
{
s.sh.sh1 = DecodePacked_5_6_5(shRaw0.x);
s.sh.sh2 = DecodePacked_5_6_5(shRaw0.x >> 16);
s.sh.sh3 = DecodePacked_5_6_5(shRaw0.y);
s.sh.sh4 = DecodePacked_5_6_5(shRaw0.y >> 16);
s.sh.sh5 = DecodePacked_5_6_5(shRaw0.z);
s.sh.sh6 = DecodePacked_5_6_5(shRaw0.z >> 16);
s.sh.sh7 = DecodePacked_5_6_5(shRaw0.w);
s.sh.sh8 = DecodePacked_5_6_5(shRaw0.w >> 16);
s.sh.sh9 = DecodePacked_5_6_5(shRaw1.x);
s.sh.sh10 = DecodePacked_5_6_5(shRaw1.x >> 16);
s.sh.sh11 = DecodePacked_5_6_5(shRaw1.y);
s.sh.sh12 = DecodePacked_5_6_5(shRaw1.y >> 16);
s.sh.sh13 = DecodePacked_5_6_5(shRaw1.z);
s.sh.sh14 = DecodePacked_5_6_5(shRaw1.z >> 16);
s.sh.sh15 = DecodePacked_5_6_5(shRaw1.w);
}
// if raw data is chunk-relative, convert to final values by interpolating between chunk min/max
uint chunkIdx = idx / kChunkSize;
if (chunkIdx < _SplatChunkCount)
{
SplatChunkInfo chunk = _SplatChunks[chunkIdx];
float3 posMin = float3(chunk.posX.x, chunk.posY.x, chunk.posZ.x);
float3 posMax = float3(chunk.posX.y, chunk.posY.y, chunk.posZ.y);
half3 sclMin = half3(f16tof32(chunk.sclX ), f16tof32(chunk.sclY ), f16tof32(chunk.sclZ ));
half3 sclMax = half3(f16tof32(chunk.sclX>>16), f16tof32(chunk.sclY>>16), f16tof32(chunk.sclZ>>16));
half4 colMin = half4(f16tof32(chunk.colR ), f16tof32(chunk.colG ), f16tof32(chunk.colB ), f16tof32(chunk.colA ));
half4 colMax = half4(f16tof32(chunk.colR>>16), f16tof32(chunk.colG>>16), f16tof32(chunk.colB>>16), f16tof32(chunk.colA>>16));
half3 shMin = half3(f16tof32(chunk.shR ), f16tof32(chunk.shG ), f16tof32(chunk.shB ));
half3 shMax = half3(f16tof32(chunk.shR>>16), f16tof32(chunk.shG>>16), f16tof32(chunk.shB>>16));
s.pos = lerp(posMin, posMax, s.pos);
s.scale = lerp(sclMin, sclMax, s.scale);
s.scale *= s.scale;
s.scale *= s.scale;
s.scale *= s.scale;
col = lerp(colMin, colMax, col);
col.a = InvSquareCentered01(col.a);
if (shFormat > VECTOR_FMT_32F && shFormat <= VECTOR_FMT_6)
{
s.sh.sh1 = lerp(shMin, shMax, s.sh.sh1 );
s.sh.sh2 = lerp(shMin, shMax, s.sh.sh2 );
s.sh.sh3 = lerp(shMin, shMax, s.sh.sh3 );
s.sh.sh4 = lerp(shMin, shMax, s.sh.sh4 );
s.sh.sh5 = lerp(shMin, shMax, s.sh.sh5 );
s.sh.sh6 = lerp(shMin, shMax, s.sh.sh6 );
s.sh.sh7 = lerp(shMin, shMax, s.sh.sh7 );
s.sh.sh8 = lerp(shMin, shMax, s.sh.sh8 );
s.sh.sh9 = lerp(shMin, shMax, s.sh.sh9 );
s.sh.sh10 = lerp(shMin, shMax, s.sh.sh10);
s.sh.sh11 = lerp(shMin, shMax, s.sh.sh11);
s.sh.sh12 = lerp(shMin, shMax, s.sh.sh12);
s.sh.sh13 = lerp(shMin, shMax, s.sh.sh13);
s.sh.sh14 = lerp(shMin, shMax, s.sh.sh14);
s.sh.sh15 = lerp(shMin, shMax, s.sh.sh15);
}
}
s.opacity = col.a;
s.sh.col = col.rgb;
return s;
}
struct SplatViewData
{
float4 pos;
float2 axis1, axis2;
uint2 color; // 4xFP16
};
#endif // GAUSSIAN_SPLATTING_HLSL

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: d4087e54957693c48a7be32de91c99e2
ShaderIncludeImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,111 @@
// SPDX-License-Identifier: MIT
Shader "Gaussian Splatting/Render Splats"
{
SubShader
{
Tags { "RenderType"="Transparent" "Queue"="Transparent" }
Pass
{
ZWrite Off
Blend OneMinusDstAlpha One
Cull Off
CGPROGRAM
#pragma vertex vert
#pragma fragment frag
#pragma require compute
#pragma use_dxc
#include "GaussianSplatting.hlsl"
StructuredBuffer<uint> _OrderBuffer;
struct v2f
{
half4 col : COLOR0;
float2 pos : TEXCOORD0;
float4 vertex : SV_POSITION;
};
StructuredBuffer<SplatViewData> _SplatViewData;
ByteAddressBuffer _SplatSelectedBits;
uint _SplatBitsValid;
v2f vert (uint vtxID : SV_VertexID, uint instID : SV_InstanceID)
{
v2f o = (v2f)0;
instID = _OrderBuffer[instID];
SplatViewData view = _SplatViewData[instID];
float4 centerClipPos = view.pos;
bool behindCam = centerClipPos.w <= 0;
if (behindCam)
{
o.vertex = asfloat(0x7fc00000); // NaN discards the primitive
}
else
{
o.col.r = f16tof32(view.color.x >> 16);
o.col.g = f16tof32(view.color.x);
o.col.b = f16tof32(view.color.y >> 16);
o.col.a = f16tof32(view.color.y);
uint idx = vtxID;
float2 quadPos = float2(idx&1, (idx>>1)&1) * 2.0 - 1.0;
quadPos *= 2;
o.pos = quadPos;
float2 deltaScreenPos = (quadPos.x * view.axis1 + quadPos.y * view.axis2) * 2 / _ScreenParams.xy;
o.vertex = centerClipPos;
o.vertex.xy += deltaScreenPos * centerClipPos.w;
// is this splat selected?
if (_SplatBitsValid)
{
uint wordIdx = instID / 32;
uint bitIdx = instID & 31;
uint selVal = _SplatSelectedBits.Load(wordIdx * 4);
if (selVal & (1 << bitIdx))
{
o.col.a = -1;
}
}
}
return o;
}
half4 frag (v2f i) : SV_Target
{
float power = -dot(i.pos, i.pos);
half alpha = exp(power);
if (i.col.a >= 0)
{
alpha = saturate(alpha * i.col.a);
}
else
{
// "selected" splat: magenta outline, increase opacity, magenta tint
half3 selectedColor = half3(1,0,1);
if (alpha > 7.0/255.0)
{
if (alpha < 10.0/255.0)
{
alpha = 1;
i.col.rgb = selectedColor;
}
alpha = saturate(alpha + 0.3);
}
i.col.rgb = lerp(i.col.rgb, selectedColor, 0.5);
}
if (alpha < 1.0/255.0)
discard;
half4 res = half4(i.col.rgb * alpha, alpha);
return res;
}
ENDCG
}
}
}

View File

@@ -0,0 +1,9 @@
fileFormatVersion: 2
guid: ed800126ae8844a67aad1974ddddd59c
ShaderImporter:
externalObjects: {}
defaultTextures: []
nonModifiableTextures: []
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,959 @@
/******************************************************************************
* SortCommon
* Common functions for GPUSorting
*
* SPDX-License-Identifier: MIT
* Copyright Thomas Smith 5/17/2024
* https://github.com/b0nes164/GPUSorting
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
******************************************************************************/
#define KEYS_PER_THREAD 15U
#define D_DIM 256U
#define PART_SIZE 3840U
#define D_TOTAL_SMEM 4096U
#define RADIX 256U //Number of digit bins
#define RADIX_MASK 255U //Mask of digit bins
#define HALF_RADIX 128U //For smaller waves where bit packing is necessary
#define HALF_MASK 127U // ''
#define RADIX_LOG 8U //log2(RADIX)
#define RADIX_PASSES 4U //(Key width) / RADIX_LOG
cbuffer cbGpuSorting : register(b0)
{
uint e_numKeys;
uint e_radixShift;
uint e_threadBlocks;
uint padding;
};
#if defined(KEY_UINT)
RWStructuredBuffer<uint> b_sort;
RWStructuredBuffer<uint> b_alt;
#elif defined(KEY_INT)
RWStructuredBuffer<int> b_sort;
RWStructuredBuffer<int> b_alt;
#elif defined(KEY_FLOAT)
RWStructuredBuffer<float> b_sort;
RWStructuredBuffer<float> b_alt;
#endif
#if defined(PAYLOAD_UINT)
RWStructuredBuffer<uint> b_sortPayload;
RWStructuredBuffer<uint> b_altPayload;
#elif defined(PAYLOAD_INT)
RWStructuredBuffer<int> b_sortPayload;
RWStructuredBuffer<int> b_altPayload;
#elif defined(PAYLOAD_FLOAT)
RWStructuredBuffer<float> b_sortPayload;
RWStructuredBuffer<float> b_altPayload;
#endif
groupshared uint g_d[D_TOTAL_SMEM]; //Shared memory for DigitBinningPass and DownSweep kernels
struct KeyStruct
{
uint k[KEYS_PER_THREAD];
};
struct OffsetStruct
{
#if defined(ENABLE_16_BIT)
uint16_t o[KEYS_PER_THREAD];
#else
uint o[KEYS_PER_THREAD];
#endif
};
struct DigitStruct
{
#if defined(ENABLE_16_BIT)
uint16_t d[KEYS_PER_THREAD];
#else
uint d[KEYS_PER_THREAD];
#endif
};
//*****************************************************************************
//HELPER FUNCTIONS
//*****************************************************************************
//Due to a bug with SPIRV pre 1.6, we cannot use WaveGetLaneCount() to get the currently active wavesize
inline uint getWaveSize()
{
#if defined(VULKAN)
GroupMemoryBarrierWithGroupSync(); //Make absolutely sure the wave is not diverged here
return dot(countbits(WaveActiveBallot(true)), uint4(1, 1, 1, 1));
#else
return WaveGetLaneCount();
#endif
}
inline uint getWaveIndex(uint gtid, uint waveSize)
{
return gtid / waveSize;
}
//Radix Tricks by Michael Herf
//http://stereopsis.com/radix.html
inline uint FloatToUint(float f)
{
uint mask = -((int) (asuint(f) >> 31)) | 0x80000000;
return asuint(f) ^ mask;
}
inline float UintToFloat(uint u)
{
uint mask = ((u >> 31) - 1) | 0x80000000;
return asfloat(u ^ mask);
}
inline uint IntToUint(int i)
{
return asuint(i ^ 0x80000000);
}
inline int UintToInt(uint u)
{
return asint(u ^ 0x80000000);
}
inline uint getWaveCountPass(uint waveSize)
{
return D_DIM / waveSize;
}
inline uint ExtractDigit(uint key)
{
return key >> e_radixShift & RADIX_MASK;
}
inline uint ExtractDigit(uint key, uint shift)
{
return key >> shift & RADIX_MASK;
}
inline uint ExtractPackedIndex(uint key)
{
return key >> (e_radixShift + 1) & HALF_MASK;
}
inline uint ExtractPackedShift(uint key)
{
return (key >> e_radixShift & 1) ? 16 : 0;
}
inline uint ExtractPackedValue(uint packed, uint key)
{
return packed >> ExtractPackedShift(key) & 0xffff;
}
inline uint SubPartSizeWGE16(uint waveSize)
{
return KEYS_PER_THREAD * waveSize;
}
inline uint SharedOffsetWGE16(uint gtid, uint waveSize)
{
return WaveGetLaneIndex() + getWaveIndex(gtid, waveSize) * SubPartSizeWGE16(waveSize);
}
inline uint SubPartSizeWLT16(uint waveSize, uint _serialIterations)
{
return KEYS_PER_THREAD * waveSize * _serialIterations;
}
inline uint SharedOffsetWLT16(uint gtid, uint waveSize, uint _serialIterations)
{
return WaveGetLaneIndex() +
(getWaveIndex(gtid, waveSize) / _serialIterations * SubPartSizeWLT16(waveSize, _serialIterations)) +
(getWaveIndex(gtid, waveSize) % _serialIterations * waveSize);
}
inline uint DeviceOffsetWGE16(uint gtid, uint waveSize, uint partIndex)
{
return SharedOffsetWGE16(gtid, waveSize) + partIndex * PART_SIZE;
}
inline uint DeviceOffsetWLT16(uint gtid, uint waveSize, uint partIndex, uint serialIterations)
{
return SharedOffsetWLT16(gtid, waveSize, serialIterations) + partIndex * PART_SIZE;
}
inline uint GlobalHistOffset()
{
return e_radixShift << 5;
}
inline uint WaveHistsSizeWGE16(uint waveSize)
{
return D_DIM / waveSize * RADIX;
}
inline uint WaveHistsSizeWLT16()
{
return D_TOTAL_SMEM;
}
//*****************************************************************************
//FUNCTIONS COMMON TO THE DOWNSWEEP / DIGIT BINNING PASS
//*****************************************************************************
//If the size of a wave is too small, we do not have enough space in
//shared memory to assign a histogram to each wave, so instead,
//some operations are peformed serially.
inline uint SerialIterations(uint waveSize)
{
return (D_DIM / waveSize + 31) >> 5;
}
inline void ClearWaveHists(uint gtid, uint waveSize)
{
const uint histsEnd = waveSize >= 16 ?
WaveHistsSizeWGE16(waveSize) : WaveHistsSizeWLT16();
for (uint i = gtid; i < histsEnd; i += D_DIM)
g_d[i] = 0;
}
inline void LoadKey(inout uint key, uint index)
{
#if defined(KEY_UINT)
key = b_sort[index];
#elif defined(KEY_INT)
key = UintToInt(b_sort[index]);
#elif defined(KEY_FLOAT)
key = FloatToUint(b_sort[index]);
#endif
}
inline void LoadDummyKey(inout uint key)
{
key = 0xffffffff;
}
inline KeyStruct LoadKeysWGE16(uint gtid, uint waveSize, uint partIndex)
{
KeyStruct keys;
[unroll]
for (uint i = 0, t = DeviceOffsetWGE16(gtid, waveSize, partIndex);
i < KEYS_PER_THREAD;
++i, t += waveSize)
{
LoadKey(keys.k[i], t);
}
return keys;
}
inline KeyStruct LoadKeysWLT16(uint gtid, uint waveSize, uint partIndex, uint serialIterations)
{
KeyStruct keys;
[unroll]
for (uint i = 0, t = DeviceOffsetWLT16(gtid, waveSize, partIndex, serialIterations);
i < KEYS_PER_THREAD;
++i, t += waveSize * serialIterations)
{
LoadKey(keys.k[i], t);
}
return keys;
}
inline KeyStruct LoadKeysPartialWGE16(uint gtid, uint waveSize, uint partIndex)
{
KeyStruct keys;
[unroll]
for (uint i = 0, t = DeviceOffsetWGE16(gtid, waveSize, partIndex);
i < KEYS_PER_THREAD;
++i, t += waveSize)
{
if (t < e_numKeys)
LoadKey(keys.k[i], t);
else
LoadDummyKey(keys.k[i]);
}
return keys;
}
inline KeyStruct LoadKeysPartialWLT16(uint gtid, uint waveSize, uint partIndex, uint serialIterations)
{
KeyStruct keys;
[unroll]
for (uint i = 0, t = DeviceOffsetWLT16(gtid, waveSize, partIndex, serialIterations);
i < KEYS_PER_THREAD;
++i, t += waveSize * serialIterations)
{
if (t < e_numKeys)
LoadKey(keys.k[i], t);
else
LoadDummyKey(keys.k[i]);
}
return keys;
}
inline uint WaveFlagsWGE16(uint waveSize)
{
return (waveSize & 31) ? (1U << waveSize) - 1 : 0xffffffff;
}
inline uint WaveFlagsWLT16(uint waveSize)
{
return (1U << waveSize) - 1;;
}
inline void WarpLevelMultiSplitWGE16(uint key, inout uint4 waveFlags)
{
[unroll]
for (uint k = 0; k < RADIX_LOG; ++k)
{
const uint currentBit = 1 << k + e_radixShift;
const bool t = (key & currentBit) != 0;
GroupMemoryBarrierWithGroupSync(); //Play on the safe side, throw in a barrier for convergence
const uint4 ballot = WaveActiveBallot(t);
if(t)
waveFlags &= ballot;
else
waveFlags &= (~ballot);
}
}
inline uint2 CountBitsWGE16(uint waveSize, uint ltMask, uint4 waveFlags)
{
uint2 count = uint2(0, 0);
for(uint wavePart = 0; wavePart < waveSize; wavePart += 32)
{
uint t = countbits(waveFlags[wavePart >> 5]);
if (WaveGetLaneIndex() >= wavePart)
{
if (WaveGetLaneIndex() >= wavePart + 32)
count.x += t;
else
count.x += countbits(waveFlags[wavePart >> 5] & ltMask);
}
count.y += t;
}
return count;
}
inline void WarpLevelMultiSplitWLT16(uint key, inout uint waveFlags)
{
[unroll]
for (uint k = 0; k < RADIX_LOG; ++k)
{
const bool t = key >> (k + e_radixShift) & 1;
waveFlags &= (t ? 0 : 0xffffffff) ^ (uint) WaveActiveBallot(t);
}
}
inline OffsetStruct RankKeysWGE16(
uint waveSize,
uint waveOffset,
KeyStruct keys)
{
OffsetStruct offsets;
const uint initialFlags = WaveFlagsWGE16(waveSize);
const uint ltMask = (1U << (WaveGetLaneIndex() & 31)) - 1;
[unroll]
for (uint i = 0; i < KEYS_PER_THREAD; ++i)
{
uint4 waveFlags = initialFlags;
WarpLevelMultiSplitWGE16(keys.k[i], waveFlags);
const uint index = ExtractDigit(keys.k[i]) + waveOffset;
const uint2 bitCount = CountBitsWGE16(waveSize, ltMask, waveFlags);
offsets.o[i] = g_d[index] + bitCount.x;
GroupMemoryBarrierWithGroupSync();
if (bitCount.x == 0)
g_d[index] += bitCount.y;
GroupMemoryBarrierWithGroupSync();
}
return offsets;
}
inline OffsetStruct RankKeysWLT16(uint waveSize, uint waveIndex, KeyStruct keys, uint serialIterations)
{
OffsetStruct offsets;
const uint ltMask = (1U << WaveGetLaneIndex()) - 1;
const uint initialFlags = WaveFlagsWLT16(waveSize);
[unroll]
for (uint i = 0; i < KEYS_PER_THREAD; ++i)
{
uint waveFlags = initialFlags;
WarpLevelMultiSplitWLT16(keys.k[i], waveFlags);
const uint index = ExtractPackedIndex(keys.k[i]) +
(waveIndex / serialIterations * HALF_RADIX);
const uint peerBits = countbits(waveFlags & ltMask);
for (uint k = 0; k < serialIterations; ++k)
{
if (waveIndex % serialIterations == k)
offsets.o[i] = ExtractPackedValue(g_d[index], keys.k[i]) + peerBits;
GroupMemoryBarrierWithGroupSync();
if (waveIndex % serialIterations == k && peerBits == 0)
{
InterlockedAdd(g_d[index],
countbits(waveFlags) << ExtractPackedShift(keys.k[i]));
}
GroupMemoryBarrierWithGroupSync();
}
}
return offsets;
}
inline uint WaveHistInclusiveScanCircularShiftWGE16(uint gtid, uint waveSize)
{
uint histReduction = g_d[gtid];
for (uint i = gtid + RADIX; i < WaveHistsSizeWGE16(waveSize); i += RADIX)
{
histReduction += g_d[i];
g_d[i] = histReduction - g_d[i];
}
return histReduction;
}
inline uint WaveHistInclusiveScanCircularShiftWLT16(uint gtid)
{
uint histReduction = g_d[gtid];
for (uint i = gtid + HALF_RADIX; i < WaveHistsSizeWLT16(); i += HALF_RADIX)
{
histReduction += g_d[i];
g_d[i] = histReduction - g_d[i];
}
return histReduction;
}
inline void WaveHistReductionExclusiveScanWGE16(uint gtid, uint waveSize, uint histReduction)
{
if (gtid < RADIX)
{
const uint laneMask = waveSize - 1;
g_d[((WaveGetLaneIndex() + 1) & laneMask) + (gtid & ~laneMask)] = histReduction;
}
GroupMemoryBarrierWithGroupSync();
if (gtid < RADIX / waveSize)
{
g_d[gtid * waveSize] =
WavePrefixSum(g_d[gtid * waveSize]);
}
GroupMemoryBarrierWithGroupSync();
uint t = WaveReadLaneAt(g_d[gtid], 0);
if (gtid < RADIX && WaveGetLaneIndex())
g_d[gtid] += t;
}
//inclusive/exclusive prefix sum up the histograms,
//use a blelloch scan for in place packed exclusive
inline void WaveHistReductionExclusiveScanWLT16(uint gtid)
{
uint shift = 1;
for (uint j = RADIX >> 2; j > 0; j >>= 1)
{
GroupMemoryBarrierWithGroupSync();
if (gtid < j)
{
g_d[((((gtid << 1) + 2) << shift) - 1) >> 1] +=
g_d[((((gtid << 1) + 1) << shift) - 1) >> 1] & 0xffff0000;
}
shift++;
}
GroupMemoryBarrierWithGroupSync();
if (gtid == 0)
g_d[HALF_RADIX - 1] &= 0xffff;
for (uint j = 1; j < RADIX >> 1; j <<= 1)
{
--shift;
GroupMemoryBarrierWithGroupSync();
if (gtid < j)
{
const uint t = ((((gtid << 1) + 1) << shift) - 1) >> 1;
const uint t2 = ((((gtid << 1) + 2) << shift) - 1) >> 1;
const uint t3 = g_d[t];
g_d[t] = (g_d[t] & 0xffff) | (g_d[t2] & 0xffff0000);
g_d[t2] += t3 & 0xffff0000;
}
}
GroupMemoryBarrierWithGroupSync();
if (gtid < HALF_RADIX)
{
const uint t = g_d[gtid];
g_d[gtid] = (t >> 16) + (t << 16) + (t & 0xffff0000);
}
}
inline void UpdateOffsetsWGE16(
uint gtid,
uint waveSize,
inout OffsetStruct offsets,
KeyStruct keys)
{
if (gtid >= waveSize)
{
const uint t = getWaveIndex(gtid, waveSize) * RADIX;
[unroll]
for (uint i = 0; i < KEYS_PER_THREAD; ++i)
{
const uint t2 = ExtractDigit(keys.k[i]);
offsets.o[i] += g_d[t2 + t] + g_d[t2];
}
}
else
{
[unroll]
for (uint i = 0; i < KEYS_PER_THREAD; ++i)
offsets.o[i] += g_d[ExtractDigit(keys.k[i])];
}
}
inline void UpdateOffsetsWLT16(
uint gtid,
uint waveSize,
uint serialIterations,
inout OffsetStruct offsets,
KeyStruct keys)
{
if (gtid >= waveSize * serialIterations)
{
const uint t = getWaveIndex(gtid, waveSize) / serialIterations * HALF_RADIX;
[unroll]
for (uint i = 0; i < KEYS_PER_THREAD; ++i)
{
const uint t2 = ExtractPackedIndex(keys.k[i]);
offsets.o[i] += ExtractPackedValue(g_d[t2 + t] + g_d[t2], keys.k[i]);
}
}
else
{
[unroll]
for (uint i = 0; i < KEYS_PER_THREAD; ++i)
offsets.o[i] += ExtractPackedValue(g_d[ExtractPackedIndex(keys.k[i])], keys.k[i]);
}
}
inline void ScatterKeysShared(OffsetStruct offsets, KeyStruct keys)
{
[unroll]
for (uint i = 0; i < KEYS_PER_THREAD; ++i)
g_d[offsets.o[i]] = keys.k[i];
}
inline uint DescendingIndex(uint deviceIndex)
{
return e_numKeys - deviceIndex - 1;
}
inline void WriteKey(uint deviceIndex, uint groupSharedIndex)
{
#if defined(KEY_UINT)
b_alt[deviceIndex] = g_d[groupSharedIndex];
#elif defined(KEY_INT)
b_alt[deviceIndex] = UintToInt(g_d[groupSharedIndex]);
#elif defined(KEY_FLOAT)
b_alt[deviceIndex] = UintToFloat(g_d[groupSharedIndex]);
#endif
}
inline void LoadPayload(inout uint payload, uint deviceIndex)
{
#if defined(PAYLOAD_UINT)
payload = b_sortPayload[deviceIndex];
#elif defined(PAYLOAD_INT) || defined(PAYLOAD_FLOAT)
payload = asuint(b_sortPayload[deviceIndex]);
#endif
}
inline void ScatterPayloadsShared(OffsetStruct offsets, KeyStruct payloads)
{
ScatterKeysShared(offsets, payloads);
}
inline void WritePayload(uint deviceIndex, uint groupSharedIndex)
{
#if defined(PAYLOAD_UINT)
b_altPayload[deviceIndex] = g_d[groupSharedIndex];
#elif defined(PAYLOAD_INT)
b_altPayload[deviceIndex] = asint(g_d[groupSharedIndex]);
#elif defined(PAYLOAD_FLOAT)
b_altPayload[deviceIndex] = asfloat(g_d[groupSharedIndex]);
#endif
}
//*****************************************************************************
//SCATTERING: FULL PARTITIONS
//*****************************************************************************
//KEYS ONLY
inline void ScatterKeysOnlyDeviceAscending(uint gtid)
{
for (uint i = gtid; i < PART_SIZE; i += D_DIM)
WriteKey(g_d[ExtractDigit(g_d[i]) + PART_SIZE] + i, i);
}
inline void ScatterKeysOnlyDeviceDescending(uint gtid)
{
if (e_radixShift == 24)
{
for (uint i = gtid; i < PART_SIZE; i += D_DIM)
WriteKey(DescendingIndex(g_d[ExtractDigit(g_d[i]) + PART_SIZE] + i), i);
}
else
{
ScatterKeysOnlyDeviceAscending(gtid);
}
}
inline void ScatterKeysOnlyDevice(uint gtid)
{
#if defined(SHOULD_ASCEND)
ScatterKeysOnlyDeviceAscending(gtid);
#else
ScatterKeysOnlyDeviceDescending(gtid);
#endif
}
//KEY VALUE PAIRS
inline void ScatterPairsKeyPhaseAscending(
uint gtid,
inout DigitStruct digits)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
{
digits.d[i] = ExtractDigit(g_d[t]);
WriteKey(g_d[digits.d[i] + PART_SIZE] + t, t);
}
}
inline void ScatterPairsKeyPhaseDescending(
uint gtid,
inout DigitStruct digits)
{
if (e_radixShift == 24)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
{
digits.d[i] = ExtractDigit(g_d[t]);
WriteKey(DescendingIndex(g_d[digits.d[i] + PART_SIZE] + t), t);
}
}
else
{
ScatterPairsKeyPhaseAscending(gtid, digits);
}
}
inline void LoadPayloadsWGE16(
uint gtid,
uint waveSize,
uint partIndex,
inout KeyStruct payloads)
{
[unroll]
for (uint i = 0, t = DeviceOffsetWGE16(gtid, waveSize, partIndex);
i < KEYS_PER_THREAD;
++i, t += waveSize)
{
LoadPayload(payloads.k[i], t);
}
}
inline void LoadPayloadsWLT16(
uint gtid,
uint waveSize,
uint partIndex,
uint serialIterations,
inout KeyStruct payloads)
{
[unroll]
for (uint i = 0, t = DeviceOffsetWLT16(gtid, waveSize, partIndex, serialIterations);
i < KEYS_PER_THREAD;
++i, t += waveSize * serialIterations)
{
LoadPayload(payloads.k[i], t);
}
}
inline void ScatterPayloadsAscending(uint gtid, DigitStruct digits)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
WritePayload(g_d[digits.d[i] + PART_SIZE] + t, t);
}
inline void ScatterPayloadsDescending(uint gtid, DigitStruct digits)
{
if (e_radixShift == 24)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
WritePayload(DescendingIndex(g_d[digits.d[i] + PART_SIZE] + t), t);
}
else
{
ScatterPayloadsAscending(gtid, digits);
}
}
inline void ScatterPairsDevice(
uint gtid,
uint waveSize,
uint partIndex,
OffsetStruct offsets)
{
DigitStruct digits;
#if defined(SHOULD_ASCEND)
ScatterPairsKeyPhaseAscending(gtid, digits);
#else
ScatterPairsKeyPhaseDescending(gtid, digits);
#endif
GroupMemoryBarrierWithGroupSync();
KeyStruct payloads;
if (waveSize >= 16)
LoadPayloadsWGE16(gtid, waveSize, partIndex, payloads);
else
LoadPayloadsWLT16(gtid, waveSize, partIndex, SerialIterations(waveSize), payloads);
ScatterPayloadsShared(offsets, payloads);
GroupMemoryBarrierWithGroupSync();
#if defined(SHOULD_ASCEND)
ScatterPayloadsAscending(gtid, digits);
#else
ScatterPayloadsDescending(gtid, digits);
#endif
}
inline void ScatterDevice(
uint gtid,
uint waveSize,
uint partIndex,
OffsetStruct offsets)
{
#if defined(SORT_PAIRS)
ScatterPairsDevice(
gtid,
waveSize,
partIndex,
offsets);
#else
ScatterKeysOnlyDevice(gtid);
#endif
}
//*****************************************************************************
//SCATTERING: PARTIAL PARTITIONS
//*****************************************************************************
//KEYS ONLY
inline void ScatterKeysOnlyDevicePartialAscending(uint gtid, uint finalPartSize)
{
for (uint i = gtid; i < PART_SIZE; i += D_DIM)
{
if (i < finalPartSize)
WriteKey(g_d[ExtractDigit(g_d[i]) + PART_SIZE] + i, i);
}
}
inline void ScatterKeysOnlyDevicePartialDescending(uint gtid, uint finalPartSize)
{
if (e_radixShift == 24)
{
for (uint i = gtid; i < PART_SIZE; i += D_DIM)
{
if (i < finalPartSize)
WriteKey(DescendingIndex(g_d[ExtractDigit(g_d[i]) + PART_SIZE] + i), i);
}
}
else
{
ScatterKeysOnlyDevicePartialAscending(gtid, finalPartSize);
}
}
inline void ScatterKeysOnlyDevicePartial(uint gtid, uint partIndex)
{
const uint finalPartSize = e_numKeys - partIndex * PART_SIZE;
#if defined(SHOULD_ASCEND)
ScatterKeysOnlyDevicePartialAscending(gtid, finalPartSize);
#else
ScatterKeysOnlyDevicePartialDescending(gtid, finalPartSize);
#endif
}
//KEY VALUE PAIRS
inline void ScatterPairsKeyPhaseAscendingPartial(
uint gtid,
uint finalPartSize,
inout DigitStruct digits)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
{
if (t < finalPartSize)
{
digits.d[i] = ExtractDigit(g_d[t]);
WriteKey(g_d[digits.d[i] + PART_SIZE] + t, t);
}
}
}
inline void ScatterPairsKeyPhaseDescendingPartial(
uint gtid,
uint finalPartSize,
inout DigitStruct digits)
{
if (e_radixShift == 24)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
{
if (t < finalPartSize)
{
digits.d[i] = ExtractDigit(g_d[t]);
WriteKey(DescendingIndex(g_d[digits.d[i] + PART_SIZE] + t), t);
}
}
}
else
{
ScatterPairsKeyPhaseAscendingPartial(gtid, finalPartSize, digits);
}
}
inline void LoadPayloadsPartialWGE16(
uint gtid,
uint waveSize,
uint partIndex,
inout KeyStruct payloads)
{
[unroll]
for (uint i = 0, t = DeviceOffsetWGE16(gtid, waveSize, partIndex);
i < KEYS_PER_THREAD;
++i, t += waveSize)
{
if (t < e_numKeys)
LoadPayload(payloads.k[i], t);
}
}
inline void LoadPayloadsPartialWLT16(
uint gtid,
uint waveSize,
uint partIndex,
uint serialIterations,
inout KeyStruct payloads)
{
[unroll]
for (uint i = 0, t = DeviceOffsetWLT16(gtid, waveSize, partIndex, serialIterations);
i < KEYS_PER_THREAD;
++i, t += waveSize * serialIterations)
{
if (t < e_numKeys)
LoadPayload(payloads.k[i], t);
}
}
inline void ScatterPayloadsAscendingPartial(
uint gtid,
uint finalPartSize,
DigitStruct digits)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
{
if (t < finalPartSize)
WritePayload(g_d[digits.d[i] + PART_SIZE] + t, t);
}
}
inline void ScatterPayloadsDescendingPartial(
uint gtid,
uint finalPartSize,
DigitStruct digits)
{
if (e_radixShift == 24)
{
[unroll]
for (uint i = 0, t = gtid; i < KEYS_PER_THREAD; ++i, t += D_DIM)
{
if (t < finalPartSize)
WritePayload(DescendingIndex(g_d[digits.d[i] + PART_SIZE] + t), t);
}
}
else
{
ScatterPayloadsAscendingPartial(gtid, finalPartSize, digits);
}
}
inline void ScatterPairsDevicePartial(
uint gtid,
uint waveSize,
uint partIndex,
OffsetStruct offsets)
{
DigitStruct digits;
const uint finalPartSize = e_numKeys - partIndex * PART_SIZE;
#if defined(SHOULD_ASCEND)
ScatterPairsKeyPhaseAscendingPartial(gtid, finalPartSize, digits);
#else
ScatterPairsKeyPhaseDescendingPartial(gtid, finalPartSize, digits);
#endif
GroupMemoryBarrierWithGroupSync();
KeyStruct payloads;
if (waveSize >= 16)
LoadPayloadsPartialWGE16(gtid, waveSize, partIndex, payloads);
else
LoadPayloadsPartialWLT16(gtid, waveSize, partIndex, SerialIterations(waveSize), payloads);
ScatterPayloadsShared(offsets, payloads);
GroupMemoryBarrierWithGroupSync();
#if defined(SHOULD_ASCEND)
ScatterPayloadsAscendingPartial(gtid, finalPartSize, digits);
#else
ScatterPayloadsDescendingPartial(gtid, finalPartSize, digits);
#endif
}
inline void ScatterDevicePartial(
uint gtid,
uint waveSize,
uint partIndex,
OffsetStruct offsets)
{
#if defined(SORT_PAIRS)
ScatterPairsDevicePartial(
gtid,
waveSize,
partIndex,
offsets);
#else
ScatterKeysOnlyDevicePartial(gtid, partIndex);
#endif
}

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: 268e5936ab6d79f4b8aeef8f5d14e7ee
ShaderIncludeImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,212 @@
// SPDX-License-Identifier: MIT
#ifndef SPHERICAL_HARMONICS_HLSL
#define SPHERICAL_HARMONICS_HLSL
// SH rotation based on https://github.com/andrewwillmott/sh-lib (Unlicense / public domain)
#define SH_MAX_ORDER 4
#define SH_MAX_COEFFS_COUNT (SH_MAX_ORDER*SH_MAX_ORDER)
float3 Dot3(int vidx, float3 v[SH_MAX_COEFFS_COUNT], float f[3])
{
return v[vidx+0] * f[0] + v[vidx+1] * f[1] + v[vidx+2] * f[2];
}
float3 Dot5(int vidx, float3 v[SH_MAX_COEFFS_COUNT], float f[5])
{
return v[vidx+0] * f[0] + v[vidx+1] * f[1] + v[vidx+2] * f[2] + v[vidx+3] * f[3] + v[vidx+4] * f[4];
}
float3 Dot7(int vidx, float3 v[SH_MAX_COEFFS_COUNT], float f[7])
{
return v[vidx+0] * f[0] + v[vidx+1] * f[1] + v[vidx+2] * f[2] + v[vidx+3] * f[3] + v[vidx+4] * f[4] + v[vidx+5] * f[5] + v[vidx+6] * f[6];
}
void RotateSH(float3x3 orient, int n, float3 coeffsIn[SH_MAX_COEFFS_COUNT], out float3 coeffs[SH_MAX_COEFFS_COUNT])
{
const float kSqrt03_02 = sqrt( 3.0 / 2.0);
const float kSqrt01_03 = sqrt( 1.0 / 3.0);
const float kSqrt02_03 = sqrt( 2.0 / 3.0);
const float kSqrt04_03 = sqrt( 4.0 / 3.0);
const float kSqrt01_04 = sqrt( 1.0 / 4.0);
const float kSqrt03_04 = sqrt( 3.0 / 4.0);
const float kSqrt01_05 = sqrt( 1.0 / 5.0);
const float kSqrt03_05 = sqrt( 3.0 / 5.0);
const float kSqrt06_05 = sqrt( 6.0 / 5.0);
const float kSqrt08_05 = sqrt( 8.0 / 5.0);
const float kSqrt09_05 = sqrt( 9.0 / 5.0);
const float kSqrt05_06 = sqrt( 5.0 / 6.0);
const float kSqrt01_06 = sqrt( 1.0 / 6.0);
const float kSqrt03_08 = sqrt( 3.0 / 8.0);
const float kSqrt05_08 = sqrt( 5.0 / 8.0);
const float kSqrt07_08 = sqrt( 7.0 / 8.0);
const float kSqrt09_08 = sqrt( 9.0 / 8.0);
const float kSqrt05_09 = sqrt( 5.0 / 9.0);
const float kSqrt08_09 = sqrt( 8.0 / 9.0);
const float kSqrt01_10 = sqrt( 1.0 / 10.0);
const float kSqrt03_10 = sqrt( 3.0 / 10.0);
const float kSqrt01_12 = sqrt( 1.0 / 12.0);
const float kSqrt04_15 = sqrt( 4.0 / 15.0);
const float kSqrt01_16 = sqrt( 1.0 / 16.0);
const float kSqrt07_16 = sqrt( 7.0 / 16.0);
const float kSqrt15_16 = sqrt(15.0 / 16.0);
const float kSqrt01_18 = sqrt( 1.0 / 18.0);
const float kSqrt03_25 = sqrt( 3.0 / 25.0);
const float kSqrt14_25 = sqrt(14.0 / 25.0);
const float kSqrt15_25 = sqrt(15.0 / 25.0);
const float kSqrt18_25 = sqrt(18.0 / 25.0);
const float kSqrt01_32 = sqrt( 1.0 / 32.0);
const float kSqrt03_32 = sqrt( 3.0 / 32.0);
const float kSqrt15_32 = sqrt(15.0 / 32.0);
const float kSqrt21_32 = sqrt(21.0 / 32.0);
const float kSqrt01_50 = sqrt( 1.0 / 50.0);
const float kSqrt03_50 = sqrt( 3.0 / 50.0);
const float kSqrt21_50 = sqrt(21.0 / 50.0);
int srcIdx = 0;
int dstIdx = 0;
// band 0
coeffs[dstIdx++] = coeffsIn[0];
if (n < 2)
return;
// band 1
srcIdx += 1;
float sh1[3][3] =
{
// NOTE: change from upstream code at https://github.com/andrewwillmott/sh-lib, some of the
// values need to have "-" in front of them.
orient._22, -orient._23, orient._21,
-orient._32, orient._33, -orient._31,
orient._12, -orient._13, orient._11
};
coeffs[dstIdx++] = Dot3(srcIdx, coeffsIn, sh1[0]);
coeffs[dstIdx++] = Dot3(srcIdx, coeffsIn, sh1[1]);
coeffs[dstIdx++] = Dot3(srcIdx, coeffsIn, sh1[2]);
if (n < 3)
return;
// band 2
srcIdx += 3;
float sh2[5][5];
sh2[0][0] = kSqrt01_04 * ((sh1[2][2] * sh1[0][0] + sh1[2][0] * sh1[0][2]) + (sh1[0][2] * sh1[2][0] + sh1[0][0] * sh1[2][2]));
sh2[0][1] = (sh1[2][1] * sh1[0][0] + sh1[0][1] * sh1[2][0]);
sh2[0][2] = kSqrt03_04 * (sh1[2][1] * sh1[0][1] + sh1[0][1] * sh1[2][1]);
sh2[0][3] = (sh1[2][1] * sh1[0][2] + sh1[0][1] * sh1[2][2]);
sh2[0][4] = kSqrt01_04 * ((sh1[2][2] * sh1[0][2] - sh1[2][0] * sh1[0][0]) + (sh1[0][2] * sh1[2][2] - sh1[0][0] * sh1[2][0]));
coeffs[dstIdx++] = Dot5(srcIdx, coeffsIn, sh2[0]);
sh2[1][0] = kSqrt01_04 * ((sh1[1][2] * sh1[0][0] + sh1[1][0] * sh1[0][2]) + (sh1[0][2] * sh1[1][0] + sh1[0][0] * sh1[1][2]));
sh2[1][1] = sh1[1][1] * sh1[0][0] + sh1[0][1] * sh1[1][0];
sh2[1][2] = kSqrt03_04 * (sh1[1][1] * sh1[0][1] + sh1[0][1] * sh1[1][1]);
sh2[1][3] = sh1[1][1] * sh1[0][2] + sh1[0][1] * sh1[1][2];
sh2[1][4] = kSqrt01_04 * ((sh1[1][2] * sh1[0][2] - sh1[1][0] * sh1[0][0]) + (sh1[0][2] * sh1[1][2] - sh1[0][0] * sh1[1][0]));
coeffs[dstIdx++] = Dot5(srcIdx, coeffsIn, sh2[1]);
sh2[2][0] = kSqrt01_03 * (sh1[1][2] * sh1[1][0] + sh1[1][0] * sh1[1][2]) + -kSqrt01_12 * ((sh1[2][2] * sh1[2][0] + sh1[2][0] * sh1[2][2]) + (sh1[0][2] * sh1[0][0] + sh1[0][0] * sh1[0][2]));
sh2[2][1] = kSqrt04_03 * sh1[1][1] * sh1[1][0] + -kSqrt01_03 * (sh1[2][1] * sh1[2][0] + sh1[0][1] * sh1[0][0]);
sh2[2][2] = sh1[1][1] * sh1[1][1] + -kSqrt01_04 * (sh1[2][1] * sh1[2][1] + sh1[0][1] * sh1[0][1]);
sh2[2][3] = kSqrt04_03 * sh1[1][1] * sh1[1][2] + -kSqrt01_03 * (sh1[2][1] * sh1[2][2] + sh1[0][1] * sh1[0][2]);
sh2[2][4] = kSqrt01_03 * (sh1[1][2] * sh1[1][2] - sh1[1][0] * sh1[1][0]) + -kSqrt01_12 * ((sh1[2][2] * sh1[2][2] - sh1[2][0] * sh1[2][0]) + (sh1[0][2] * sh1[0][2] - sh1[0][0] * sh1[0][0]));
coeffs[dstIdx++] = Dot5(srcIdx, coeffsIn, sh2[2]);
sh2[3][0] = kSqrt01_04 * ((sh1[1][2] * sh1[2][0] + sh1[1][0] * sh1[2][2]) + (sh1[2][2] * sh1[1][0] + sh1[2][0] * sh1[1][2]));
sh2[3][1] = sh1[1][1] * sh1[2][0] + sh1[2][1] * sh1[1][0];
sh2[3][2] = kSqrt03_04 * (sh1[1][1] * sh1[2][1] + sh1[2][1] * sh1[1][1]);
sh2[3][3] = sh1[1][1] * sh1[2][2] + sh1[2][1] * sh1[1][2];
sh2[3][4] = kSqrt01_04 * ((sh1[1][2] * sh1[2][2] - sh1[1][0] * sh1[2][0]) + (sh1[2][2] * sh1[1][2] - sh1[2][0] * sh1[1][0]));
coeffs[dstIdx++] = Dot5(srcIdx, coeffsIn, sh2[3]);
sh2[4][0] = kSqrt01_04 * ((sh1[2][2] * sh1[2][0] + sh1[2][0] * sh1[2][2]) - (sh1[0][2] * sh1[0][0] + sh1[0][0] * sh1[0][2]));
sh2[4][1] = (sh1[2][1] * sh1[2][0] - sh1[0][1] * sh1[0][0]);
sh2[4][2] = kSqrt03_04 * (sh1[2][1] * sh1[2][1] - sh1[0][1] * sh1[0][1]);
sh2[4][3] = (sh1[2][1] * sh1[2][2] - sh1[0][1] * sh1[0][2]);
sh2[4][4] = kSqrt01_04 * ((sh1[2][2] * sh1[2][2] - sh1[2][0] * sh1[2][0]) - (sh1[0][2] * sh1[0][2] - sh1[0][0] * sh1[0][0]));
coeffs[dstIdx++] = Dot5(srcIdx, coeffsIn, sh2[4]);
if (n < 4)
return;
// band 3
srcIdx += 5;
float sh3[7][7];
sh3[0][0] = kSqrt01_04 * ((sh1[2][2] * sh2[0][0] + sh1[2][0] * sh2[0][4]) + (sh1[0][2] * sh2[4][0] + sh1[0][0] * sh2[4][4]));
sh3[0][1] = kSqrt03_02 * (sh1[2][1] * sh2[0][0] + sh1[0][1] * sh2[4][0]);
sh3[0][2] = kSqrt15_16 * (sh1[2][1] * sh2[0][1] + sh1[0][1] * sh2[4][1]);
sh3[0][3] = kSqrt05_06 * (sh1[2][1] * sh2[0][2] + sh1[0][1] * sh2[4][2]);
sh3[0][4] = kSqrt15_16 * (sh1[2][1] * sh2[0][3] + sh1[0][1] * sh2[4][3]);
sh3[0][5] = kSqrt03_02 * (sh1[2][1] * sh2[0][4] + sh1[0][1] * sh2[4][4]);
sh3[0][6] = kSqrt01_04 * ((sh1[2][2] * sh2[0][4] - sh1[2][0] * sh2[0][0]) + (sh1[0][2] * sh2[4][4] - sh1[0][0] * sh2[4][0]));
coeffs[dstIdx++] = Dot7(srcIdx, coeffsIn, sh3[0]);
sh3[1][0] = kSqrt01_06 * (sh1[1][2] * sh2[0][0] + sh1[1][0] * sh2[0][4]) + kSqrt01_06 * ((sh1[2][2] * sh2[1][0] + sh1[2][0] * sh2[1][4]) + (sh1[0][2] * sh2[3][0] + sh1[0][0] * sh2[3][4]));
sh3[1][1] = sh1[1][1] * sh2[0][0] + (sh1[2][1] * sh2[1][0] + sh1[0][1] * sh2[3][0]);
sh3[1][2] = kSqrt05_08 * sh1[1][1] * sh2[0][1] + kSqrt05_08 * (sh1[2][1] * sh2[1][1] + sh1[0][1] * sh2[3][1]);
sh3[1][3] = kSqrt05_09 * sh1[1][1] * sh2[0][2] + kSqrt05_09 * (sh1[2][1] * sh2[1][2] + sh1[0][1] * sh2[3][2]);
sh3[1][4] = kSqrt05_08 * sh1[1][1] * sh2[0][3] + kSqrt05_08 * (sh1[2][1] * sh2[1][3] + sh1[0][1] * sh2[3][3]);
sh3[1][5] = sh1[1][1] * sh2[0][4] + (sh1[2][1] * sh2[1][4] + sh1[0][1] * sh2[3][4]);
sh3[1][6] = kSqrt01_06 * (sh1[1][2] * sh2[0][4] - sh1[1][0] * sh2[0][0]) + kSqrt01_06 * ((sh1[2][2] * sh2[1][4] - sh1[2][0] * sh2[1][0]) + (sh1[0][2] * sh2[3][4] - sh1[0][0] * sh2[3][0]));
coeffs[dstIdx++] = Dot7(srcIdx, coeffsIn, sh3[1]);
sh3[2][0] = kSqrt04_15 * (sh1[1][2] * sh2[1][0] + sh1[1][0] * sh2[1][4]) + kSqrt01_05 * (sh1[0][2] * sh2[2][0] + sh1[0][0] * sh2[2][4]) + -sqrt(1.0 / 60.0) * ((sh1[2][2] * sh2[0][0] + sh1[2][0] * sh2[0][4]) - (sh1[0][2] * sh2[4][0] + sh1[0][0] * sh2[4][4]));
sh3[2][1] = kSqrt08_05 * sh1[1][1] * sh2[1][0] + kSqrt06_05 * sh1[0][1] * sh2[2][0] + -kSqrt01_10 * (sh1[2][1] * sh2[0][0] - sh1[0][1] * sh2[4][0]);
sh3[2][2] = sh1[1][1] * sh2[1][1] + kSqrt03_04 * sh1[0][1] * sh2[2][1] + -kSqrt01_16 * (sh1[2][1] * sh2[0][1] - sh1[0][1] * sh2[4][1]);
sh3[2][3] = kSqrt08_09 * sh1[1][1] * sh2[1][2] + kSqrt02_03 * sh1[0][1] * sh2[2][2] + -kSqrt01_18 * (sh1[2][1] * sh2[0][2] - sh1[0][1] * sh2[4][2]);
sh3[2][4] = sh1[1][1] * sh2[1][3] + kSqrt03_04 * sh1[0][1] * sh2[2][3] + -kSqrt01_16 * (sh1[2][1] * sh2[0][3] - sh1[0][1] * sh2[4][3]);
sh3[2][5] = kSqrt08_05 * sh1[1][1] * sh2[1][4] + kSqrt06_05 * sh1[0][1] * sh2[2][4] + -kSqrt01_10 * (sh1[2][1] * sh2[0][4] - sh1[0][1] * sh2[4][4]);
sh3[2][6] = kSqrt04_15 * (sh1[1][2] * sh2[1][4] - sh1[1][0] * sh2[1][0]) + kSqrt01_05 * (sh1[0][2] * sh2[2][4] - sh1[0][0] * sh2[2][0]) + -sqrt(1.0 / 60.0) * ((sh1[2][2] * sh2[0][4] - sh1[2][0] * sh2[0][0]) - (sh1[0][2] * sh2[4][4] - sh1[0][0] * sh2[4][0]));
coeffs[dstIdx++] = Dot7(srcIdx, coeffsIn, sh3[2]);
sh3[3][0] = kSqrt03_10 * (sh1[1][2] * sh2[2][0] + sh1[1][0] * sh2[2][4]) + -kSqrt01_10 * ((sh1[2][2] * sh2[3][0] + sh1[2][0] * sh2[3][4]) + (sh1[0][2] * sh2[1][0] + sh1[0][0] * sh2[1][4]));
sh3[3][1] = kSqrt09_05 * sh1[1][1] * sh2[2][0] + -kSqrt03_05 * (sh1[2][1] * sh2[3][0] + sh1[0][1] * sh2[1][0]);
sh3[3][2] = kSqrt09_08 * sh1[1][1] * sh2[2][1] + -kSqrt03_08 * (sh1[2][1] * sh2[3][1] + sh1[0][1] * sh2[1][1]);
sh3[3][3] = sh1[1][1] * sh2[2][2] + -kSqrt01_03 * (sh1[2][1] * sh2[3][2] + sh1[0][1] * sh2[1][2]);
sh3[3][4] = kSqrt09_08 * sh1[1][1] * sh2[2][3] + -kSqrt03_08 * (sh1[2][1] * sh2[3][3] + sh1[0][1] * sh2[1][3]);
sh3[3][5] = kSqrt09_05 * sh1[1][1] * sh2[2][4] + -kSqrt03_05 * (sh1[2][1] * sh2[3][4] + sh1[0][1] * sh2[1][4]);
sh3[3][6] = kSqrt03_10 * (sh1[1][2] * sh2[2][4] - sh1[1][0] * sh2[2][0]) + -kSqrt01_10 * ((sh1[2][2] * sh2[3][4] - sh1[2][0] * sh2[3][0]) + (sh1[0][2] * sh2[1][4] - sh1[0][0] * sh2[1][0]));
coeffs[dstIdx++] = Dot7(srcIdx, coeffsIn, sh3[3]);
sh3[4][0] = kSqrt04_15 * (sh1[1][2] * sh2[3][0] + sh1[1][0] * sh2[3][4]) + kSqrt01_05 * (sh1[2][2] * sh2[2][0] + sh1[2][0] * sh2[2][4]) + -sqrt(1.0 / 60.0) * ((sh1[2][2] * sh2[4][0] + sh1[2][0] * sh2[4][4]) + (sh1[0][2] * sh2[0][0] + sh1[0][0] * sh2[0][4]));
sh3[4][1] = kSqrt08_05 * sh1[1][1] * sh2[3][0] + kSqrt06_05 * sh1[2][1] * sh2[2][0] + -kSqrt01_10 * (sh1[2][1] * sh2[4][0] + sh1[0][1] * sh2[0][0]);
sh3[4][2] = sh1[1][1] * sh2[3][1] + kSqrt03_04 * sh1[2][1] * sh2[2][1] + -kSqrt01_16 * (sh1[2][1] * sh2[4][1] + sh1[0][1] * sh2[0][1]);
sh3[4][3] = kSqrt08_09 * sh1[1][1] * sh2[3][2] + kSqrt02_03 * sh1[2][1] * sh2[2][2] + -kSqrt01_18 * (sh1[2][1] * sh2[4][2] + sh1[0][1] * sh2[0][2]);
sh3[4][4] = sh1[1][1] * sh2[3][3] + kSqrt03_04 * sh1[2][1] * sh2[2][3] + -kSqrt01_16 * (sh1[2][1] * sh2[4][3] + sh1[0][1] * sh2[0][3]);
sh3[4][5] = kSqrt08_05 * sh1[1][1] * sh2[3][4] + kSqrt06_05 * sh1[2][1] * sh2[2][4] + -kSqrt01_10 * (sh1[2][1] * sh2[4][4] + sh1[0][1] * sh2[0][4]);
sh3[4][6] = kSqrt04_15 * (sh1[1][2] * sh2[3][4] - sh1[1][0] * sh2[3][0]) + kSqrt01_05 * (sh1[2][2] * sh2[2][4] - sh1[2][0] * sh2[2][0]) + -sqrt(1.0 / 60.0) * ((sh1[2][2] * sh2[4][4] - sh1[2][0] * sh2[4][0]) + (sh1[0][2] * sh2[0][4] - sh1[0][0] * sh2[0][0]));
coeffs[dstIdx++] = Dot7(srcIdx, coeffsIn, sh3[4]);
sh3[5][0] = kSqrt01_06 * (sh1[1][2] * sh2[4][0] + sh1[1][0] * sh2[4][4]) + kSqrt01_06 * ((sh1[2][2] * sh2[3][0] + sh1[2][0] * sh2[3][4]) - (sh1[0][2] * sh2[1][0] + sh1[0][0] * sh2[1][4]));
sh3[5][1] = sh1[1][1] * sh2[4][0] + (sh1[2][1] * sh2[3][0] - sh1[0][1] * sh2[1][0]);
sh3[5][2] = kSqrt05_08 * sh1[1][1] * sh2[4][1] + kSqrt05_08 * (sh1[2][1] * sh2[3][1] - sh1[0][1] * sh2[1][1]);
sh3[5][3] = kSqrt05_09 * sh1[1][1] * sh2[4][2] + kSqrt05_09 * (sh1[2][1] * sh2[3][2] - sh1[0][1] * sh2[1][2]);
sh3[5][4] = kSqrt05_08 * sh1[1][1] * sh2[4][3] + kSqrt05_08 * (sh1[2][1] * sh2[3][3] - sh1[0][1] * sh2[1][3]);
sh3[5][5] = sh1[1][1] * sh2[4][4] + (sh1[2][1] * sh2[3][4] - sh1[0][1] * sh2[1][4]);
sh3[5][6] = kSqrt01_06 * (sh1[1][2] * sh2[4][4] - sh1[1][0] * sh2[4][0]) + kSqrt01_06 * ((sh1[2][2] * sh2[3][4] - sh1[2][0] * sh2[3][0]) - (sh1[0][2] * sh2[1][4] - sh1[0][0] * sh2[1][0]));
coeffs[dstIdx++] = Dot7(srcIdx, coeffsIn, sh3[5]);
sh3[6][0] = kSqrt01_04 * ((sh1[2][2] * sh2[4][0] + sh1[2][0] * sh2[4][4]) - (sh1[0][2] * sh2[0][0] + sh1[0][0] * sh2[0][4]));
sh3[6][1] = kSqrt03_02 * (sh1[2][1] * sh2[4][0] - sh1[0][1] * sh2[0][0]);
sh3[6][2] = kSqrt15_16 * (sh1[2][1] * sh2[4][1] - sh1[0][1] * sh2[0][1]);
sh3[6][3] = kSqrt05_06 * (sh1[2][1] * sh2[4][2] - sh1[0][1] * sh2[0][2]);
sh3[6][4] = kSqrt15_16 * (sh1[2][1] * sh2[4][3] - sh1[0][1] * sh2[0][3]);
sh3[6][5] = kSqrt03_02 * (sh1[2][1] * sh2[4][4] - sh1[0][1] * sh2[0][4]);
sh3[6][6] = kSqrt01_04 * ((sh1[2][2] * sh2[4][4] - sh1[2][0] * sh2[4][0]) - (sh1[0][2] * sh2[0][4] - sh1[0][0] * sh2[0][0]));
coeffs[dstIdx++] = Dot7(srcIdx, coeffsIn, sh3[6]);
}
#endif // SPHERICAL_HARMONICS_HLSL

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: 0e45617a7c5ba4b4eb55e897761dcb31
ShaderIncludeImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,757 @@
// SPDX-License-Identifier: MIT
#define GROUP_SIZE 1024
#pragma kernel CSSetIndices
#pragma kernel CSCalcDistances
#pragma kernel CSCalcViewData
#pragma kernel CSUpdateEditData
#pragma kernel CSInitEditData
#pragma kernel CSClearBuffer
#pragma kernel CSInvertSelection
#pragma kernel CSSelectAll
#pragma kernel CSOrBuffers
#pragma kernel CSSelectionUpdate
#pragma kernel CSTranslateSelection
#pragma kernel CSRotateSelection
#pragma kernel CSScaleSelection
#pragma kernel CSExportData
#pragma kernel CSCopySplats
// DeviceRadixSort
#pragma multi_compile __ KEY_UINT KEY_INT KEY_FLOAT
#pragma multi_compile __ PAYLOAD_UINT PAYLOAD_INT PAYLOAD_FLOAT
#pragma multi_compile __ SHOULD_ASCEND
#pragma multi_compile __ SORT_PAIRS
#pragma multi_compile __ VULKAN
#pragma kernel InitDeviceRadixSort
#pragma kernel Upsweep
#pragma kernel Scan
#pragma kernel Downsweep
// GPU sorting needs wave ops
#pragma require wavebasic
#pragma require waveballot
#pragma use_dxc
#include "DeviceRadixSort.hlsl"
#include "GaussianSplatting.hlsl"
#include "UnityCG.cginc"
float4x4 _MatrixObjectToWorld;
float4x4 _MatrixWorldToObject;
float4x4 _MatrixMV;
float4 _VecScreenParams;
float4 _VecWorldSpaceCameraPos;
int _SelectionMode;
RWStructuredBuffer<uint> _SplatSortDistances;
RWStructuredBuffer<uint> _SplatSortKeys;
uint _SplatCount;
// radix sort etc. friendly, see http://stereopsis.com/radix.html
uint FloatToSortableUint(float f)
{
uint fu = asuint(f);
uint mask = -((int)(fu >> 31)) | 0x80000000;
return fu ^ mask;
}
[numthreads(GROUP_SIZE,1,1)]
void CSSetIndices (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
_SplatSortKeys[idx] = idx;
}
[numthreads(GROUP_SIZE,1,1)]
void CSCalcDistances (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
uint origIdx = _SplatSortKeys[idx];
float3 pos = LoadSplatPos(origIdx);
pos = mul(_MatrixMV, float4(pos.xyz, 1)).xyz;
_SplatSortDistances[idx] = FloatToSortableUint(pos.z);
}
RWStructuredBuffer<SplatViewData> _SplatViewData;
float _SplatScale;
float _SplatOpacityScale;
uint _SHOrder;
uint _SHOnly;
uint _SplatCutoutsCount;
#define SPLAT_CUTOUT_TYPE_ELLIPSOID 0
#define SPLAT_CUTOUT_TYPE_BOX 1
struct GaussianCutoutShaderData // match GaussianCutout.ShaderData in C#
{
float4x4 mat;
uint typeAndFlags;
};
StructuredBuffer<GaussianCutoutShaderData> _SplatCutouts;
RWByteAddressBuffer _SplatSelectedBits;
ByteAddressBuffer _SplatDeletedBits;
uint _SplatBitsValid;
void DecomposeCovariance(float3 cov2d, out float2 v1, out float2 v2)
{
#if 0 // does not quite give the correct results?
// https://jsfiddle.net/mattrossman/ehxmtgw6/
// References:
// - https://www.youtube.com/watch?v=e50Bj7jn9IQ
// - https://en.wikipedia.org/wiki/Eigenvalue_algorithm#2%C3%972_matrices
// - https://people.math.harvard.edu/~knill/teaching/math21b2004/exhibits/2dmatrices/index.html
float a = cov2d.x;
float b = cov2d.y;
float d = cov2d.z;
float det = a * d - b * b; // matrix is symmetric, so "c" is same as "b"
float trace = a + d;
float mean = 0.5 * trace;
float dist = sqrt(mean * mean - det);
float lambda1 = mean + dist; // 1st eigenvalue
float lambda2 = mean - dist; // 2nd eigenvalue
if (b == 0) {
// https://twitter.com/the_ross_man/status/1706342719776551360
if (a > d) v1 = float2(1, 0);
else v1 = float2(0, 1);
} else
v1 = normalize(float2(b, d - lambda2));
v1.y = -v1.y;
// The 2nd eigenvector is just a 90 degree rotation of the first since Gaussian axes are orthogonal
v2 = float2(v1.y, -v1.x);
// scaling components
v1 *= sqrt(lambda1);
v2 *= sqrt(lambda2);
float radius = 1.5;
v1 *= radius;
v2 *= radius;
#else
// same as in antimatter15/splat
float diag1 = cov2d.x, diag2 = cov2d.z, offDiag = cov2d.y;
float mid = 0.5f * (diag1 + diag2);
float radius = length(float2((diag1 - diag2) / 2.0, offDiag));
float lambda1 = mid + radius;
float lambda2 = max(mid - radius, 0.1);
float2 diagVec = normalize(float2(offDiag, lambda1 - diag1));
diagVec.y = -diagVec.y;
float maxSize = 4096.0;
v1 = min(sqrt(2.0 * lambda1), maxSize) * diagVec;
v2 = min(sqrt(2.0 * lambda2), maxSize) * float2(diagVec.y, -diagVec.x);
#endif
}
bool IsSplatCut(float3 pos)
{
bool finalCut = false;
for (uint i = 0; i < _SplatCutoutsCount; ++i)
{
GaussianCutoutShaderData cutData = _SplatCutouts[i];
uint type = cutData.typeAndFlags & 0xFF;
if (type == 0xFF) // invalid/null cutout, ignore
continue;
bool invert = (cutData.typeAndFlags & 0xFF00) != 0;
float3 cutoutPos = mul(cutData.mat, float4(pos, 1)).xyz;
if (type == SPLAT_CUTOUT_TYPE_ELLIPSOID)
{
if (dot(cutoutPos, cutoutPos) <= 1) return invert;
}
if (type == SPLAT_CUTOUT_TYPE_BOX)
{
if (all(abs(cutoutPos) <= 1)) return invert;
}
finalCut |= !invert;
}
return finalCut;
}
[numthreads(GROUP_SIZE,1,1)]
void CSCalcViewData (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
SplatData splat = LoadSplatData(idx);
SplatViewData view = (SplatViewData)0;
float3 centerWorldPos = mul(_MatrixObjectToWorld, float4(splat.pos,1)).xyz;
float4 centerClipPos = mul(UNITY_MATRIX_VP, float4(centerWorldPos, 1));
half opacityScale = _SplatOpacityScale;
float splatScale = _SplatScale;
// deleted?
if (_SplatBitsValid)
{
uint wordIdx = idx / 32;
uint bitIdx = idx & 31;
uint wordVal = _SplatDeletedBits.Load(wordIdx * 4);
if (wordVal & (1 << bitIdx))
{
centerClipPos.w = 0;
}
}
// cutouts
if (IsSplatCut(splat.pos))
{
centerClipPos.w = 0;
}
view.pos = centerClipPos;
bool behindCam = centerClipPos.w <= 0;
if (!behindCam)
{
float4 boxRot = splat.rot;
float3 boxSize = splat.scale;
float3x3 splatRotScaleMat = CalcMatrixFromRotationScale(boxRot, boxSize);
float3 cov3d0, cov3d1;
CalcCovariance3D(splatRotScaleMat, cov3d0, cov3d1);
float splatScale2 = splatScale * splatScale;
cov3d0 *= splatScale2;
cov3d1 *= splatScale2;
float3 cov2d = CalcCovariance2D(splat.pos, cov3d0, cov3d1, _MatrixMV, UNITY_MATRIX_P, _VecScreenParams);
DecomposeCovariance(cov2d, view.axis1, view.axis2);
float3 worldViewDir = _VecWorldSpaceCameraPos.xyz - centerWorldPos;
float3 objViewDir = mul((float3x3)_MatrixWorldToObject, worldViewDir);
objViewDir = normalize(objViewDir);
half4 col;
col.rgb = ShadeSH(splat.sh, objViewDir, _SHOrder, _SHOnly != 0);
col.a = min(splat.opacity * opacityScale, 65000);
view.color.x = (f32tof16(col.r) << 16) | f32tof16(col.g);
view.color.y = (f32tof16(col.b) << 16) | f32tof16(col.a);
}
_SplatViewData[idx] = view;
}
RWByteAddressBuffer _DstBuffer;
ByteAddressBuffer _SrcBuffer;
uint _BufferSize;
uint2 GetSplatIndicesFromWord(uint idx)
{
uint idxStart = idx * 32;
uint idxEnd = min(idxStart + 32, _SplatCount);
return uint2(idxStart, idxEnd);
}
[numthreads(GROUP_SIZE,1,1)]
void CSUpdateEditData (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _BufferSize)
return;
uint valSel = _SplatSelectedBits.Load(idx * 4);
uint valDel = _SplatDeletedBits.Load(idx * 4);
valSel &= ~valDel; // don't count deleted splats as selected
uint2 splatIndices = GetSplatIndicesFromWord(idx);
// update selection bounds
float3 bmin = 1.0e38;
float3 bmax = -1.0e38;
uint mask = 1;
uint valCut = 0;
for (uint sidx = splatIndices.x; sidx < splatIndices.y; ++sidx, mask <<= 1)
{
float3 spos = LoadSplatPos(sidx);
// don't count cut splats as selected
if (IsSplatCut(spos))
{
valSel &= ~mask;
valCut |= mask;
}
if (valSel & mask)
{
bmin = min(bmin, spos);
bmax = max(bmax, spos);
}
}
valCut &= ~valDel; // don't count deleted splats as cut
if (valSel != 0)
{
_DstBuffer.InterlockedMin(12, FloatToSortableUint(bmin.x));
_DstBuffer.InterlockedMin(16, FloatToSortableUint(bmin.y));
_DstBuffer.InterlockedMin(20, FloatToSortableUint(bmin.z));
_DstBuffer.InterlockedMax(24, FloatToSortableUint(bmax.x));
_DstBuffer.InterlockedMax(28, FloatToSortableUint(bmax.y));
_DstBuffer.InterlockedMax(32, FloatToSortableUint(bmax.z));
}
uint sumSel = countbits(valSel);
uint sumDel = countbits(valDel);
uint sumCut = countbits(valCut);
_DstBuffer.InterlockedAdd(0, sumSel);
_DstBuffer.InterlockedAdd(4, sumDel);
_DstBuffer.InterlockedAdd(8, sumCut);
}
[numthreads(1,1,1)]
void CSInitEditData (uint3 id : SV_DispatchThreadID)
{
_DstBuffer.Store3(0, uint3(0,0,0)); // selected, deleted, cut counts
uint initMin = FloatToSortableUint(1.0e38);
uint initMax = FloatToSortableUint(-1.0e38);
_DstBuffer.Store3(12, uint3(initMin, initMin, initMin));
_DstBuffer.Store3(24, uint3(initMax, initMax, initMax));
}
[numthreads(GROUP_SIZE,1,1)]
void CSClearBuffer (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _BufferSize)
return;
_DstBuffer.Store(idx * 4, 0);
}
[numthreads(GROUP_SIZE,1,1)]
void CSInvertSelection (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _BufferSize)
return;
uint v = _DstBuffer.Load(idx * 4);
v = ~v;
// do not select splats that are cut
uint2 splatIndices = GetSplatIndicesFromWord(idx);
uint mask = 1;
for (uint sidx = splatIndices.x; sidx < splatIndices.y; ++sidx, mask <<= 1)
{
float3 spos = LoadSplatPos(sidx);
if (IsSplatCut(spos))
v &= ~mask;
}
_DstBuffer.Store(idx * 4, v);
}
[numthreads(GROUP_SIZE,1,1)]
void CSSelectAll (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _BufferSize)
return;
uint v = ~0;
// do not select splats that are cut
uint2 splatIndices = GetSplatIndicesFromWord(idx);
uint mask = 1;
for (uint sidx = splatIndices.x; sidx < splatIndices.y; ++sidx, mask <<= 1)
{
float3 spos = LoadSplatPos(sidx);
if (IsSplatCut(spos))
v &= ~mask;
}
_DstBuffer.Store(idx * 4, v);
}
[numthreads(GROUP_SIZE,1,1)]
void CSOrBuffers (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _BufferSize)
return;
uint a = _SrcBuffer.Load(idx * 4);
uint b = _DstBuffer.Load(idx * 4);
_DstBuffer.Store(idx * 4, a | b);
}
float4 _SelectionRect;
[numthreads(GROUP_SIZE,1,1)]
void CSSelectionUpdate (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
float3 pos = LoadSplatPos(idx);
if (IsSplatCut(pos))
return;
float3 centerWorldPos = mul(_MatrixObjectToWorld, float4(pos,1)).xyz;
float4 centerClipPos = mul(UNITY_MATRIX_VP, float4(centerWorldPos, 1));
bool behindCam = centerClipPos.w <= 0;
if (behindCam)
return;
float2 pixelPos = (centerClipPos.xy / centerClipPos.w * float2(0.5, -0.5) + 0.5) * _VecScreenParams.xy;
if (pixelPos.x < _SelectionRect.x || pixelPos.x > _SelectionRect.z ||
pixelPos.y < _SelectionRect.y || pixelPos.y > _SelectionRect.w)
{
return;
}
uint wordIdx = idx / 32;
uint bitIdx = idx & 31;
if (_SelectionMode)
_SplatSelectedBits.InterlockedOr(wordIdx * 4, 1u << bitIdx); // +
else
_SplatSelectedBits.InterlockedAnd(wordIdx * 4, ~(1u << bitIdx)); // -
}
float3 _SelectionDelta;
bool IsSplatSelected(uint idx)
{
uint wordIdx = idx / 32;
uint bitIdx = idx & 31;
uint selVal = _SplatSelectedBits.Load(wordIdx * 4);
return (selVal & (1 << bitIdx)) != 0;
}
[numthreads(GROUP_SIZE,1,1)]
void CSTranslateSelection (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
if (!IsSplatSelected(idx))
return;
uint fmt = _SplatFormat & 0xFF;
if (_SplatChunkCount == 0 && fmt == VECTOR_FMT_32F)
{
uint stride = 12;
float3 pos = asfloat(_SplatPos.Load3(idx * stride));
pos += _SelectionDelta;
_SplatPos.Store3(idx * stride, asuint(pos));
}
}
float3 _SelectionCenter;
float4 _SelectionDeltaRot;
ByteAddressBuffer _SplatPosMouseDown;
ByteAddressBuffer _SplatOtherMouseDown;
[numthreads(GROUP_SIZE,1,1)]
void CSRotateSelection (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
if (!IsSplatSelected(idx))
return;
uint posFmt = _SplatFormat & 0xFF;
if (_SplatChunkCount == 0 && posFmt == VECTOR_FMT_32F)
{
uint posStride = 12;
float3 pos = asfloat(_SplatPosMouseDown.Load3(idx * posStride));
pos -= _SelectionCenter;
pos = mul(_MatrixObjectToWorld, float4(pos,1)).xyz;
pos = QuatRotateVector(pos, _SelectionDeltaRot);
pos = mul(_MatrixWorldToObject, float4(pos,1)).xyz;
pos += _SelectionCenter;
_SplatPos.Store3(idx * posStride, asuint(pos));
}
uint scaleFmt = (_SplatFormat >> 8) & 0xFF;
uint shFormat = (_SplatFormat >> 16) & 0xFF;
if (_SplatChunkCount == 0 && scaleFmt == VECTOR_FMT_32F && shFormat == VECTOR_FMT_32F)
{
uint otherStride = 4 + 12;
uint rotVal = _SplatOtherMouseDown.Load(idx * otherStride);
float4 rot = DecodeRotation(DecodePacked_10_10_10_2(rotVal));
//@TODO: correct rotation
rot = QuatMul(rot, _SelectionDeltaRot);
rotVal = EncodeQuatToNorm10(PackSmallest3Rotation(rot));
_SplatOther.Store(idx * otherStride, rotVal);
}
//@TODO: rotate SHs
}
//@TODO: maybe scale the splat scale itself too?
[numthreads(GROUP_SIZE,1,1)]
void CSScaleSelection (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
if (!IsSplatSelected(idx))
return;
uint fmt = _SplatFormat & 0xFF;
if (_SplatChunkCount == 0 && fmt == VECTOR_FMT_32F)
{
uint stride = 12;
float3 pos = asfloat(_SplatPosMouseDown.Load3(idx * stride));
pos -= _SelectionCenter;
pos = mul(_MatrixObjectToWorld, float4(pos,1)).xyz;
pos *= _SelectionDelta;
pos = mul(_MatrixWorldToObject, float4(pos,1)).xyz;
pos += _SelectionCenter;
_SplatPos.Store3(idx * stride, asuint(pos));
}
}
struct ExportSplatData
{
float3 pos;
float3 nor;
float3 dc0;
float4 shR14; float4 shR58; float4 shR9C; float3 shRDF;
float4 shG14; float4 shG58; float4 shG9C; float3 shGDF;
float4 shB14; float4 shB58; float4 shB9C; float3 shBDF;
float opacity;
float3 scale;
float4 rot;
};
RWStructuredBuffer<ExportSplatData> _ExportBuffer;
float3 ColorToSH0(float3 col)
{
return (col - 0.5) / 0.2820948;
}
float InvSigmoid(float v)
{
return log(v / max(1 - v, 1.0e-6));
}
// SH rotation
#include "SphericalHarmonics.hlsl"
void RotateSH(inout SplatSHData sh, float3x3 rot)
{
float3 shin[16];
float3 shout[16];
shin[0] = sh.col;
shin[1] = sh.sh1;
shin[2] = sh.sh2;
shin[3] = sh.sh3;
shin[4] = sh.sh4;
shin[5] = sh.sh5;
shin[6] = sh.sh6;
shin[7] = sh.sh7;
shin[8] = sh.sh8;
shin[9] = sh.sh9;
shin[10] = sh.sh10;
shin[11] = sh.sh11;
shin[12] = sh.sh12;
shin[13] = sh.sh13;
shin[14] = sh.sh14;
shin[15] = sh.sh15;
RotateSH(rot, 4, shin, shout);
sh.col = shout[0];
sh.sh1 = shout[1];
sh.sh2 = shout[2];
sh.sh3 = shout[3];
sh.sh4 = shout[4];
sh.sh5 = shout[5];
sh.sh6 = shout[6];
sh.sh7 = shout[7];
sh.sh8 = shout[8];
sh.sh9 = shout[9];
sh.sh10 = shout[10];
sh.sh11 = shout[11];
sh.sh12 = shout[12];
sh.sh13 = shout[13];
sh.sh14 = shout[14];
sh.sh15 = shout[15];
}
float3x3 CalcSHRotMatrix(float4x4 objToWorld)
{
float3x3 m = (float3x3)objToWorld;
float sx = length(float3(m[0][0], m[0][1], m[0][2]));
float sy = length(float3(m[1][0], m[1][1], m[1][2]));
float sz = length(float3(m[2][0], m[2][1], m[2][2]));
float invSX = 1.0 / sx;
float invSY = 1.0 / sy;
float invSZ = 1.0 / sz;
m[0][0] *= invSX;
m[0][1] *= invSX;
m[0][2] *= invSX;
m[1][0] *= invSY;
m[1][1] *= invSY;
m[1][2] *= invSY;
m[2][0] *= invSZ;
m[2][1] *= invSZ;
m[2][2] *= invSZ;
return m;
}
float4 _ExportTransformRotation;
float3 _ExportTransformScale;
uint _ExportTransformFlags;
[numthreads(GROUP_SIZE,1,1)]
void CSExportData (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _SplatCount)
return;
SplatData src = LoadSplatData(idx);
bool isCut = IsSplatCut(src.pos);
// transform splat by matrix, if needed
if (_ExportTransformFlags != 0)
{
src.pos = mul(_MatrixObjectToWorld, float4(src.pos,1)).xyz;
// note: this only handles axis flips from scale, not any arbitrary scaling
if (_ExportTransformScale.x < 0)
src.rot.yz = -src.rot.yz;
if (_ExportTransformScale.y < 0)
src.rot.xz = -src.rot.xz;
if (_ExportTransformScale.z < 0)
src.rot.xy = -src.rot.xy;
src.rot = QuatMul(_ExportTransformRotation, src.rot);
src.scale *= abs(_ExportTransformScale);
float3x3 shRot = CalcSHRotMatrix(_MatrixObjectToWorld);
RotateSH(src.sh, shRot);
}
ExportSplatData dst;
dst.pos = src.pos;
dst.nor = 0;
dst.dc0 = ColorToSH0(src.sh.col);
dst.shR14 = float4(src.sh.sh1.r, src.sh.sh2.r, src.sh.sh3.r, src.sh.sh4.r);
dst.shR58 = float4(src.sh.sh5.r, src.sh.sh6.r, src.sh.sh7.r, src.sh.sh8.r);
dst.shR9C = float4(src.sh.sh9.r, src.sh.sh10.r, src.sh.sh11.r, src.sh.sh12.r);
dst.shRDF = float3(src.sh.sh13.r, src.sh.sh14.r, src.sh.sh15.r);
dst.shG14 = float4(src.sh.sh1.g, src.sh.sh2.g, src.sh.sh3.g, src.sh.sh4.g);
dst.shG58 = float4(src.sh.sh5.g, src.sh.sh6.g, src.sh.sh7.g, src.sh.sh8.g);
dst.shG9C = float4(src.sh.sh9.g, src.sh.sh10.g, src.sh.sh11.g, src.sh.sh12.g);
dst.shGDF = float3(src.sh.sh13.g, src.sh.sh14.g, src.sh.sh15.g);
dst.shB14 = float4(src.sh.sh1.b, src.sh.sh2.b, src.sh.sh3.b, src.sh.sh4.b);
dst.shB58 = float4(src.sh.sh5.b, src.sh.sh6.b, src.sh.sh7.b, src.sh.sh8.b);
dst.shB9C = float4(src.sh.sh9.b, src.sh.sh10.b, src.sh.sh11.b, src.sh.sh12.b);
dst.shBDF = float3(src.sh.sh13.b, src.sh.sh14.b, src.sh.sh15.b);
dst.opacity = InvSigmoid(src.opacity);
dst.scale = log(src.scale);
dst.rot = src.rot.wxyz;
if (isCut)
dst.nor = 1; // mark as skipped for export
_ExportBuffer[idx] = dst;
}
RWByteAddressBuffer _CopyDstPos;
RWByteAddressBuffer _CopyDstOther;
RWByteAddressBuffer _CopyDstSH;
RWByteAddressBuffer _CopyDstEditDeleted;
RWTexture2D<float4> _CopyDstColor;
uint _CopyDstSize, _CopySrcStartIndex, _CopyDstStartIndex, _CopyCount;
float4x4 _CopyTransformMatrix;
float4 _CopyTransformRotation;
float3 _CopyTransformScale;
[numthreads(GROUP_SIZE,1,1)]
void CSCopySplats (uint3 id : SV_DispatchThreadID)
{
uint idx = id.x;
if (idx >= _CopyCount)
return;
uint srcIdx = _CopySrcStartIndex + idx;
uint dstIdx = _CopyDstStartIndex + idx;
if (srcIdx >= _SplatCount || dstIdx >= _CopyDstSize)
return;
SplatData src = LoadSplatData(idx);
// transform the splat
src.pos = mul(_CopyTransformMatrix, float4(src.pos,1)).xyz;
// note: this only handles axis flips from scale, not any arbitrary scaling
if (_CopyTransformScale.x < 0)
src.rot.yz = -src.rot.yz;
if (_CopyTransformScale.y < 0)
src.rot.xz = -src.rot.xz;
if (_CopyTransformScale.z < 0)
src.rot.xy = -src.rot.xy;
src.rot = QuatMul(_CopyTransformRotation, src.rot);
src.scale *= abs(_CopyTransformScale);
float3x3 shRot = CalcSHRotMatrix(_CopyTransformMatrix);
RotateSH(src.sh, shRot);
// output data into destination:
// pos
uint posStride = 12;
_CopyDstPos.Store3(dstIdx * posStride, asuint(src.pos));
// rot + scale
uint otherStride = 4 + 12;
uint rotVal = EncodeQuatToNorm10(PackSmallest3Rotation(src.rot));
_CopyDstOther.Store4(dstIdx * otherStride, uint4(
rotVal,
asuint(src.scale.x),
asuint(src.scale.y),
asuint(src.scale.z)));
// color
uint3 pixelIndex = SplatIndexToPixelIndex(dstIdx);
_CopyDstColor[pixelIndex.xy] = float4(src.sh.col, src.opacity);
// SH
uint shStride = 192; // 15*3 fp32, rounded up to multiple of 16
uint shOffset = dstIdx * shStride;
_CopyDstSH.Store3(shOffset + 12 * 0, asuint(src.sh.sh1));
_CopyDstSH.Store3(shOffset + 12 * 1, asuint(src.sh.sh2));
_CopyDstSH.Store3(shOffset + 12 * 2, asuint(src.sh.sh3));
_CopyDstSH.Store3(shOffset + 12 * 3, asuint(src.sh.sh4));
_CopyDstSH.Store3(shOffset + 12 * 4, asuint(src.sh.sh5));
_CopyDstSH.Store3(shOffset + 12 * 5, asuint(src.sh.sh6));
_CopyDstSH.Store3(shOffset + 12 * 6, asuint(src.sh.sh7));
_CopyDstSH.Store3(shOffset + 12 * 7, asuint(src.sh.sh8));
_CopyDstSH.Store3(shOffset + 12 * 8, asuint(src.sh.sh9));
_CopyDstSH.Store3(shOffset + 12 * 9, asuint(src.sh.sh10));
_CopyDstSH.Store3(shOffset + 12 * 10, asuint(src.sh.sh11));
_CopyDstSH.Store3(shOffset + 12 * 11, asuint(src.sh.sh12));
_CopyDstSH.Store3(shOffset + 12 * 12, asuint(src.sh.sh13));
_CopyDstSH.Store3(shOffset + 12 * 13, asuint(src.sh.sh14));
_CopyDstSH.Store3(shOffset + 12 * 14, asuint(src.sh.sh15));
// deleted bits
uint srcWordIdx = srcIdx / 32;
uint srcBitIdx = srcIdx & 31;
if (_SplatDeletedBits.Load(srcWordIdx * 4) & (1u << srcBitIdx))
{
uint dstWordIdx = dstIdx / 32;
uint dstBitIdx = dstIdx & 31;
_CopyDstEditDeleted.InterlockedOr(dstWordIdx * 4, 1u << dstBitIdx);
}
}

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: ec84f78b836bd4f96a105d6b804f08bd
ComputeShaderImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

View File

@@ -0,0 +1,13 @@
{
"author": "Aras Pranckevicius",
"dependencies": { "com.unity.burst": "1.8.8", "com.unity.collections": "2.1.4", "com.unity.mathematics": "1.2.6" },
"description": "3D Gaussian Splatting rendering and tools",
"displayName": "Gaussian Splatting",
"keywords": [ "unity" ],
"license": "MIT",
"name": "org.nesnausk.gaussian-splatting",
"repository": "github:aras-p/UnityGaussianSplatting",
"unity": "2022.3",
"unityRelease": "7f1",
"version": "1.0.0"
}

View File

@@ -0,0 +1,7 @@
fileFormatVersion: 2
guid: 110db01d54641354399a53474743e47b
PackageManifestImporter:
externalObjects: {}
userData:
assetBundleName:
assetBundleVariant:

BIN
MVS/3DGS-Unity/room.ply LFS Normal file

Binary file not shown.

Binary file not shown.

After

Width:  |  Height:  |  Size: 1.8 MiB

View File

@@ -0,0 +1,16 @@
# Blender MTL File: 'None'
# Material Count: 1
newmtl Scene_-_Root
Ns 225.000000
Ka 1.000000 1.000000 1.000000
Kd 0.800000 0.800000 0.800000
Ks 0.500000 0.500000 0.500000
Ke 0.0 0.0 0.0
Ni 1.450000
d 1.000000
illum 2
map_Kd diffuse.jpg
map_Bump normal.png
map_Ks specular.jpg

Binary file not shown.

After

Width:  |  Height:  |  Size: 5.8 MiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 14 MiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 4.2 MiB

View File

@@ -0,0 +1,3 @@
Model by Berk Gedik, from: https://sketchfab.com/3d-models/survival-guitar-backpack-low-poly-799f8c4511f84fab8c3f12887f7e6b36
Modified material assignment (Joey de Vries) for easier load in OpenGL model loading chapter, and renamed albedo to diffuse and metallic to specular to match non-PBR lighting setup.

Binary file not shown.

After

Width:  |  Height:  |  Size: 6.4 MiB