diff --git a/CMakeLists.txt b/CMakeLists.txt index 807cb48..f0c3e2e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -45,7 +45,7 @@ set(HOST_SOURCES main.cpp core/VolumeReader.cpp core/lights/lights.cpp gui/AddAreaLightDialog.cpp gui/AddAreaLightDialog.h) -set(DEVICE_SOURCES pathtracer.cu) +set(DEVICE_SOURCES pathtracer.cu raycasting.cu) cuda_compile(DEVICE_OBJS ${DEVICE_SOURCES}) diff --git a/core/cuda_volume.h b/core/cuda_volume.h index d0b6cf5..1b6588b 100644 --- a/core/cuda_volume.h +++ b/core/cuda_volume.h @@ -34,6 +34,8 @@ class cudaVolume __host__ __device__ void SetYClipPlane(const glm::vec2& clip) {y_clip = clip;} __host__ __device__ void SetZClipPlane(const glm::vec2& clip) {z_clip = clip;} + __host__ __device__ void SetDensityScale(float s = 1.f) {densityScale = s;} + __device__ float operator ()(const glm::vec3& pointInWorld) const { return GetIntensity(pointInWorld); @@ -83,7 +85,7 @@ class cudaVolume { #ifdef __CUDACC__ auto texCoord = GetNormalizedTexCoord(pointInWorld); - return tex3D(tex, texCoord.x, texCoord.y, texCoord.z); + return tex3D(tex, texCoord.x, texCoord.y, texCoord.z) * densityScale; #else return 0.f; #endif @@ -101,6 +103,7 @@ class cudaVolume private: cudaBBox bbox; cudaTextureObject_t tex; + float densityScale; glm::vec3 spacing; glm::vec3 invSpacing; glm::vec2 x_clip; diff --git a/gui/canvas.cpp b/gui/canvas.cpp index bdea8ca..d7704cd 100644 --- a/gui/canvas.cpp +++ b/gui/canvas.cpp @@ -26,6 +26,7 @@ void Canvas::LoadVolume(std::string filename) volumeReader.Read(filename); volumeReader.CreateDeviceVolume(&deviceVolume); deviceVolume.SetClipPlane(glm::vec2(-1.f, 1.f), glm::vec2(-1.f, 1.f), glm::vec2(-1.f, 1.f)); + deviceVolume.SetDensityScale(1.f); setup_volume(deviceVolume); ZoomToExtent(); @@ -85,7 +86,7 @@ void Canvas::paintGL() glVertex2f(1.f, -1.f); glEnd(); - render_raycasting(img, volumeReader.GetElementBoundingSphereRadius()); + render_raycasting(img, deviceVolume, transferFunction, camera, volumeReader.GetElementBoundingSphereRadius()); } else { diff --git a/gui/canvas.h b/gui/canvas.h index 6b57129..76ff693 100644 --- a/gui/canvas.h +++ b/gui/canvas.h @@ -36,12 +36,6 @@ class Canvas : public QGLWidget explicit Canvas(const QGLFormat& format, QWidget* parent = 0); virtual ~Canvas(); - void SetTransferFunction(const cudaTextureObject_t& tex, float maxOpacity) - { - transferFunction.Set(tex, maxOpacity); - setup_transferfunction(transferFunction); - ReStartRender(); - }; void LoadVolume(std::string filename); void StartTimer() {timerId = this->startTimer(0);} void KillTimer() {this->killTimer(timerId);} @@ -52,6 +46,20 @@ class Canvas : public QGLWidget renderParams.frameNo = 0; } + void SetTransferFunction(const cudaTextureObject_t& tex, float maxOpacity) + { + transferFunction.Set(tex, maxOpacity); + setup_transferfunction(transferFunction); + ReStartRender(); + } + + void SetDensityScale(double s) + { + deviceVolume.SetDensityScale(s); + setup_volume(deviceVolume); + ReStartRender(); + } + void SetScatterTimes(double val) { renderParams.traceDepth = val; diff --git a/gui/mainwindow.cpp b/gui/mainwindow.cpp index 47f7c10..71a23a8 100644 --- a/gui/mainwindow.cpp +++ b/gui/mainwindow.cpp @@ -72,6 +72,8 @@ void MainWindow::ConfigureTransferFunction() ui->colorTransferFunc->view()->setAxesToChartBounds(); connect(tf, SIGNAL(Changed()), this, SLOT(onTransferFunctionChanged())); + + connect(ui->SliderWidget_DensityScale, SIGNAL(valueChanged(double)), this, SLOT(onDensityScaleChanged(double))); } void MainWindow::onTransferFunctionChanged() @@ -465,3 +467,8 @@ void MainWindow::onZClipChanged(double min, double max) { canvas->SetZClipPlane(min, max); } + +void MainWindow::onDensityScaleChanged(double s) +{ + canvas->SetDensityScale(s); +} diff --git a/gui/mainwindow.h b/gui/mainwindow.h index 6c7db77..c940b93 100644 --- a/gui/mainwindow.h +++ b/gui/mainwindow.h @@ -42,6 +42,7 @@ class MainWindow : public QMainWindow private slots: void onTransferFunctionChanged(); + void onDensityScaleChanged(double s); void onScatterTimesChanged(double val); void onEnvLightUOffsetChanged(double u); diff --git a/gui/mainwindow.ui b/gui/mainwindow.ui index 0611148..5d6f189 100644 --- a/gui/mainwindow.ui +++ b/gui/mainwindow.ui @@ -110,9 +110,9 @@ 0 - 0 + -144 363 - 694 + 730 @@ -179,6 +179,30 @@ Transfer Function + + + + + + Density scale + + + + + + + 3 + + + 0.001000000000000 + + + 1.000000000000000 + + + + + diff --git a/pathtracer.cu b/pathtracer.cu index c4481e7..8931008 100644 --- a/pathtracer.cu +++ b/pathtracer.cu @@ -298,67 +298,4 @@ extern "C" void render_pathtracer(glm::u8vec4* img, const RenderParams& renderPa kernel_pathtracer<<>>(renderParams, wangHash(renderParams.frameNo)); hdr_to_ldr<<>>(img, renderParams); -} - -// front to back composite -__global__ void kernel_raycasting(glm::u8vec4* img, float stepSize) -{ - auto idx = blockDim.x * blockIdx.x + threadIdx.x; - auto idy = blockDim.y * blockIdx.y + threadIdx.y; - auto offset = idy * WIDTH + idx; - - cudaRay ray; - camera.GenerateRay(idx, idy, &ray); - - glm::vec4 L = glm::vec4(0.f); - - float tNear, tFar, t; - if(volume.Intersect(ray, &tNear, &tFar)) - { - t = tNear; - while(t <= tFar) - { - auto ptInWorld = ray.PointOnRay(t); - auto intensity = volume(ptInWorld); - auto color_opacity = transferFunction(intensity); - - // apply lighting - auto gradient = volume.Gradient_CentralDiff(ptInWorld); - auto gradientMagnitude = sqrtf(glm::dot(gradient, gradient)); - float cosTerm = 1.f; - float specularTerm = 0.f; - if(gradientMagnitude > 1e-3) - { - auto normal = glm::normalize(gradient); - auto lightDir = glm::normalize(camera.pos - ptInWorld); - cosTerm = fabsf(glm::dot(normal, lightDir)); - - specularTerm = powf(cosTerm, 30.f); - } - - color_opacity.x = color_opacity.x * color_opacity.w * cosTerm * 0.8f + color_opacity.w * specularTerm * 0.2f; - color_opacity.y = color_opacity.y * color_opacity.w * cosTerm * 0.8f + color_opacity.w * specularTerm * 0.2f; - color_opacity.z = color_opacity.z * color_opacity.w * cosTerm * 0.8f + color_opacity.w * specularTerm * 0.2f; - - L += (1.f - L.w) * color_opacity; - - if(L.w > 0.95f) break; - - t += stepSize * 0.5f; - } - - } - - L.x = fminf(L.x, 1.f); - L.y = fminf(L.y, 1.f); - L.z = fminf(L.z, 1.f); - img[offset] = glm::u8vec4(L.x * 255, L.y * 255, L.z * 255, 255 * L.w); -} - -extern "C" void render_raycasting(glm::u8vec4* img, float stepSize) -{ - dim3 blockSize(16, 16); - dim3 gridSize(WIDTH / blockSize.x, HEIGHT / blockSize.y); - - kernel_raycasting<<>>(img, stepSize); } \ No newline at end of file diff --git a/raycasting.cu b/raycasting.cu new file mode 100644 index 0000000..671dfb7 --- /dev/null +++ b/raycasting.cu @@ -0,0 +1,75 @@ +#define GLM_FORCE_NO_CTOR_INIT +#define GLM_FORCE_INLINE +#include + +#include +#include +#include + +#include "common.h" +#include "utils/helper_cuda.h" +#include "core/cuda_camera.h" +#include "core/cuda_transfer_function.h" +#include "core/cuda_volume.h" + +__global__ void kernel_raycasting(glm::u8vec4* img, cudaVolume volume, cudaTransferFunction transferFunction, cudaCamera camera, float stepSize) +{ + auto idx = blockDim.x * blockIdx.x + threadIdx.x; + auto idy = blockDim.y * blockIdx.y + threadIdx.y; + auto offset = idy * WIDTH + idx; + + cudaRay ray; + camera.GenerateRay(idx, idy, &ray); + + glm::vec4 L = glm::vec4(0.f); + + float tNear, tFar, t; + if(volume.Intersect(ray, &tNear, &tFar)) + { + t = tNear; + while(t <= tFar) + { + auto ptInWorld = ray.PointOnRay(t); + auto intensity = volume(ptInWorld); + auto color_opacity = transferFunction(intensity); + + // apply lighting + auto gradient = volume.Gradient_CentralDiff(ptInWorld); + auto gradientMagnitude = sqrtf(glm::dot(gradient, gradient)); + float cosTerm = 1.f; + float specularTerm = 0.f; + if(gradientMagnitude > 1e-3) + { + auto normal = glm::normalize(gradient); + auto lightDir = glm::normalize(camera.pos - ptInWorld); + cosTerm = fabsf(glm::dot(normal, lightDir)); + + specularTerm = powf(cosTerm, 30.f); + } + + color_opacity.x = color_opacity.x * color_opacity.w * cosTerm * 0.8f + color_opacity.w * specularTerm * 0.2f; + color_opacity.y = color_opacity.y * color_opacity.w * cosTerm * 0.8f + color_opacity.w * specularTerm * 0.2f; + color_opacity.z = color_opacity.z * color_opacity.w * cosTerm * 0.8f + color_opacity.w * specularTerm * 0.2f; + + L += (1.f - L.w) * color_opacity; + + if(L.w > 0.95f) break; + + t += stepSize * 0.5f; + } + + } + + L.x = fminf(L.x, 1.f); + L.y = fminf(L.y, 1.f); + L.z = fminf(L.z, 1.f); + img[offset] = glm::u8vec4(L.x * 255, L.y * 255, L.z * 255, 255 * L.w); +} + +extern "C" void render_raycasting(glm::u8vec4* img, cudaVolume& volume, cudaTransferFunction& transferFunction, cudaCamera& camera, float stepSize) +{ + dim3 blockSize(16, 16); + dim3 gridSize(WIDTH / blockSize.x, HEIGHT / blockSize.y); + + kernel_raycasting<<>>(img, volume, transferFunction, camera, stepSize); +} \ No newline at end of file diff --git a/raycasting.h b/raycasting.h index 9a2f306..51606f8 100644 --- a/raycasting.h +++ b/raycasting.h @@ -5,6 +5,6 @@ #ifndef SUNVOLUMERENDER_RAYCASTING_H_H #define SUNVOLUMERENDER_RAYCASTING_H_H -extern "C" void render_raycasting(glm::u8vec4* img, float stepSize); +extern "C" void render_raycasting(glm::u8vec4* img, cudaVolume& volume, cudaTransferFunction& transferFunction, cudaCamera& camera, float stepSize); #endif //SUNVOLUMERENDER_RAYCASTING_H_H