add environment

CMakeLists.txt

@@ -18,7 +18,7 @@ set(QT_USE_QTOPENGL TRUE)
 set(CMAKE_AUTOMOC ON)
 set(CMAKE_AUTOUIC ON)
 set(CMAKE_AUTORCC ON)
-set(RESOURCES_FILES qdarkstyle/style.qrc)
+set(RESOURCES_FILES qdarkstyle/style.qrc resources/images.qrc)
 
 #VTK
 set(VTK_DIR /Volumes/Free/SDK/CTK/build/VTK-build)
@@ -40,7 +40,9 @@ set(HOST_SOURCES main.cpp
         gui/mainwindow.cpp
         gui/canvas.cpp
         gui/transferfunction.cpp
-        core/VolumeReader.cpp)
+        gui/qcustomplot.cpp
+        core/VolumeReader.cpp
+        core/lights/lights.cpp)
 
 set(DEVICE_SOURCES pathtracer.cu)
 

core/VolumeReader.cpp

@@ -43,14 +43,6 @@ void VolumeReader::Read(std::string filename)
     imageCast->Update();
     auto imageData = imageCast->GetOutput();
 
-    //auto imageGradientMagnitude = vtkSmartPointer<vtkImageGradientMagnitude>::New();
-    //imageGradientMagnitude->SetDimensionality(3);
-    //imageGradientMagnitude->SetInput(imageCast->GetOutput());
-    //imageGradientMagnitude->Update();
-    //auto magData = imageGradientMagnitude->GetOutput();
-    //auto magRange = magData->GetScalarRange();
-    //maxMagnitude = magRange[2];
-
     auto dataExtent = imageData->GetExtent();
     this->dim = glm::ivec3(dataExtent[1] + 1, dataExtent[3] + 1, dataExtent[5] + 1);
 
@@ -62,6 +54,27 @@ void VolumeReader::Read(std::string filename)
     auto dataRange = imageData->GetScalarRange();
     Rescale<short, unsigned short>(reinterpret_cast<short*>(imageData->GetScalarPointer()), noEles, dataRange[0], dataRange[1]);
 
+    auto hist = vtkSmartPointer<vtkImageAccumulate>::New();
+    hist->SetInputConnection(imageCast->GetOutputPort());
+    hist->SetComponentExtent(0, dataRange[1] - dataRange[0] - 1, 0, 0, 0, 0);
+    hist->SetComponentOrigin(dataRange[0], 0, 0);
+    hist->SetComponentSpacing(1, 0, 0);
+    hist->IgnoreZeroOn();
+    hist->Update();
+
+    auto histDims = hist->GetOutput()->GetDimensions();
+    histogram.clear();
+    histogram.resize(histDims[0]);
+    memcpy(&histogram[0], hist->GetOutput()->GetScalarPointer(), sizeof(uint32_t) * histDims[0]);
+
+    //auto imageGradientMagnitude = vtkSmartPointer<vtkImageGradientMagnitude>::New();
+    //imageGradientMagnitude->SetDimensionality(3);
+    //imageGradientMagnitude->SetInput(imageCast->GetOutput());
+    //imageGradientMagnitude->Update();
+    //auto magData = imageGradientMagnitude->GetOutput();
+    //auto magRange = magData->GetScalarRange();
+    //maxMagnitude = magRange[2];
+
     //auto ptr = reinterpret_cast<unsigned short*>(data) + 60 * 512 * 512;
     //unsigned char* tmp = new unsigned char[512 * 512 * 3];
     //for(auto i = 0; i < 512; ++i)

core/VolumeReader.h

@@ -13,6 +13,7 @@
 #include <vtkImageData.h>
 #include <vtkImageCast.h>
 #include <vtkImageMagnitude.h>
+#include <vtkImageAccumulate.h>
 #include <vtkSmartPointer.h>
 #include <vtkErrorCode.h>
 
@@ -37,6 +38,9 @@ class VolumeReader
     void Rescale(T* dataPtr, size_t size, float dataMin, float dataMax);
     void CreateTextures();
 
