From c23a33ebef85336d45217b65f279442d875cc744 Mon Sep 17 00:00:00 2001 From: djairoh Date: Fri, 3 Jan 2025 19:57:10 +0100 Subject: [PATCH] feat (wip): real-time gui with rendering in opengl --- src/consts.cu | 2 +- src/consts.h | 6 +- src/gui/MainWindow.cpp | 125 ++++++++++++++++ src/gui/MainWindow.h | 32 +++++ src/gui/Quad.cpp | 136 ++++++++++++++++++ src/gui/Quad.h | 36 +++++ src/gui/shaders/rendertype_accumulate.frag | 19 +++ src/gui/shaders/rendertype_accumulate.vert | 11 ++ src/gui/shaders/rendertype_screen.frag | 12 ++ src/gui/shaders/rendertype_screen.vert | 11 ++ src/illumination/FrameBuffer.cu | 13 ++ src/illumination/FrameBuffer.h | 20 +++ .../{raycaster.h => Raycaster.cu} | 121 ++++++++++++++-- src/illumination/Raycaster.h | 33 +++++ src/illumination/illumination.h | 4 +- src/illumination/shading.cu | 14 ++ src/illumination/shading.h | 14 +- src/linalg/mat.cu | 46 ++++++ src/linalg/mat.h | 26 +--- src/linalg/vec.h | 2 +- src/main.cpp | 8 ++ src/main.cu | 27 ++-- 22 files changed, 656 insertions(+), 62 deletions(-) create mode 100644 src/gui/MainWindow.cpp create mode 100644 src/gui/MainWindow.h create mode 100644 src/gui/Quad.cpp create mode 100644 src/gui/Quad.h create mode 100644 src/gui/shaders/rendertype_accumulate.frag create mode 100644 src/gui/shaders/rendertype_accumulate.vert create mode 100644 src/gui/shaders/rendertype_screen.frag create mode 100644 src/gui/shaders/rendertype_screen.vert create mode 100644 src/illumination/FrameBuffer.cu create mode 100644 src/illumination/FrameBuffer.h rename src/illumination/{raycaster.h => Raycaster.cu} (58%) create mode 100644 src/illumination/Raycaster.h create mode 100644 src/illumination/shading.cu create mode 100644 src/linalg/mat.cu create mode 100644 src/main.cpp diff --git a/src/consts.cu b/src/consts.cu index 18cf40a..b432f6d 100644 --- a/src/consts.cu +++ b/src/consts.cu @@ -15,4 +15,4 @@ void copyConstantsToDevice() { cudaMemcpyToSymbol(d_cameraDir, &h_cameraDir, sizeof(Vec3)); cudaMemcpyToSymbol(d_cameraUp, &h_cameraUp, sizeof(Vec3)); cudaMemcpyToSymbol(d_lightPos, &h_lightPos, sizeof(Point3)); -} \ No newline at end of file +} diff --git a/src/consts.h b/src/consts.h index 774c297..2bb8caa 100644 --- a/src/consts.h +++ b/src/consts.h @@ -9,8 +9,8 @@ const int VOLUME_WIDTH = 49; const int VOLUME_HEIGHT = 51; const int VOLUME_DEPTH = 42; -const int IMAGE_WIDTH = 2560; -const int IMAGE_HEIGHT = 1440; +const int IMAGE_WIDTH = 800; +const int IMAGE_HEIGHT = 600; const double epsilon = 1e-10f; const double infty = 1e15f; // This vlalue is used to represent missing values in data @@ -41,4 +41,4 @@ extern __device__ Point3 d_lightPos; // --------------------------- Functions for handling external constants --------------------------- void copyConstantsToDevice(); -#endif // CONSTS_H \ No newline at end of file +#endif // CONSTS_H diff --git a/src/gui/MainWindow.cpp b/src/gui/MainWindow.cpp new file mode 100644 index 0000000..2878a25 --- /dev/null +++ b/src/gui/MainWindow.cpp @@ -0,0 +1,125 @@ +#include "MainWindow.h" + +#include +#include + +Window::Window(unsigned int w, unsigned int h) { + Window::w = w; + Window::h = h; +} + +void framebuffer_size_callback(GLFWwindow* window, int w, int h) { + // This function is called by glfw when the window is reized. + glViewport(0 , 0, w, h); + Window* newWin = reinterpret_cast(glfwGetWindowUserPointer(window)); + newWin->resize(w, h); + +} + +int Window::init() { + // init glfw + glfwInit(); + // requesting context version 1.0 makes glfw try to provide the latest version if possible + glfwWindowHint(GLFW_CONTEXT_VERSION_MAJOR, 1); + glfwWindowHint(GLFW_CONTEXT_VERSION_MINOR, 0); + + this->window = glfwCreateWindow(this->w, this->h, "CUDA ray tracing", NULL, NULL); + + //hide cursor + glfwSetInputMode(this->window, GLFW_CURSOR, GLFW_CURSOR_DISABLED); + glfwSetWindowUserPointer(this->window, reinterpret_cast(this)); + + if (this->window == NULL) { + std::cout << "Failed to create window\n"; + glfwTerminate(); + return -1; + } + + glfwMakeContextCurrent(this->window); + + + // init glad(opengl) + if (!gladLoadGLLoader((GLADloadproc)glfwGetProcAddress)) { + std::cout << "Failed to initialize GLAD\n"; + return -1; + } + + + // init framebuffer + glViewport(0, 0, this->w, this->h); + if (glfwSetFramebufferSizeCallback(this->window, framebuffer_size_callback) != 0) return -1; + + + if (init_quad()) return -1; + this->last_frame = std::chrono::steady_clock::now(); + + while (!glfwWindowShouldClose(window)) { + Window::tick(); + } + + Window::free(); + return 0; +} + + +int Window::init_quad() { + this->current_quad = std::make_unique(this->w, this->h); + this->current_quad->cuda_init(); + + // TODO: default shaders + return 0; +} + + +void Window::free() { + // To preserve the proper destruction order we forcefully set the quads to null (calling their destructor in the process) + // Not strictly necessary, but i saw some weird errors on exit without this so best to keep it in. + this->current_quad = nullptr; + + glfwDestroyWindow(window); + glfwTerminate(); +} + + +void Window::tick() { + // manually track time diff + std::chrono::steady_clock::time_point now = std::chrono::steady_clock::now(); + float diff = (float) std::chrono::duration_cast(now - this->last_frame).count(); + this->last_frame = now; + + // TODO: remove debug line at some point + std::cout << 1000.0/diff << " fps\n"; + + // TODO: code input logic and class/struct and stuff + // ticking input probably involves 4? steps: + // * check if window needs to be closed (escape/q pressed) + // * check if camera moved (wasd/hjkl pressed) + // (phase 3/do later): check if we switched from realtime tracing to that other option - maybe a pause function? (p pressed?) + // * if moved -> update camera (raytracing will involve some logic here too? see when i get there) + + // tick render + glBindFramebuffer(GL_FRAMEBUFFER, 0); + glDisable(GL_DEPTH_TEST); + + glClearColor(0.2f, 0.3f, 0.3f, 1.0f); + glClear(GL_COLOR_BUFFER_BIT); + + // render frame + glBindFramebuffer(GL_FRAMEBUFFER, this->current_quad->fb); + this->current_quad->render(); + glBindVertexArray(this->current_quad->VAO); + glBindTexture(GL_TEXTURE_2D, this->current_quad->tex); + glDrawArrays(GL_TRIANGLES, 0, 6); // draw current frame to texture + + // check for events + // + swap buffers; TODO: check if necessary? + glfwSwapBuffers(this->window); + glfwPollEvents(); + +} + +void Window::resize(unsigned int w, unsigned int h) { + this->w = w; + this->h = h; + this->current_quad->resize(w, h); +} diff --git a/src/gui/MainWindow.h b/src/gui/MainWindow.h new file mode 100644 index 0000000..bd7ea79 --- /dev/null +++ b/src/gui/MainWindow.h @@ -0,0 +1,32 @@ +#ifndef MAINWINDOW_H +#define MAINWINDOW_H + +#include "Quad.h" +#include +#include +#include + + +class Window { +public: + unsigned int w; + unsigned int h; + + Window(unsigned int w, unsigned int h); + + int init(); + void free(); + void resize(unsigned int w, unsigned int h); + +private: + GLFWwindow* window; + std::unique_ptr current_quad; + + std::chrono::steady_clock::time_point last_frame; + + void tick(); + int init_quad(); + + // std::unique_ptr shader; +}; +#endif // MAINWINDOW_H diff --git a/src/gui/Quad.cpp b/src/gui/Quad.cpp new file mode 100644 index 0000000..5a3ffd4 --- /dev/null +++ b/src/gui/Quad.cpp @@ -0,0 +1,136 @@ +#include "Quad.h" + +#include +#include +#include "cuda_runtime.h" +#include "device_launch_parameters.h" +#include +#include + +Quad::Quad(unsigned int w, unsigned int h) { + this->w = w; + this->h = h; + + std::vector vertices = { + -1.0f, 1.0f, 0.0f, 1.0f, + -1.0f, -1.0f, 0.0f, 0.0f, + 1.0f, -1.0f, 1.0f, 0.0f, + + -1.0f, 1.0f, 0.0f, 1.0f, + 1.0f, -1.0f, 1.0f, 0.0f, + 1.0f, 1.0f, 1.0f, 1.0f + }; + + glGenBuffers(1, &VBO); + glGenVertexArrays(1, &VAO); + + glBindVertexArray(VAO); + glBindBuffer(GL_ARRAY_BUFFER, VBO); + + // copy vertex data to buffer on gpu + glBufferData(GL_ARRAY_BUFFER, vertices.size() * sizeof(float), vertices.data(), GL_STATIC_DRAW); + + // set our vertex attributes pointers + glEnableVertexAttribArray(0); + glVertexAttribPointer(0, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float), (void*)0); + + glEnableVertexAttribArray(1); + glVertexAttribPointer(1, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float), (void*)(2 * sizeof(float))); + + glBindBuffer(GL_ARRAY_BUFFER, 0); + glBindVertexArray(0); + + // texture stuff + glGenBuffers(1, &PBO); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, PBO); + glBufferData(GL_PIXEL_UNPACK_BUFFER, w * h * 4, NULL, GL_DYNAMIC_COPY); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + + glEnable(GL_TEXTURE_2D); + + glGenTextures(1, &tex); + + glBindTexture(GL_TEXTURE_2D, tex); + + // parameters for texture + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, w, h, 0, GL_BGRA, GL_UNSIGNED_BYTE, NULL); + glBindTexture(GL_TEXTURE_2D, 0); + + // register the FBO + glGenFramebuffers(1, &fb); + glBindFramebuffer(GL_FRAMEBUFFER, fb); + glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, this->tex, 0); + glBindFramebuffer(GL_FRAMEBUFFER, 0); +}; + + +Quad::~Quad() { + int res = cudaGraphicsUnregisterResource(CGR); + if (res) { + std::cout << "CUDA error while deregistering the graphics resource: " << res; + cudaDeviceReset(); + exit(1); + } +}; + + +void Quad::cuda_init() { + int res = cudaGraphicsGLRegisterBuffer(&this->CGR, this->PBO, cudaGraphicsRegisterFlagsNone); + if (res) { + std::cout << "CUDA error while registering the graphics resource: " << res; + cudaDeviceReset(); + exit(1); + } + this->renderer = std::make_unique(this->CGR, this->w, this->h); +}; + + +void Quad::render() { + glBindTexture(GL_TEXTURE_2D, 0); + this->renderer->render(); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, this->PBO); + glBindTexture(GL_TEXTURE_2D, this->tex); + glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, this->w, this->h, GL_BGRA, GL_UNSIGNED_BYTE, NULL); +}; + + +void Quad::resize(unsigned int w, unsigned int h) { + this->w = w; + this->h = h; + + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, this->PBO); + glBufferData(GL_PIXEL_UNPACK_BUFFER, w * h * 4, NULL, GL_DYNAMIC_COPY); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + + glBindTexture(GL_TEXTURE_2D, this->tex); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, w, h, 0, GL_BGRA, GL_UNSIGNED_BYTE, NULL); + glBindTexture(GL_TEXTURE_2D, 0); + + glBindFramebuffer(GL_FRAMEBUFFER, this->fb); + glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, this->tex, 0); + glBindFramebuffer(GL_FRAMEBUFFER, 0); + + if (this->renderer != nullptr) { + // TODO: probably make a function for the cuda error checking + int res = cudaGraphicsUnregisterResource(CGR); + if (res) { + std::cout << "CUDA error while deregistering the graphics resource: " << res; + cudaDeviceReset(); + exit(1); + } + + res = cudaGraphicsGLRegisterBuffer(&this->CGR, this->PBO, cudaGraphicsRegisterFlagsNone); + if (res) { + std::cout << "CUDA error while registering the graphics resource: " << res; + cudaDeviceReset(); + exit(1); + } + this->renderer->resources = this->CGR; + this->renderer->resize(w, h); + } +}; diff --git a/src/gui/Quad.h b/src/gui/Quad.h new file mode 100644 index 0000000..6e6f9ec --- /dev/null +++ b/src/gui/Quad.h @@ -0,0 +1,36 @@ +#ifndef QUAD_H +#define QUAD_H + +#include +#include +#include + +#include "illumination/Raycaster.h" + +#include +#include + +class Quad { +public: + unsigned int VAO; + unsigned int VBO; + unsigned int PBO; + cudaGraphicsResource_t CGR; + + unsigned int tex; + unsigned int fb; + + unsigned int w; + unsigned int h; + + std::unique_ptr renderer; + + Quad(unsigned int w, unsigned int h); + ~Quad(); + + void render(); + void resize(unsigned int w, unsigned int h); + void cuda_init(); + +}; +#endif // QUAD_H diff --git a/src/gui/shaders/rendertype_accumulate.frag b/src/gui/shaders/rendertype_accumulate.frag new file mode 100644 index 0000000..cce423c --- /dev/null +++ b/src/gui/shaders/rendertype_accumulate.frag @@ -0,0 +1,19 @@ +#version 330 core +out vec4 FragColor; + +in vec2 TexCoords; + +uniform sampler2D currentFrameTex; +uniform sampler2D lastFrameTex; +uniform int frameCount; + +#define MAX_FRAMES 500.0f + +void main() +{ + vec3 col = texture(currentFrameTex, TexCoords).rgb; + vec3 col2 = texture(lastFrameTex, TexCoords).rgb; + + col = mix(col, col2, min(frameCount/MAX_FRAMES,1.0f)); + FragColor = vec4(col, 1.0); +} \ No newline at end of file diff --git a/src/gui/shaders/rendertype_accumulate.vert b/src/gui/shaders/rendertype_accumulate.vert new file mode 100644 index 0000000..737bdb3 --- /dev/null +++ b/src/gui/shaders/rendertype_accumulate.vert @@ -0,0 +1,11 @@ +#version 330 core +layout(location = 0) in vec2 aPos; +layout(location = 1) in vec2 aTexCoords; + +out vec2 TexCoords; + +void main() +{ + TexCoords = aTexCoords; + gl_Position = vec4(aPos.x, aPos.y, 0.0, 1.0); +} \ No newline at end of file diff --git a/src/gui/shaders/rendertype_screen.frag b/src/gui/shaders/rendertype_screen.frag new file mode 100644 index 0000000..23016a3 --- /dev/null +++ b/src/gui/shaders/rendertype_screen.frag @@ -0,0 +1,12 @@ +#version 330 core +out vec4 FragColor; + +in vec2 TexCoords; + +uniform sampler2D screenTexture; + +void main() +{ + vec3 col = texture(screenTexture, TexCoords).rgb; + FragColor = vec4(col, 1.0); +} \ No newline at end of file diff --git a/src/gui/shaders/rendertype_screen.vert b/src/gui/shaders/rendertype_screen.vert new file mode 100644 index 0000000..737bdb3 --- /dev/null +++ b/src/gui/shaders/rendertype_screen.vert @@ -0,0 +1,11 @@ +#version 330 core +layout(location = 0) in vec2 aPos; +layout(location = 1) in vec2 aTexCoords; + +out vec2 TexCoords; + +void main() +{ + TexCoords = aTexCoords; + gl_Position = vec4(aPos.x, aPos.y, 0.0, 1.0); +} \ No newline at end of file diff --git a/src/illumination/FrameBuffer.cu b/src/illumination/FrameBuffer.cu new file mode 100644 index 0000000..0edcc58 --- /dev/null +++ b/src/illumination/FrameBuffer.cu @@ -0,0 +1,13 @@ +#include "FrameBuffer.h" +#include "linalg/linalg.h" + + +__host__ FrameBuffer::FrameBuffer(unsigned int w, unsigned int h) : w(w), h(h) {} + + +__device__ void FrameBuffer::writePixel(int x, int y, float r, float g, float b) { + int i = y * this->w + x; + + // the opengl buffer uses BGRA format; dunno why + this->buffer[i] = packUnorm4x8(b, g, r, 1.0f); +} diff --git a/src/illumination/FrameBuffer.h b/src/illumination/FrameBuffer.h new file mode 100644 index 0000000..a2577fe --- /dev/null +++ b/src/illumination/FrameBuffer.h @@ -0,0 +1,20 @@ +#ifndef FRAMEBUFFER_H +#define FRAMEBUFFER_H + +#include "cuda_runtime.h" +#include "linalg/linalg.h" +#include + + +class FrameBuffer { +public: + uint32_t* buffer; + std::size_t buffer_size; + unsigned int w; + unsigned int h; + + __host__ FrameBuffer(unsigned int w, unsigned int h); + __device__ void writePixel(int x, int y, float r, float g, float b); +}; + +#endif // FRAMEBUFFER_H diff --git a/src/illumination/raycaster.h b/src/illumination/Raycaster.cu similarity index 58% rename from src/illumination/raycaster.h rename to src/illumination/Raycaster.cu index a1c4292..4f7b88e 100644 --- a/src/illumination/raycaster.h +++ b/src/illumination/Raycaster.cu @@ -1,14 +1,16 @@ -#ifndef RAYCASTER_H -#define RAYCASTER_H +#include "Raycaster.h" + +#include "cuda_runtime.h" +#include "device_launch_parameters.h" -#include #include "linalg/linalg.h" #include "consts.h" #include "shading.h" +#include +#include "objs/sphere.h" -// Raycast + phong, TODO: Consider wrapping in a class -__global__ void raycastKernel(float* volumeData, unsigned char* framebuffer) { +__global__ void raycastKernel(float* volumeData, FrameBuffer framebuffer) { int px = blockIdx.x * blockDim.x + threadIdx.x; int py = blockIdx.y * blockDim.y + threadIdx.y; if (px >= IMAGE_WIDTH || py >= IMAGE_HEIGHT) return; @@ -117,10 +119,109 @@ __global__ void raycastKernel(float* volumeData, unsigned char* framebuffer) { accumB /= (float)SAMPLES_PER_PIXEL; // Final colour - int fbIndex = (py * IMAGE_WIDTH + px) * 3; - framebuffer[fbIndex + 0] = (unsigned char)(fminf(accumR, 1.f) * 255); - framebuffer[fbIndex + 1] = (unsigned char)(fminf(accumG, 1.f) * 255); - framebuffer[fbIndex + 2] = (unsigned char)(fminf(accumB, 1.f) * 255); + framebuffer.writePixel(px, py, accumR, accumG, accumB); + // int fbIndex = (py * IMAGE_WIDTH + px) * 3; + // framebuffer[fbIndex + 0] = (unsigned char)(fminf(accumR, 1.f) * 255); + // framebuffer[fbIndex + 1] = (unsigned char)(fminf(accumG, 1.f) * 255); + // framebuffer[fbIndex + 2] = (unsigned char)(fminf(accumB, 1.f) * 255); } -#endif // RAYCASTER_H \ No newline at end of file + +Raycaster::Raycaster(cudaGraphicsResource_t resources, int w, int h) { + this->resources = resources; + this->w = h; + this->w = h; + + this->fb = new FrameBuffer(w, h); + + // camera_info = CameraInfo(Vec3(0.0f, 0.0f, 0.0f), Vec3(0.0f, 0.0f, 0.0f), 90.0f, (float) w, (float) h); + // d_camera = thrust::device_new(); + + int res = cudaDeviceSynchronize(); + if (res) { + std::cout << "CUDA error while synchronizing device: " << res; + cudaDeviceReset(); + exit(1); + } + + res = cudaDeviceSynchronize(); + if (res) { + std::cout << "CUDA error while synchronizing device: " << res; + cudaDeviceReset(); + exit(1); + } +} + + +void Raycaster::render() { + int res = cudaGraphicsMapresources(1, this->resources); + if (res) { + std::cout << "CUDA error while mapping graphic resource: " << res; + cudaDeviceReset(); + exit(1); + } + + // check_cuda_errors(cudaGraphicsResourceGetMappedPointer((void**)&(frame_buffer->device_ptr), &(frame_buffer->buffer_size), resources)); + res = cudaGraphicsResourceGetMappedPointer((void**)(this->fb->buffer), &this->fb->buffer_size, this->resources); + if (res) { + std::cout << "CUDA error while fetching resource pointer: " << res; + cudaDeviceReset(); + exit(1); + } + + // FIXME: might not be the best parallelization configuraiton + int tx = 32; + int ty = 32; + + dim3 blocks(this->w / tx + 1, this->h / ty + 1); + dim3 threads(tx, ty); + + // TODO: pass camera info at some point + // TODO: pass float volume data. + // frame buffer is implicitly copied to the device each frame + raycastKernel<<>> (nullptr, this->fb); + + res = cudaGetLastError(); + if (res) { + std::cout << "CUDA error while raycasting: " << res; + cudaDeviceReset(); + exit(1); + } + + res = cudaDeviceSynchronize(); + if (res) { + std::cout << "CUDA error while synchronizing device: " << res; + cudaDeviceReset(); + exit(1); + } + + res = cudaGraphicsUnmapResources(1, &this->resources); + if (res) { + std::cout << "CUDA error while unmapping a resource: " << res; + cudaDeviceReset(); + exit(1); + } +} + + +void Raycaster::resize(int w, int h) { + this->w = w; + this->h = h; + + delete fb; + this->fb = new FrameBuffer(w, h); + + // TODO: should be globals probably + int tx = 8; + int ty = 8; + + dim3 blocks(w / tx + 1, h / ty + 1); + dim3 threads(tx, ty); + + int res = cudaDeviceSynchronize(); + if (res != 0) { + std::cout << "CUDA error while synchronizing device: " << res; + cudaDeviceReset(); + exit(1); + } +} diff --git a/src/illumination/Raycaster.h b/src/illumination/Raycaster.h new file mode 100644 index 0000000..170eb41 --- /dev/null +++ b/src/illumination/Raycaster.h @@ -0,0 +1,33 @@ +#ifndef RAYCASTER_H +#define RAYCASTER_H + +// #include "Camera.h" +#include "cuda_runtime.h" +#include "FrameBuffer.h" +#include "linalg/linalg.h" + +// #include + +__global__ void raycastKernel(float* volumeData, unsigned char* framebuffer); + +struct Raycaster { + + // thrust::device_ptr d_camera; + // CameraInfo camera_info; + + cudaGraphicsResource_t resources; + FrameBuffer* fb; + + int w; + int h; + + Raycaster() {}; + Raycaster(cudaGraphicsResource_t resources, int nx, int ny); + // ~Raycaster(); + + void set_camera(Vec3 position, Vec3 forward, Vec3 up); + void render(); + void resize(int nx, int ny); + // void raycastKernel(float* volumeData, unsigned char* framebuffer); // TODO: proper framebuffer class +}; +#endif // RAYCASTER_H diff --git a/src/illumination/illumination.h b/src/illumination/illumination.h index 6e9b6f0..3d2e90d 100644 --- a/src/illumination/illumination.h +++ b/src/illumination/illumination.h @@ -1,7 +1,7 @@ #ifndef ILLUMINATION_H #define ILLUMINATION_H -#include "raycaster.h" +#include "Raycaster.h" #include "shading.h" -#endif // ILLUMINATION_H \ No newline at end of file +#endif // ILLUMINATION_H diff --git a/src/illumination/shading.cu b/src/illumination/shading.cu new file mode 100644 index 0000000..41141f0 --- /dev/null +++ b/src/illumination/shading.cu @@ -0,0 +1,14 @@ +#include "shading.h" + +// TODO: Consider wrapping this in a class (?) +__device__ Vec3 phongShading(const Vec3& normal, const Vec3& lightDir, const Vec3& viewDir, const Vec3& baseColor) { + Vec3 ambient = baseColor * ambientStrength; + double diff = fmax(normal.dot(lightDir), 0.0); + Vec3 diffuse = baseColor * (diffuseStrength * diff); + + Vec3 reflectDir = (normal * (2.0 * normal.dot(lightDir)) - lightDir).normalize(); + double spec = pow(fmax(viewDir.dot(reflectDir), 0.0), shininess); + Vec3 specular = Vec3::init(1.0, 1.0, 1.0) * (specularStrength * spec); + + return ambient + diffuse + specular; +}; diff --git a/src/illumination/shading.h b/src/illumination/shading.h index fa3ee8e..8957c71 100644 --- a/src/illumination/shading.h +++ b/src/illumination/shading.h @@ -4,17 +4,7 @@ #include "linalg/linalg.h" #include "consts.h" -// TODO: Consider wrapping this in a class (?) -__device__ Vec3 phongShading(const Vec3& normal, const Vec3& lightDir, const Vec3& viewDir, const Vec3& baseColor) { - Vec3 ambient = baseColor * ambientStrength; - double diff = fmax(normal.dot(lightDir), 0.0); - Vec3 diffuse = baseColor * (diffuseStrength * diff); +__device__ Vec3 phongShading(const Vec3& normal, const Vec3& lightDir, const Vec3& viewDir, const Vec3& baseColor); - Vec3 reflectDir = (normal * (2.0 * normal.dot(lightDir)) - lightDir).normalize(); - double spec = pow(fmax(viewDir.dot(reflectDir), 0.0), shininess); - Vec3 specular = Vec3::init(1.0, 1.0, 1.0) * (specularStrength * spec); - return ambient + diffuse + specular; -} - -#endif // SHADING_H \ No newline at end of file +#endif // SHADING_H diff --git a/src/linalg/mat.cu b/src/linalg/mat.cu new file mode 100644 index 0000000..d5237e1 --- /dev/null +++ b/src/linalg/mat.cu @@ -0,0 +1,46 @@ +#include "mat.h" +#include +#include + +using namespace std; + +__device__ Vec3 computeGradient(float* volumeData, const int volW, const int volH, const int volD, int x, int y, int z) { + // Finite difference for partial derivatives. + // For boundary voxels - clamp to the boundary. + // Normal should point from higher to lower intensities + + int xm = max(x - 1, 0); + int xp = min(x + 1, volW - 1); + int ym = max(y - 1, 0); + int yp = min(y + 1, volH - 1); + int zm = max(z - 1, 0); + int zp = min(z + 1, volD - 1); + + // Note: Assuming data is linearized (idx = z*w*h + y*w + x) TODO: Unlinearize if data not linear + float gx = volumeData[z * volW * volH + y * volW + xp] + - volumeData[z * volW * volH + y * volW + xm]; + float gy = volumeData[z * volW * volH + yp * volW + x ] + - volumeData[z * volW * volH + ym * volW + x ]; + float gz = volumeData[zp * volW * volH + y * volW + x ] + - volumeData[zm * volW * volH + y * volW + x ]; + + return Vec3::init(gx, gy, gz); +}; + +// TESTING: haven't tested this function at all tbh +__device__ unsigned int packUnorm4x8(float r, float g, float b, float a) { + union { + unsigned char in[4]; + uint out; + } u; + + double len = sqrt(r*r + g*g + b*b + a*a); + + // This is a Vec4 but i can't be bothered to make that its own struct/class; FIXME: maybe do that if we need to? + std::vector v{r/len, g/len, b/len, a/len}; + for (int i = 0; i < v.size(); i++) { + u.in[i] = round(std::clamp(v[i], 0.0f, 1.0f) * 255.0f); + } + + return u.out; +} diff --git a/src/linalg/mat.h b/src/linalg/mat.h index 92f6fdc..cd070d7 100644 --- a/src/linalg/mat.h +++ b/src/linalg/mat.h @@ -1,24 +1,10 @@ -#pragma once +#ifndef MAT_H +#define MAT_H -__device__ Vec3 computeGradient(float* volumeData, const int volW, const int volH, const int volD, int x, int y, int z) { - // Finite difference for partial derivatives. - // For boundary voxels - clamp to the boundary. - // Normal should point from higher to lower intensities +#include "vec.h" - int xm = max(x - 1, 0); - int xp = min(x + 1, volW - 1); - int ym = max(y - 1, 0); - int yp = min(y + 1, volH - 1); - int zm = max(z - 1, 0); - int zp = min(z + 1, volD - 1); +__device__ Vec3 computeGradient(float* volumeData, const int volW, const int volH, const int volD, int x, int y, int z); - // Note: Assuming data is linearized (idx = z*w*h + y*w + x) TODO: Unlinearize if data not linear - float gx = volumeData[z * volW * volH + y * volW + xp] - - volumeData[z * volW * volH + y * volW + xm]; - float gy = volumeData[z * volW * volH + yp * volW + x ] - - volumeData[z * volW * volH + ym * volW + x ]; - float gz = volumeData[zp * volW * volH + y * volW + x ] - - volumeData[zm * volW * volH + y * volW + x ]; +__device__ unsigned int packUnorm4x8(float r, float g, float b, float a); - return Vec3::init(gx, gy, gz); -} +#endif // MAT_H diff --git a/src/linalg/vec.h b/src/linalg/vec.h index 9fc6745..f4bf332 100644 --- a/src/linalg/vec.h +++ b/src/linalg/vec.h @@ -27,4 +27,4 @@ struct Vec3 { // TODO: Maybe make this into a class ... maybe }; typedef Vec3 Point3; -typedef Vec3 Color3; \ No newline at end of file +typedef Vec3 Color3; diff --git a/src/main.cpp b/src/main.cpp new file mode 100644 index 0000000..60211cb --- /dev/null +++ b/src/main.cpp @@ -0,0 +1,8 @@ +#include "gui/MainWindow.h" +#include "consts.h" + + +int main() { + Window window(IMAGE_WIDTH, IMAGE_HEIGHT); + return window.init(); +} diff --git a/src/main.cu b/src/main.cu index 3c7b38a..4a0ceb4 100644 --- a/src/main.cu +++ b/src/main.cu @@ -7,7 +7,6 @@ #include "hurricanedata/datareader.h" #include "linalg/linalg.h" -#include "objs/sphere.h" #include "img/handler.h" #include "consts.h" #include "illumination/illumination.h" @@ -46,7 +45,8 @@ void getSpeed(std::vector& speedData, int idx = 0) { } } -int main(int argc, char** argv) { +// TODO: incorporate this main into main.cpp +int unmain(int argc, char** argv) { std::vector data; // getTemperature(data); getSpeed(data); @@ -83,16 +83,17 @@ int main(int argc, char** argv) { // Copy external constants from consts.h to cuda copyConstantsToDevice(); - // Launch kernel - dim3 blockSize(16, 16); // TODO: Figure out a good size for parallelization - dim3 gridSize((IMAGE_WIDTH + blockSize.x - 1)/blockSize.x, - (IMAGE_HEIGHT + blockSize.y - 1)/blockSize.y); - - raycastKernel<<>>( - d_volume, - d_framebuffer - ); - cudaDeviceSynchronize(); + // NOTE: this shold be done within the rayTracer class + // // Launch kernel + // dim3 blockSize(16, 16); + // dim3 gridSize((IMAGE_WIDTH + blockSize.x - 1)/blockSize.x, + // (IMAGE_HEIGHT + blockSize.y - 1)/blockSize.y); + // + // raycastKernel<<>>( + // d_volume, + // d_framebuffer + // ); + // cudaDeviceSynchronize(); // Copy framebuffer back to CPU unsigned char* hostFramebuffer = new unsigned char[IMAGE_WIDTH * IMAGE_HEIGHT * 3]; @@ -109,4 +110,4 @@ int main(int argc, char** argv) { std::cout << "Phong-DVR rendering done. Image saved to output.ppm" << std::endl; return 0; -} \ No newline at end of file +}