Add point cloud renderer.

CMakeLists.txt

@@ -2,6 +2,9 @@ cmake_minimum_required(VERSION 3.9.4)
 project(ProtoVoxel VERSION 0.1.0)
 
 find_package(OpenGL REQUIRED)
+find_package(CUDA REQUIRED)
+
+enable_language(CUDA)
 
 add_library(glad ${CMAKE_SOURCE_DIR}/glad/glad.c)
 target_include_directories(glad PRIVATE ${CMAKE_SOURCE_DIR})
@@ -14,6 +17,8 @@ file(GLOB_RECURSE SRC_FILES ${CMAKE_SOURCE_DIR}/src/*.cpp)
 add_executable(ProtoVoxel main.cpp ${SRC_FILES})
 set_property(TARGET ProtoVoxel PROPERTY CXX_STANDARD 17)
 set_property(TARGET ProtoVoxel PROPERTY INTERPROCEDURAL_OPTIMIZATION TRUE)
+set_property(TARGET ProtoVoxel PROPERTY CUDA_SEPARABLE_COMPILATION ON)
+set_property(TARGET ProtoVoxel PROPERTY CUDA_ARCHITECTURES 80)
 
 set(TINYGLTF_HEADER_ONLY ON CACHE INTERNAL "" FORCE)
 set(TINYGLTF_INSTALL OFF CACHE INTERNAL "" FORCE)

CMakeSettings.json

@@ -4,7 +4,7 @@
       "name": "x64-Debug",
       "generator": "Ninja",
       "configurationType": "Debug",
-      "inheritEnvironments": [ "clang_cl_x64_x64" ],
+      "inheritEnvironments": [ "msvc_x64_x64" ],
       "buildRoot": "${projectDir}\\out\\build\\${name}",
       "installRoot": "${projectDir}\\out\\install\\${name}",
       "cmakeCommandArgs": "",

main.cpp

@@ -3,11 +3,9 @@
 #include "graphics/gpubuffer.h"
 #include "graphics/graphicsdevice.h"
 #include "scenegraph/scenebase.h"
-#include "voxel/chunkmanager.h"
-#include "voxel/chunkjobmanager.h"
-#include "voxel/mortoncode.h"
+#include "pointcloud/chunkmanager.h"
+#include "pointcloud/chunkjobmanager.h"
 
-#include "voxel/PerlinNoise.h"
 #include "misc/model.h"
 
 #include "imgui/imgui.h"
@@ -21,7 +19,7 @@
 
 namespace PVC = ProtoVoxel::Core;
 namespace PVG = ProtoVoxel::Graphics;
-namespace PVV = ProtoVoxel::Voxel;
+namespace PPC = ProtoVoxel::PointCloud;
 namespace PVSG = ProtoVoxel::SceneGraph;
 namespace PVM = ProtoVoxel::Misc;
 
@@ -32,21 +30,13 @@ class TestScene : public PVSG::SceneBase
     TestScene() {}
     ~TestScene() {}
 
-    PVV::ChunkPalette palette;
-    PVV::ChunkManager manager;
+    PPC::ChunkManager manager;
 
     void Initialize() override
     {
         PVSG::SceneBase::Initialize();
 
-        PVM::Model mdl;
-        PVM::Model::LoadModel("test.glb", mdl);
-
-        palette.Initialize();
-        for (int i = 255; i >= 0; i--)
-            palette.Register(glm::vec4(0, i / 255.0, 0, 1));
-
-        manager.Initialize(palette);
+        manager.Initialize();
     }
 
     void Update(double time) override

shaders/pointcloud/splat.glsl

@@ -0,0 +1,84 @@
+//ComputeShader
+#version 460
+#extension GL_KHR_shader_subgroup_vote : enable
+#extension GL_KHR_shader_subgroup_arithmetic : enable
+//#extension GL_KHR_shader_subgroup : enable
+
+layout (local_size_x = 128, local_size_y = 1) in;
+
+// Values that stay constant for the whole mesh.
+layout(std140, binding = 0) uniform GlobalParams_t {
+        mat4 proj;
+        mat4 view;
+        mat4 vp;
+        mat4 ivp;
+        mat4 prev_proj;
+        mat4 prev_view;
+        mat4 prev_vp;
+        mat4 prev_ivp;
+        vec4 prev_eyePos;
+        vec4 prev_eyeUp;
+        vec4 prev_eyeDir;
+        vec4 eyePos;
+        vec4 eyeUp;
+        vec4 eyeDir;
+        vec4 eyeRight;
+} GlobalParams;
+
+layout(std430, binding = 0) readonly restrict buffer Voxels_t {
+    ivec4 v[];
+} vxls;
+
+//splat smaller voxels as 2x2 pixels with atomicMin
+uniform layout(binding = 0, r32ui) restrict uimage2D pointBuffer;
+uniform layout(binding = 1, rgba8) restrict image2D colorBuffer;
+
+void main(){
+	uint DrawID = gl_GlobalInvocationID.x;
+    //if (DrawID >= in_draws.cnt)
+    //    return;
+    //if (in_draws.cmds[DrawID].count < gl_GlobalInvocationID.y)
+    //    return;
+
+    ivec2 imgSize = imageSize(pointBuffer);
+    uint VertexID = DrawID;//in_draws.cmds[DrawID].baseVertex + gl_GlobalInvocationID.y;
+
+            ivec4 vID = vxls.v[VertexID];
+            float x = float(vID.x) / 10000.0;
+            float y = float(vID.y) / 10000.0;
+            float z = float(vID.z) / 10000.0;
+            vec4 color = unpackUnorm4x8(vID.w);
+
+            vec3 UV;
+            UV.x = x;// + in_draws.cmds[DrawID].pos.x;
+            UV.y = y;// + in_draws.cmds[DrawID].pos.y;
+            UV.z = z;// + in_draws.cmds[DrawID].pos.z;
+
+            vec4 ppos = GlobalParams.vp * vec4(UV, 1);
+            ppos /= ppos.w;
+
+            ppos.xy = ppos.xy * 0.5f + 0.5f;
+            if (ppos.x >= 0.0f && ppos.y >= 0.0f && ppos.x < 1.0f && ppos.y < 1.0f){
+                ivec2 ppos_pxl = ivec2(ppos.xy * imgSize);
+                uint test_val = uint(ppos.z * 16777216);
+
+                //get all threads in workgroup which are writing to the same pixel
+                //and find the one with the largest depth
+                if (subgroupAllEqual(ppos_pxl))
+                {
+                    if( subgroupMax(test_val) == test_val)
+                    {
+                        if(test_val > imageAtomicMax(pointBuffer, ppos_pxl, test_val))
+                        {
+                            imageStore(colorBuffer, ppos_pxl, vec4(0.0f, 1.0f, 0.0f, 1.0f));
+                        }
+                    }
+                }
+                else if(test_val > imageAtomicMax(pointBuffer, ppos_pxl, test_val))
+                {
+                    imageStore(colorBuffer, ppos_pxl, color);
+                }
+            }
+
+            VertexID ++;
+}

src/graphics/texture.cpp

@@ -49,7 +49,7 @@ void PVG::Texture::SetStorage(GLenum target, int levels, int internalFormat,
     SetStorage(target, levels, internalFormat, w, 1, 1);
 }
 
-void ProtoVoxel::Graphics::Texture::Clear(int lv, int internalFormat)
+void ProtoVoxel::Graphics::Texture::Clear(int lv, int internalFormat, int type)
 {
-    glClearTexImage(id, lv, internalFormat, GL_BYTE, nullptr);
+    glClearTexImage(id, lv, internalFormat, type, nullptr);
 }

src/graphics/texture.h

@@ -19,7 +19,7 @@ namespace ProtoVoxel::Graphics {
         void SetStorage(GLenum target, int levels, int internalFormat, size_t w,
             size_t h);
         void SetStorage(GLenum target, int levels, int internalFormat, size_t w);
-        void Clear(int lv, int internalFormat);
+        void Clear(int lv, int internalFormat, int type = GL_BYTE);
 
         inline void SetName(const char* name);
 

src/pointcloud/chunk.cpp

@@ -0,0 +1,32 @@
+#include "chunk.h"
+#include <string.h>
+#include <immintrin.h>
+
+namespace PPC = ProtoVoxel::PointCloud;
+
+PPC::Chunk::Chunk()
+{
+    compressed_len = 0;
+    compressed_data = nullptr;
+}
+
+void PPC::Chunk::Initialize(uint32_t id)
+{
+    this->id = id;
+    compressed_len = 0;
+    compressed_data = nullptr;
+}
+
+void PPC::Chunk::SetPosition(glm::ivec3 &pos)
+{
+    position = pos;
+}
+
+glm::ivec3 &PPC::Chunk::GetPosition()
+{
+    return position;
+}
+
+PPC::Chunk::~Chunk()
+{
+}

src/pointcloud/chunk.h

@@ -0,0 +1,63 @@
+#pragma once
+
+#include "glm/glm.hpp"
+#include <stdint.h>
+#include <vector>
+
+namespace ProtoVoxel::PointCloud
+{
+    enum class ChunkStatus {
+        None,
+        SurfaceUpdatePending,
+    };
+
+    class ChunkUpdater;
+    class ChunkJobManager;
+
+    class Chunk
+    {
+    public:
+        static const int ChunkSide = 32;
+        static const int ChunkLen = ChunkSide * ChunkSide * ChunkSide;
+        static inline uint32_t Encode(uint8_t x, uint8_t y, uint8_t z)
+        {
+            return y | ((uint32_t)z << 5) | ((uint32_t)x << 10);
+        }
+
+        static inline uint32_t EncodeXZ(uint8_t x, uint8_t z)
+        {
+            return ((uint32_t)z) | ((uint32_t)x << 5);
+        }
+
+        static inline uint32_t EncodeXZ_Y(uint32_t xz, uint8_t y)
+        {
+            return y | (xz << 5);
+        }
+
+    private:
+        glm::ivec3 position;
+        glm::ivec3 min_bound;
+        glm::ivec3 max_bound;
+        uint8_t *compressed_data;
+        uint32_t compressed_len;
+
+        ChunkStatus status;
+        uint32_t loopback_cnt;
+        uint32_t mesh_area_ptr;
+        uint32_t mesh_area_len;
+
+        uint32_t id;
+
+        friend class ChunkUpdater;
+        friend class ChunkJobManager;
+
+    public:
+        Chunk();
+        ~Chunk();
+
+        void Initialize(uint32_t id);
+
+        void SetPosition(glm::ivec3 &pos);
+        glm::ivec3 &GetPosition();
+    };
+} // namespace ProtoVoxel::Voxel