diff --git a/.gitignore b/.gitignore index 48c406d..698a114 100644 --- a/.gitignore +++ b/.gitignore @@ -11,3 +11,4 @@ data/ # Resulting images *.ppm + diff --git a/CMakeLists.txt b/CMakeLists.txt index b988528..4bd8a78 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -29,7 +29,7 @@ include_directories(${CMAKE_SOURCE_DIR}/include/imgui) add_executable(${PROJECT_NAME} ${HEADER_FILES} ${SOURCE_FILES} ${IMGUI_FILES}) # CUDA has specific architectures - set it to the system's architecture if available (or 70 by default) -set_target_properties(${PROJECT_NAME} PROPERTIES CUDA_ARCHITECTURES 70) +set_target_properties(${PROJECT_NAME} PROPERTIES CUDA_ARCHITECTURES 52) set_target_properties(${PROJECT_NAME} PROPERTIES CUDA_SEPARABLE_COMPILATION ON @@ -60,7 +60,7 @@ include_directories("${CUDA_INCLUDE_DIRS}") # netcdf find_package(netCDF REQUIRED) -message(STATUS "Found netcdf in ${NETCDF_INCLUDE_DIR}") +message(STATUS "Found netcdf in ${NETCDF_LIB}") execute_process( COMMAND nc-config --includedir @@ -74,9 +74,16 @@ execute_process( OUTPUT_STRIP_TRAILING_WHITESPACE ) +execute_process( + COMMAND ncxx4-config --prefix + OUTPUT_VARIABLE NETCDFCXX_PREFIX_DIR + OUTPUT_STRIP_TRAILING_WHITESPACE +) + target_include_directories(cuda-raytracer PUBLIC ${netCDF_INCLUDE_DIR}) -find_library(NETCDF_LIB NAMES netcdf-cxx4 netcdf_c++4 PATHS ${NETCDFCXX_LIB_DIR} NO_DEFAULT_PATH) +find_library(NETCDF_LIB NAMES netcdf-cxx4 netcdf_c++4 PATHS ${NETCDFCXX_LIB_DIR} ${NETCDFCXX_PREFIX_DIR}/lib NO_DEFAULT_PATH) +message(STATUS "Found NetCDFCXX in ${NETCDF_LIB}") set(LIBS ${GLFW3_LIBRARY} ${OPENGL_LIBRARY} GLAD ${CMAKE_DL_LIBS} ${CUDA_LIBRARIES} ${NETCDF_LIB}) diff --git a/default.nix b/default.nix new file mode 100644 index 0000000..e8ea029 --- /dev/null +++ b/default.nix @@ -0,0 +1,21 @@ +with import { config.allowUnfree = true; }; + +stdenv.mkDerivation { + name = "cuda-raytracer"; + src = ./.; + nativeBuildInputs = with cudaPackages; [cmake autoAddDriverRunpath autoPatchelfHook ]; + buildInputs = with pkgs; with cudaPackages; [ libGL glfw netcdf netcdfcxx4 cuda_nvcc cuda_cudart cuda_cccl libcurand libdecor]; + + postConfigure = '' + export netCDFCxx_DIR=${netcdfcxx4}/lib/cmake/netCDFCxx + ''; + + installTargets = "preinstall"; + + postInstall = '' + mkdir -p $out/bin + cp cuda-raytracer $out/bin + ''; + + LD_LIBRARY_PATH="/run/opengl-driver/lib/"; +} diff --git a/src/consts.cu b/src/consts.cu index 30afa91..415d3f4 100644 --- a/src/consts.cu +++ b/src/consts.cu @@ -1,16 +1,41 @@ #include "consts.h" +#include "cuda_error.h" // ----------------------- Colour mapping ----------------------- -__constant__ ColorStop d_stopsPythonLike[5]; -__constant__ ColorStop d_stopsGrayscale[2]; -__constant__ ColorStop d_stopsBluePurleRed[3]; +// __device__ __constant__ ColorStop d_stopsPythonLike[5]; +__device__ __constant__ ColorStop d_stopsPythonLike[11]; +__device__ __constant__ ColorStop d_stopsGrayscale[2]; +__device__ __constant__ ColorStop d_stopsBluePurleRed[3]; +// const ColorStop h_stopsPythonLike[] = { +// { 0.0f, Color3::init(0.2298057f, 0.29871797f, 0.75368315f) }, // Dark Blue +// { 0.25f, Color3::init(0.23437708f, 0.30554173f, 0.75967953f) }, // Mid Blue +// { 0.5f, Color3::init(0.27582712f, 0.36671692f, 0.81255294f) }, // White +// { 0.75f, Color3::init(0.79606387f, 0.84869321f, 0.93347147f) }, // Light Orange +// { 1.0f, Color3::init(0.70567316f, 0.01555616f, 0.15023281f) } // Red +// }; + +// const ColorStop h_stopsPythonLike[] = { +// { 0.0f, Color3::init(0.2298057f, 0.29871797f, 0.75368315f) }, // Dark Blue +// { 0.85f, Color3::init(0.23437708f, 0.30554173f, 0.75967953f) }, // Mid Blue +// { 0.90f, Color3::init(0.27582712f, 0.36671692f, 0.81255294f) }, // White +// { 0.95f, Color3::init(0.79606387f, 0.84869321f, 0.93347147f) }, // Light Orange +// { 1.0f, Color3::init(0.70567316f, 0.01555616f, 0.15023281f) } // Red +// }; + +// Python "jet" colour scheme const ColorStop h_stopsPythonLike[] = { - { 0.0f, Color3::init(0.2298057f, 0.29871797f, 0.75368315f) }, // Dark Blue - { 0.25f, Color3::init(0.23437708f, 0.30554173f, 0.75967953f) }, // Mid Blue - { 0.5f, Color3::init(0.27582712f, 0.36671692f, 0.81255294f) }, // White - { 0.75f, Color3::init(0.79606387f, 0.84869321f, 0.93347147f) }, // Light Orange - { 1.0f, Color3::init(0.70567316f, 0.01555616f, 0.15023281f) } // Red + { 0.00f, Color3::init(0.00000000f, 0.00000000f, 0.50000000f) }, + { 0.82f, Color3::init(0.00000000f, 0.00000000f, 0.94563280f) }, + { 0.84f, Color3::init(0.00000000f, 0.30000000f, 1.00000000f) }, + { 0.86f, Color3::init(0.00000000f, 0.69215686f, 1.00000000f) }, + { 0.88f, Color3::init(0.16129032f, 1.00000000f, 0.80645161f) }, + { 0.90f, Color3::init(0.49019608f, 1.00000000f, 0.47754586f) }, + { 0.92f, Color3::init(0.80645161f, 1.00000000f, 0.16129032f) }, + { 0.94f, Color3::init(1.00000000f, 0.77051561f, 0.00000000f) }, + { 0.96f, Color3::init(1.00000000f, 0.40740741f, 0.00000000f) }, + { 0.98f, Color3::init(0.94563280f, 0.02977487f, 0.00000000f) }, + { 1.00f, Color3::init(0.50000000f, 0.00000000f, 0.00000000f) }, }; const ColorStop h_stopsGrayscale[] = { @@ -37,8 +62,10 @@ Vec3 h_cameraUp = Vec3::init(0.0, 1.0, 0.0).normalize(); // Copy the above values to the device void copyConstantsToDevice() { + check_cuda_errors(cudaGetLastError()); // ----------------------- Colour mapping ----------------------- cudaMemcpyToSymbol(d_stopsPythonLike, h_stopsPythonLike, sizeof(h_stopsPythonLike)); + check_cuda_errors(cudaGetLastError()); cudaMemcpyToSymbol(d_stopsGrayscale, h_stopsGrayscale, sizeof(h_stopsGrayscale)); cudaMemcpyToSymbol(d_stopsBluePurleRed, h_stopsBluePurleRed, sizeof(h_stopsBluePurleRed)); diff --git a/src/consts.h b/src/consts.h index 3480c4a..4d2fd7b 100644 --- a/src/consts.h +++ b/src/consts.h @@ -13,13 +13,15 @@ const double infty = 1e15f; // This value is used to represent missing values // --------------------------- Dataset Constants --------------------------- // const int VOLUME_WIDTH = 576; // lon -const int VOLUME_WIDTH = 97; // lon +// const int VOLUME_WIDTH = 97; // lon +const int VOLUME_WIDTH = 57; // lon // const int VOLUME_HEIGHT = 361; // lat -const int VOLUME_HEIGHT = 71; // lat +// const int VOLUME_HEIGHT = 71; // lat +const int VOLUME_HEIGHT = 121; // lat const int VOLUME_DEPTH = 42; // lev -const float DLON = 60.0f / VOLUME_WIDTH; // 60 for current trimmed data set range -const float DLAT = 35.0f / VOLUME_HEIGHT; // 35 for current trimmed data set range +const float DLON = 35.0f / VOLUME_WIDTH; // 35 for current trimmed data set range +const float DLAT = 60.0f / VOLUME_HEIGHT; // 60 for current trimmed data set range const float DLEV = 1000.0f / VOLUME_DEPTH; // 1000 from max pressure (hPa) but not sure here const float MIN_TEMP = 210.0f; @@ -77,10 +79,10 @@ extern __device__ int d_samplesPerPixel; extern __device__ bool d_showSilhouettes; extern __device__ float d_silhouettesThreshold; -const int lenStopsPythonLike = 5; +const int lenStopsPythonLike = 11; const int lenStopsGrayscale = 2; const int lenStopsBluePurpleRed = 3; -extern __constant__ ColorStop d_stopsPythonLike[5]; +extern __constant__ ColorStop d_stopsPythonLike[11]; extern __constant__ ColorStop d_stopsGrayscale[2]; extern __constant__ ColorStop d_stopsBluePurleRed[3]; diff --git a/src/gui/MainWindow.cpp b/src/gui/MainWindow.cpp index be75dea..43399ec 100644 --- a/src/gui/MainWindow.cpp +++ b/src/gui/MainWindow.cpp @@ -7,6 +7,7 @@ #include "Shader.h" #include "input/Widget.h" +#include "cuda_error.h" Window::Window(unsigned int w, unsigned int h) { @@ -98,6 +99,7 @@ void Window::free(float* data) { void Window::tick() { + std::cout << "tick tick\n"; // 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(); @@ -111,6 +113,7 @@ void Window::tick() { this->widget->dateChanged = false; } + std::cout << "setting tick render\n"; // tick render glBindFramebuffer(GL_FRAMEBUFFER, 0); glDisable(GL_DEPTH_TEST); @@ -138,11 +141,14 @@ void Window::tick() { glDrawArrays(GL_TRIANGLES, 0, 6); // draw current frame to texture // render ImGui context + std::cout << "rendering\n"; this->widget->render(); + std::cout << "done rendering\n"; // check for events glfwSwapBuffers(this->window); glfwPollEvents(); + std::cout << "done tick tick\n"; } diff --git a/src/gui/Quad.cpp b/src/gui/Quad.cpp index d918763..5451a71 100644 --- a/src/gui/Quad.cpp +++ b/src/gui/Quad.cpp @@ -83,6 +83,7 @@ void Quad::cuda_init(float* data) { void Quad::render() { + check_cuda_errors(cudaGetLastError()); glBindTexture(GL_TEXTURE_2D, 0); this->renderer->render(); glBindTexture(GL_TEXTURE_2D, this->tex); diff --git a/src/gui/input/Widget.cpp b/src/gui/input/Widget.cpp index 8af0b98..8eeb371 100644 --- a/src/gui/input/Widget.cpp +++ b/src/gui/input/Widget.cpp @@ -4,6 +4,7 @@ #include #include #include +#include // probably not the cleanest way to do this but it works void parseDate(char* string, int dayOfYear) { @@ -53,7 +54,7 @@ Widget::Widget(GLFWwindow* window) : opacityK(0), sigmoidShift(0.5f), sigmoidExp(-250.0f), - tfComboSelected(0), + tfComboSelected(2), dateChanged(false), paused(true), renderOnce(false), @@ -62,8 +63,8 @@ Widget::Widget(GLFWwindow* window) : samplesPerPixel(1), alphaAcumLimit(0.4f), opacityConst(100), - showSilhouettes(true), - silhouettesThreshold(0.2f) + showSilhouettes(false), + silhouettesThreshold(0.02f) { IMGUI_CHECKVERSION(); ImGui::CreateContext(); @@ -72,6 +73,15 @@ Widget::Widget(GLFWwindow* window) : ImGui_ImplGlfw_InitForOpenGL(window, true); ImGui_ImplOpenGL3_Init(); + this->cameraPos = Point3::init(59.0f, 77.5f, -18.0f); // Camera for partially trimmed data set + this->pitch = 0.7; + this->yaw = 4.85; + this->roll = -0.010; + this->cameraDir = Vec3::getDirectionFromEuler(pitch, yaw, roll); + + this->bgColor = Color3::init(0.1f, 0.1f, 0.1f); + this->lightPos = Point3::init(1.5, 2.0, -1.0); + this->fps = (char*)malloc(512*sizeof(char)); this->dateString = (char*)malloc(512*sizeof(char)); sprintf(this->dateString, "January 1st"); @@ -84,6 +94,7 @@ Widget::Widget(GLFWwindow* window) : // TODO: can be marginally improved by only copying changed values to device - however we're dealing with individual floats here so i dont think the benefit would be all that obvious. // TODO: wrap basically all ImGui calls in if statements; better form + allows for checking return values / errors. void Widget::tick(double fps) { + if (this->renderOnce) { this->renderOnce = false; this->paused = true; @@ -169,12 +180,23 @@ void Widget::tick(double fps) { ImGui::DragScalar("X position", ImGuiDataType_Double, &this->cameraPos.x, 0.5f, &min, &max, "%.3f"); ImGui::DragScalar("Y position", ImGuiDataType_Double, &this->cameraPos.z, 0.5f, &min, &max, "%.3f"); ImGui::DragScalar("Z position", ImGuiDataType_Double, &this->cameraPos.y, 0.5f, &min, &max, "%.3f"); - ImGui::DragScalar("X direction", ImGuiDataType_Double, &this->cameraDir.x, 0.005f, &min, &max, "%.3f"); // TODO: should wrap around once fully rotated - ImGui::DragScalar("Y direction", ImGuiDataType_Double, &this->cameraDir.z, 0.005f, &min, &max, "%.3f"); - ImGui::DragScalar("Z direction", ImGuiDataType_Double, &this->cameraDir.y, 0.005f, &min, &max, "%.3f"); + // ImGui::DragScalar("X direction", ImGuiDataType_Double, &this->cameraDir.x, 0.005f, &min, &max, "%.3f"); // TODO: should wrap around once fully rotated + // ImGui::DragScalar("Y direction", ImGuiDataType_Double, &this->cameraDir.z, 0.005f, &min, &max, "%.3f"); + // ImGui::DragScalar("Z direction", ImGuiDataType_Double, &this->cameraDir.y, 0.005f, &min, &max, "%.3f"); + ImGui::DragScalar("Pitch", ImGuiDataType_Double, &this->pitch, 0.005f, &min, &max, "%.3f"); + ImGui::DragScalar("Yaw", ImGuiDataType_Double, &this->yaw, 0.005f, &min, &max, "%.3f"); + ImGui::DragScalar("Roll", ImGuiDataType_Double, &this->roll, 0.005f, &min, &max, "%.3f"); if (ImGui::Button("Reset camera")) resetCamera(); ImGui::End(); - + + this->cameraDir.setDirectionFromEuler(pitch, yaw, roll); + + // Calculate upCamera + Vec3 arbitraryVector = Vec3::init(1, 0, 0); + cameraUp = arbitraryVector.cross(cameraDir); + cameraUp.normalize(); + cameraUp.rotateAroundAxis(cameraDir, roll); + copyToDevice(); } @@ -198,6 +220,7 @@ void Widget::resetLight() { void Widget::copyToDevice() { cudaMemcpyToSymbol(&d_cameraPos, &this->cameraPos, sizeof(Point3)); + cudaMemcpyToSymbol(&d_cameraUp, &this->cameraUp, sizeof(Point3)); cudaMemcpyToSymbol(&d_cameraDir, &this->cameraDir, sizeof(Vec3)); cudaMemcpyToSymbol(&d_lightPos, &this->lightPos, sizeof(Point3)); cudaMemcpyToSymbol(&d_backgroundColor, &this->bgColor, sizeof(Color3)); diff --git a/src/gui/input/Widget.h b/src/gui/input/Widget.h index aec0fbc..58e69e7 100644 --- a/src/gui/input/Widget.h +++ b/src/gui/input/Widget.h @@ -24,6 +24,8 @@ private: // camera controls Point3 cameraDir; Vec3 cameraPos; + double pitch, yaw, roll; + Vec3 cameraUp; // simulation cotnrols bool renderOnce; diff --git a/src/illumination/Raycaster.cu b/src/illumination/Raycaster.cu index f66a9d9..8086e38 100644 --- a/src/illumination/Raycaster.cu +++ b/src/illumination/Raycaster.cu @@ -38,7 +38,7 @@ __global__ void raycastKernel(float* volumeData, FrameBuffer framebuffer, const u *= tanHalfFov; v *= tanHalfFov; - // Find ray direction + Vec3 cameraRight = (d_cameraDir.cross(d_cameraUp)).normalize(); d_cameraUp = (cameraRight.cross(d_cameraDir)).normalize(); Vec3 rayDir = (d_cameraDir + cameraRight*u + d_cameraUp*v).normalize(); @@ -156,6 +156,7 @@ Raycaster::Raycaster(cudaGraphicsResource_t resources, int w, int h, float* data void Raycaster::render() { + check_cuda_errors(cudaGetLastError()); check_cuda_errors(cudaGraphicsMapResources(1, &this->resources)); check_cuda_errors(cudaGraphicsResourceGetMappedPointer((void**)&(this->fb->buffer), &(this->fb->buffer_size), resources)); diff --git a/src/illumination/transferFunction.cu b/src/illumination/transferFunction.cu index a6b9087..6fc6fc8 100644 --- a/src/illumination/transferFunction.cu +++ b/src/illumination/transferFunction.cu @@ -111,7 +111,8 @@ __device__ float4 transferFunction(float density, const Vec3& grad, const Point3 result.x = 0.0f; result.y = 0.0f; result.z = 0.0f; - result.w = alpha; + // result.w = alpha; + result.w = d_opacityConst; } return result; diff --git a/src/linalg/vec.h b/src/linalg/vec.h index a243cb9..a017ba1 100644 --- a/src/linalg/vec.h +++ b/src/linalg/vec.h @@ -31,6 +31,33 @@ struct Vec3 { __host__ __device__ Vec3 cross(const Vec3& b) const { return Vec3::init(y * b.z - z * b.y, z * b.x - x * b.z, x * b.y - y * b.x); } __host__ __device__ Vec3 normalize() const { double len = sqrt(x * x + y * y + z * z); return Vec3::init(x / len, y / len, z / len); } __host__ __device__ double length() const { return sqrt(x * x + y * y + z * z); } + + __host__ __device__ void setDirectionFromEuler(double pitch, double yaw, double roll) { + // Compute the direction vector using the Euler angles in radians + double cosPitch = cos(pitch); + double sinPitch = sin(pitch); + double cosYaw = cos(yaw); + double sinYaw = sin(yaw); + + // Direction vector components + x = cosPitch * cosYaw; + y = cosPitch * sinYaw; + z = sinPitch; + } + + static __host__ __device__ Vec3 getDirectionFromEuler(double pitch, double yaw, double roll) { + Vec3 v = Vec3::init(1,0,0); + v.setDirectionFromEuler(pitch, yaw, roll); + return v; + } + + __host__ __device__ void rotateAroundAxis(const Vec3& axis, double angle) { + double cosA = cos(angle); + double sinA = sin(angle); + + Vec3 rotated = *this * cosA + axis.cross(*this) * sinA + axis * axis.dot(*this) * (1 - cosA); + *this = rotated; + } }; typedef Vec3 Point3; diff --git a/src/main.cu b/src/main.cu index aec5022..d9b7108 100644 --- a/src/main.cu +++ b/src/main.cu @@ -21,7 +21,7 @@ static float* d_volume = nullptr; // * time controls - arbitrary skipping to specified point (would require some changes to gpubuffer) (could have) void getTemperature(std::vector& temperatureData, int idx = 0) { - std::string path = "data/trimmed"; + std::string path = "../../../../hurricane-sandy-data/trimmed"; // std::string path = "data"; std::string variable = "T"; DataReader dataReader(path, variable); @@ -31,7 +31,7 @@ void getTemperature(std::vector& temperatureData, int idx = 0) { } void getSpeed(std::vector& speedData, int idx = 0) { - std::string path = "data/trimmed"; + std::string path = "../../../../hurricane-sandy-data/trimmed"; // std::string path = "data"; std::string varU = "U"; std::string varV = "V"; @@ -55,7 +55,7 @@ void getSpeed(std::vector& speedData, int idx = 0) { int main() { std::vector data; - getTemperature(data, 0); + getTemperature(data, 7); // 20121028 // getSpeed(data, 294); std::cout << "DATA size: " << data.size() << std::endl;