+public:
+    std::vector<uint32_t> histogram;
+
 private:
     glm::vec3 spacing = glm::vec3(glm::uninitialize);
     glm::ivec3 dim = glm::ivec3(glm::uninitialize);

core/cuda_volume.h

@@ -44,7 +44,7 @@ class cudaVolume
         auto ydiff = GetIntensity(pointInWorld + glm::vec3(0.f, spacing.y, 0.f)) - GetIntensity(pointInWorld - glm::vec3(0.f, spacing.y, 0.f));
         auto zdiff = GetIntensity(pointInWorld + glm::vec3(0.f, 0.f, spacing.z)) - GetIntensity(pointInWorld - glm::vec3(0.f, 0.f, spacing.z));
 
-        return glm::vec3(xdiff, ydiff, zdiff);
+        return glm::vec3(xdiff, ydiff, zdiff) * 0.5f * invSpacing;
     }
 
     __device__ glm::vec3 NormalizedGradient(const glm::vec3& pointInWorld) const
@@ -57,6 +57,11 @@ class cudaVolume
         return bbox.IsInside(ptInWorld);
     }
 
+    __host__ __device__ glm::vec3 GetSize() const
+    {
+        return 1.f / bbox.invSize;
+    }
+
 private:
     __device__ glm::vec3 GetNormalizedTexCoord(const glm::vec3& pointInWorld) const
     {

core/lights/cuda_environment_light.h

@@ -0,0 +1,65 @@
+//
+// Created by 孙万捷 on 16/5/4.
+//
+
+#ifndef SUNPATHTRACER_CUDA_ENVIRONMENT_LIGHT_H
+#define SUNPATHTRACER_CUDA_ENVIRONMENT_LIGHT_H
+
+#define GLM_FORCE_INLINE
+#include <glm/glm.hpp>
+
+#include <cuda_runtime.h>
+
+class cudaEnvironmentLight
+{
+public:
+    __host__ __device__ cudaEnvironmentLight() {}
+
+    __host__ __device__ void Set(cudaTextureObject_t tex)
+    {
+        this->tex = tex;
+    }
+
+    __host__ __device__ void Set(const glm::vec3& radiance)
+    {
+        tex = 0;
+        defaultRadiance = radiance;
+    }
+
+    __host__ __device__ cudaTextureObject_t Get()
+    {
+        return tex;
+    }
+
+    __device__ glm::vec3 GetEnvRadiance(const glm::vec2& texcoord)
+    {
+#ifdef __CUDACC__
+        auto val = tex2D<float4>(tex, texcoord.x, texcoord.y);
+        return tex ? glm::vec3(val.x, val.y, val.z) : defaultRadiance;
+#else
+        return glm::vec3(0.f);
+#endif
+    }
+
+    __device__ glm::vec3 GetEnvRadiance(const glm::vec3& dir, float u_offset = 0.f, float v_offset = 0.f)
+    {
+        float theta = acosf(dir.y);
+        float phi = atan2f(dir.x, dir.z);
+        phi = phi < 0.f ? phi + 2.f * M_PI : phi;
+        float u = phi * 0.5f * M_1_PI;
+        float v = theta * M_1_PI;
+
+#ifdef __CUDACC__
+        auto val = tex2D<float4>(tex, u + u_offset, v + v_offset);
+        return tex ? glm::vec3(val.x, val.y, val.z) : defaultRadiance;
+#else
+        return glm::vec3(0.f);
+#endif
+    }
+
+private:
+    cudaTextureObject_t tex;
+    glm::vec3 defaultRadiance;
+};
+
+#endif //SUNPATHTRACER_CUDA_ENVIRONMENT_LIGHT_H

core/lights/cuda_lights.h

@@ -0,0 +1,27 @@
+//
+// Created by 孙万捷 on 16/5/25.
+//
+
+#ifndef SUNVOLUMERENDER_CUDALIGHTS_H
+#define SUNVOLUMERENDER_CUDALIGHTS_H
+
+#include <cuda_runtime.h>
+
+#define GLM_FORCE_INLINE
+#include <glm/glm.hpp>
+
+#include "cuda_environment_light.h"
+
+class cudaLights
+{
+public:
+    __device__ glm::vec3 GetEnvironmentRadiance(const glm::vec3& dir, float u_offset = 0.f, float v_offset = 0.f)
+    {
+        return environmentLight.GetEnvRadiance(dir, u_offset, v_offset);
+    }
+
+public:
+    cudaEnvironmentLight environmentLight;
+};
+
+#endif //SUNVOLUMERENDER_CUDALIGHTS_H

core/lights/lights.cpp

@@ -0,0 +1,80 @@
+//
+// Created by 孙万捷 on 16/5/25.
+//
+
+#include "lights.h"
+#include "../../utils/helper_cuda.h"
+
+#define STB_IMAGE_IMPLEMENTATION
+#include "../../utils/stb_image.h"
+
+Lights::Lights()
+{
+    environmentLight.Set(glm::vec3(0.03f));
+}
+
+Lights::~Lights()
+{
+    if(!environmentLight.Get())
+    {
+        checkCudaErrors(cudaDestroyTextureObject(environmentLight.Get()));
+        environmentLight.Set(0);
+    }
+
+    if(envMapArray)
+    {
+        checkCudaErrors(cudaFreeArray(envMapArray));
+        envMapArray = nullptr;
+    }
+}
+
+void Lights::SetEnvironmentLight(std::string filename)
+{
+    int w = 0, h = 0, n = 0;
+    float* data = stbi_loadf(filename.c_str(), &w, &h, &n, 0);
+    if(!data)
+    {
+        std::cerr<<"Unable to load environment map: "<<filename<<std::endl;
+        exit(0);
+    }
+
+    //create channel desc
+    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();
+    //create cudaArray
+    checkCudaErrors(cudaMallocArray(&envMapArray, &channelDesc, w, h));
+    if(n == 3)
+    {
+        uint32_t count = w * h;
+        std::vector<float4> ext_data;
+        ext_data.reserve(count);
+        for(auto i = 0; i < count; ++i)
+            ext_data.push_back(make_float4(data[i * 3], data[i * 3 + 1], data[i * 3 + 2], 0.f));
+
+        checkCudaErrors(cudaMemcpyToArray(envMapArray, 0, 0, ext_data.data(), sizeof(float4) * w * h, cudaMemcpyHostToDevice));
+    }
+    else
+        checkCudaErrors(cudaMemcpyToArray(envMapArray, 0, 0, data, sizeof(float4) * w * h, cudaMemcpyHostToDevice));
+    //create resource desc
+    cudaResourceDesc resDesc;
+    memset(&resDesc, 0, sizeof(resDesc));
+    resDesc.resType = cudaResourceTypeArray;
+    resDesc.res.array.array = envMapArray;
+    //create texture desc
+    cudaTextureDesc texDesc;
+    memset(&texDesc, 0, sizeof(texDesc));
+    texDesc.addressMode[0] = cudaAddressModeWrap;
+    texDesc.addressMode[1] = cudaAddressModeWrap;
+    texDesc.filterMode = cudaFilterModeLinear;
+    texDesc.readMode = cudaReadModeElementType;
+    texDesc.normalizedCoords = true;
+    //create cudaTextureObject
+    cudaTextureObject_t tex;
+    checkCudaErrors(cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL));
+
+    environmentLight.Set(tex);
+}
+
+void Lights::SetEnvionmentLight(const glm::vec3 &radiance)
+{
+    environmentLight.Set(radiance);
+}

