basic framework

.gitignore

@@ -0,0 +1,70 @@
+CMakeCache.txt
+CMakeFiles
+CMakeScripts
+Makefile
+cmake_install.cmake
+install_manifest.txt
+# Compiled Object files
+*.slo
+*.lo
+*.o
+*.obj
+
+# Precompiled Headers
+*.gch
+*.pch
+
+# Compiled Dynamic libraries
+*.so
+*.dylib
+*.dll
+
+# Fortran module files
+*.mod
+
+# Compiled Static libraries
+*.lai
+*.la
+*.a
+*.lib
+
+# Executables
+*.exe
+*.out
+*.app
+
+*.i
+*.ii
+*.gpu
+*.ptx
+*.cubin
+*.fatbin
+
+# Qt-es
+
+/.qmake.cache
+/.qmake.stash
+*.pro.user
+*.pro.user.*
+*.qbs.user
+*.qbs.user.*
+*.moc
+moc_*.cpp
+qrc_*.cpp
+ui_*.h
+Makefile*
+*build-*
+
+# QtCreator
+
+*.autosave
+
+# QtCtreator Qml
+*.qmlproject.user
+*.qmlproject.user.*
+
+# QtCtreator CMake
+CMakeLists.txt.user
+
+build/*
+.idea/*

CMakeLists.txt

@@ -0,0 +1,49 @@
+cmake_minimum_required(VERSION 3.3)
+project(SunVolumeRender)
+
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11")
+set(CMAKE_INCLUDE_CURRENT_DIR ON)
+
+# CUDA settings
+find_package(CUDA QUIET REQUIRED)
+list(APPEND CUDA_NVCC_FLAGS --compiler-options -fno-strict-aliasing -use_fast_math -Xptxas -v -maxrregcount=32)
+list(APPEND CUDA_NVCC_FLAGS -gencode arch=compute_30,code=sm_30)
+set(CUDA_INCLUDE_DIRS /usr/local/cuda/include/)
+include_directories(${CUDA_INCLUDE_DIRS})
+
+#Qt
+find_package(Qt4 COMPONENTS QtCore QtGui QtOpenGL REQUIRED)
+include(${QT_USE_FILE})
+set(QT_USE_QTOPENGL TRUE)
+set(CMAKE_AUTOMOC ON)
+set(CMAKE_AUTOUIC ON)
+set(CMAKE_AUTORCC ON)
+set(RESOURCES_FILES qdarkstyle/style.qrc)
+
+#VTK
+set(VTK_DIR /Volumes/Free/SDK/CTK/build/VTK-build)
+find_package(VTK REQUIRED)
+include(${VTK_USE_FILE})
+
+#CTK
+set(CTK_DIR /Volumes/Free/SDK/CTK/build/)
+#set(CTK_DIR /Users/sunwj/Desktop/SDK/CTK/build)
+find_package(CTK REQUIRED)
+include(${CTK_USE_FILE})
+set(CTK_LIBRARIES CTKCore CTKWidgets CTKVisualizationVTKCore CTKVisualizationVTKWidgets)
+
+#GLM
+find_package(GLM REQUIRED)
+include_directories(${GLM_INCLUDE_DIRS})
+
+set(HOST_SOURCES main.cpp
+        gui/mainwindow.cpp
+        gui/canvas.cpp
+        gui/transferfunction.cpp)
+
+set(DEVICE_SOURCES pathtracer.cu)
+
+cuda_compile(DEVICE_OBJS ${DEVICE_SOURCES})
+
+add_executable(SunVolumeRender ${HOST_SOURCES} ${RESOURCES_FILES} ${DEVICE_OBJS})
+target_link_libraries(SunVolumeRender ${QT_LIBRARIES} ${VTK_LIBRARIES} ${CTK_LIBRARIES} ${CUDA_LIBRARIES})

common.h

@@ -0,0 +1,11 @@
+//
+// Created by 孙万捷 on 16/5/19.
+//
+
+#ifndef SUNVOLUMERENDER_COMMON_H
+#define SUNVOLUMERENDER_COMMON_H
+
+#define WIDTH 640
+#define HEIGHT 640
+
+#endif //SUNVOLUMERENDER_COMMON_H

core/cuda_bbox.h

@@ -0,0 +1,56 @@
+//
+// Created by 孙万捷 on 16/5/19.
+//
+
+#ifndef SUNVOLUMERENDER_CUDA_BBOX_H
+#define SUNVOLUMERENDER_CUDA_BBOX_H
+
+#define GLM_FORCE_INLINE
+#include <glm/glm.hpp>
+
+#include <cuda_runtime.h>
+
+#include "cuda_ray.h"
+
+class cudaBBox
+{
+public:
+    __host__ cudaBBox() {}
+    __host__ cudaBBox(const glm::vec3& vmin, const glm::vec3& vmax)
+    {
+        Set(vmin, vmax);
+    }
+
+    __host__ void Set(const glm::vec3& vmin, const glm::vec3& vmax)
+    {
+        this->vmin = vmin;
+        this->vmax = vmax;
+
+        invSize = 1.f / (vmax - vmin);
+    }
+
+    __device__ bool Intersect(const cudaRay& ray, float* tNear, float* tFar) const
+    {
+        auto invDir = 1.f / ray.dir;
+        auto tbot = invDir * (vmin - ray.orig);
+        auto ttop = invDir * (vmax - ray.orig);
+
+        auto tmin = glm::min(tbot, ttop);
+        auto tmax = glm::max(tbot, ttop);
+
+        float largest_tmin = fmaxf(tmin.x, fmaxf(tmin.y, tmin.z));
+        float smallest_tmax = fminf(tmax.x, fminf(tmax.y, tmax.z));
+
+        *tNear = largest_tmin;
+        *tFar = smallest_tmax;
+
+        return smallest_tmax > largest_tmin;
+    }
+
+public:
+    glm::vec3 vmin = glm::vec3(glm::uninitialize);
+    glm::vec3 vmax = glm::vec3(glm::uninitialize);
+    glm::vec3 invSize = glm::vec3(glm::uninitialize);
+};
+
+#endif //SUNVOLUMERENDER_CUDA_BBOX_H

core/cuda_camera.h

@@ -0,0 +1,76 @@
+//
+// Created by 孙万捷 on 16/2/6.
+//
+
+#ifndef SUNPATHTRACER_CAMERA_H
+#define SUNPATHTRACER_CAMERA_H
+
+#include <cuda_runtime.h>
+#include <curand_kernel.h>
+
+#define GLM_FORCE_INLINE
+#include <glm/glm.hpp>
+
+#include "cuda_ray.h"
+
+class cudaCamera
+{
+public:
+    __host__ __device__ cudaCamera() {}
+
+    __host__ __device__ cudaCamera(const glm::vec3& _pos, const glm::vec3& _u, const glm::vec3& _v, const glm::vec3& _w, float fovx = 45.f, unsigned int _imageW = 640, unsigned int _imageH = 480)
+    {
+        Setup(_pos, _u, _v, _w, fovx, _imageW, _imageH);
+    }
+
+    __host__ __device__ cudaCamera(const glm::vec3& _pos, const glm::vec3& target, const glm::vec3& up, float fovx = 45.f, unsigned int _imageW = 640, unsigned int _imageH = 480)
+    {
+        Setup(_pos, target, up, fovx, _imageW, _imageH);
+    }
+
+    __host__ __device__ void Setup(const glm::vec3& _pos, const glm::vec3& _u, const glm::vec3& _v, const glm::vec3& _w, float fovx, unsigned int _imageW, unsigned int _imageH)
+    {
+        pos = _pos;
+        u = _u;
+        v = _v;
+        w = _w;
+        imageW = _imageW;
+        imageH = _imageH;
+        aspectRatio = (float)imageW / (float)imageH;
+        tanFovxOverTwo = tanf(fovx * 0.5f * M_PI / 180.f);
+    }
+
+    __host__ __device__ void Setup(const glm::vec3& _pos, const glm::vec3& target, const glm::vec3& up, float fovx, unsigned int _imageW, unsigned int _imageH)
+    {
+        pos = _pos;
+        w = normalize(pos - target);
+        u = cross(up, w);
+        v = cross(w, u);
+        imageW = _imageW;
+        imageH = _imageH;
+        aspectRatio = (float)imageW / (float)imageH;
+        tanFovxOverTwo = tanf(fovx * 0.5f * M_PI / 180.f);
+    }
+
+    // TODO: depth of field
+    __device__ void GenerateRay(unsigned int x, unsigned int y, curandState& rng, cudaRay* ray) const
+    {
+        float nx = 2.f * ((x + curand_uniform(&rng)) / (imageW - 1.f)) - 1.f;
+        float ny = 2.f * ((y + curand_uniform(&rng)) / (imageH - 1.f)) - 1.f;
+
+        nx = nx * aspectRatio * tanFovxOverTwo;
+        ny = ny * tanFovxOverTwo;
+
+        ray->orig = pos;
+        ray->dir = normalize(nx * u + ny * v - w);
+    }
+
+public:
+    unsigned int imageW, imageH;
+    float aspectRatio;
+    float tanFovxOverTwo;
+    glm::vec3 pos;
+    glm::vec3 u, v, w;
+};
+
+#endif //SUNPATHTRACER_CAMERA_H

core/cuda_onb.h

@@ -0,0 +1,52 @@
+//
+// Created by 孙万捷 on 16/2/27.
+//
+
+#ifndef SUNPATHTRACER_ONB_H
+#define SUNPATHTRACER_ONB_H
+
+#define GLM_FORCE_INLINE
+#include <glm/glm.hpp>
+
+#include <cuda_runtime.h>
+
+class cudaONB
+{
+public:
+    __device__ cudaONB(const glm::vec3& _w)
+    {
+        InitFromW(_w);
+    }
+
+    __device__ cudaONB(const glm::vec3& _v, const glm::vec3& _w)
+    {
+        InitFromVW(_v, _w);
+    }
+
+    __device__ void InitFromW(const glm::vec3& _w)
+    {
+        w = _w;
+        if(fabsf(w.x) > fabsf(w.y))
+        {
+            float invLength = rsqrtf(w.x * w.x + w.z * w.z);
+            v = glm::vec3(-w.z * invLength, 0.f, w.x * invLength);
+        }
+        else
+        {
+            float  invLength = rsqrtf(w.y * w.y + w.z * w.z);
+            v = glm::vec3(0.f, w.z * invLength, -w.y * invLength);
+        }
+        u = cross(v, w);
+    }
+
+    __device__ void InitFromVW(const glm::vec3& _v, const glm::vec3& _w)
+    {
+        w = _w;
+        u = cross(_v, w);
+        v = cross(w, u);
+    }
+public:
+    glm::vec3 u, v, w;
+};
+
+#endif //SUNPATHTRACER_ONB_H

core/cuda_ray.h

@@ -0,0 +1,44 @@
+//
+// Created by 孙万捷 on 16/2/6.
+//
+
+#ifndef SUNPATHTRACER_RAY_H
+#define SUNPATHTRACER_RAY_H
+
+#include <float.h>
+
+#define GLM_FORCE_INLINE
+#include <glm/glm.hpp>
+
+#include <cuda_runtime.h>
+
+class cudaRay
+{
+public:
+    __device__ cudaRay()
+    {
+        tMin = 1e-8;
+        tMax = FLT_MAX;
+    }
+
+    __device__ cudaRay(const glm::vec3& orig, const glm::vec3& dir, float tMin = 1e-8, float tMax = FLT_MAX)
+    {
+        this->orig = orig;
+        this->dir = dir;
+        this->tMin = tMin;
+        this->tMax = tMax;
+    }
+
+    __device__ glm::vec3 PointOnRay(float t) const
+    {
+        return orig + t * dir;
+    }
+
+public:
+    glm::vec3 orig;
+    glm::vec3 dir;
+    mutable float tMin;
+    mutable float tMax;
+};
+
+#endif //SUNPATHTRACER_RAY_H