core/lights/lights.h

@@ -0,0 +1,29 @@
+//
+// Created by 孙万捷 on 16/5/25.
+//
+
+#ifndef SUNVOLUMERENDER_LIGHTS_H
+#define SUNVOLUMERENDER_LIGHTS_H
+
+#include <iostream>
+#include <vector>
+#include <string>
+
+#include <glm/glm.hpp>
+
+#include "cuda_environment_light.h"
+
+class Lights
+{
+public:
+    Lights();
+    ~Lights();
+    void SetEnvironmentLight(std::string filename);
+    void SetEnvionmentLight(const glm::vec3& radiance);
+
+public:
+    cudaArray* envMapArray = nullptr;
+    cudaEnvironmentLight environmentLight;
+};
+
+#endif //SUNVOLUMERENDER_LIGHTS_H

core/pathtracer.h

@@ -20,5 +20,6 @@ extern "C" void rendering(glm::u8vec4* img, const RenderParams& renderParams);
 extern "C" void setup_volume(const cudaVolume& vol);
 extern "C" void setup_transferfunction(const cudaTransferFunction& tf);
 extern "C" void setup_camera(const cudaCamera& cam);
+extern "C" void setup_lights(const Lights& hostLights);
 
 #endif //SUNVOLUMERENDER_PATHTRACER_H

core/render_parameters.h

@@ -37,6 +37,7 @@ class RenderParams
     float exposure = 1.f;
     float maxOpacity = 0.5f;
     glm::vec3* hdrBuffer = nullptr;
+    glm::vec2 envLightOffset = glm::vec2(0.f);
 };
 
 #endif //SUNVOLUMERENDER_RENDER_PARAMETERS_H

gui/canvas.cpp

@@ -6,18 +6,13 @@
 
 Canvas::Canvas(const QGLFormat &format, QWidget *parent) : QGLWidget(format, parent)
 {
-    volumeReader.Read("../manix_small.mhd");
-    volumeReader.CreateDeviceVolume(&deviceVolume);
-    setup_volume(deviceVolume);
-
-    ZoomToExtent();
-    camera.Setup(glm::vec3(0.f, 0.f, eyeDist), glm::vec3(0.f, 0.f, 0.f), glm::vec3(0.f, 1.f, 0.f), fov, WIDTH, HEIGHT);
-    viewMat = glm::lookAt(glm::vec3(0.f, 0.f, eyeDist), glm::vec3(0.f), glm::vec3(0.f, 1.f, 0.f));
-    UpdateCamera();
+    // lights
+    lights.SetEnvironmentLight("LA_Downtown_Helipad_GoldenHour_Env.hdr");
+    setup_lights(lights);
 
     // render params
     renderParams.SetupHDRBuffer(WIDTH, HEIGHT);
-    renderParams.exposure = 3.f;
+    renderParams.exposure = 1.f;
     renderParams.traceDepth = 5;
 }
 
@@ -26,6 +21,21 @@ Canvas::~Canvas()
     renderParams.Clear();
 }
 
+void Canvas::LoadVolume(std::string filename)
+{
+    volumeReader.Read(filename);
+    volumeReader.CreateDeviceVolume(&deviceVolume);
+    setup_volume(deviceVolume);
+
+    ZoomToExtent();
+    camera.Setup(glm::vec3(0.f, 0.f, eyeDist), glm::vec3(0.f, 0.f, 0.f), glm::vec3(0.f, 1.f, 0.f), fov, WIDTH, HEIGHT);
+    viewMat = glm::lookAt(glm::vec3(0.f, 0.f, eyeDist), glm::vec3(0.f), glm::vec3(0.f, 1.f, 0.f));
+    UpdateCamera();
+
+    ready = true;
+    StartTimer();
+}
+
 void Canvas::initializeGL()
 {
     makeCurrent();
@@ -39,8 +49,6 @@ void Canvas::initializeGL()
     glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
 
     checkCudaErrors(cudaGraphicsGLRegisterBuffer(&resource, pbo, cudaGraphicsMapFlagsNone));
-
-    this->startTimer(0);
 }
 
 void Canvas::resizeGL(int w, int h)
@@ -52,6 +60,8 @@ void Canvas::paintGL()
 {
     glClear(GL_COLOR_BUFFER_BIT);
 
+    if(!ready) return;
+
     size_t size;
     checkCudaErrors(cudaGraphicsMapResources(1, &resource, 0));
     checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void**)&img, &size, resource));
@@ -103,7 +113,7 @@ void Canvas::mouseMoveEvent(QMouseEvent *e)
     // translation
     if(e->buttons() & Qt::MidButton)
     {
-        constexpr float baseTranslate = 5.f;
+        float baseTranslate = glm::length(deviceVolume.GetSize()) * 0.5f;
         cameraTranslate.x += static_cast<float>(delta.x() * baseTranslate);
         cameraTranslate.y += static_cast<float>(delta.y() * baseTranslate);