djairoh
/
cuda-based-raytrace
mirror of https://github.com/MartinOpat/cuda-based-raytrace.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #2 from MartinOpat/gui 

Gui
This commit is contained in:
Djairo 2025-01-08 15:24:28 +00:00 committed by GitHub
parent c92d0bc2a9 b3db9a7045
commit 6f22230ead
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
31 changed files with 12993 additions and 133 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

101
CMakeLists.txt Normal file
View File

@ -0,0 +1,101 @@
cmake_minimum_required (VERSION 3.8)
project ("cuda-raytracer" LANGUAGES CUDA CXX C)
# get the shader files in the build directory
file(COPY ${CMAKE_SOURCE_DIR}/src/gui/shaders DESTINATION .)
set(CMAKE_BUILD_TYPE Debug)
# source files
file(GLOB_RECURSE SOURCE_FILES
${CMAKE_SOURCE_DIR}/src/*.c
${CMAKE_SOURCE_DIR}/src/*.cpp
${CMAKE_SOURCE_DIR}/src/*.cu)
# header files
file(GLOB_RECURSE HEADER_FILES
${CMAKE_SOURCE_DIR}/src/*.h
${CMAKE_SOURCE_DIR}/src/*.hpp)
add_executable(${PROJECT_NAME} ${HEADER_FILES} ${SOURCE_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_SEPARABLE_COMPILATION ON
CUDA_RESOLVE_DEVICE_SYMBOLS ON
)
# Add .lib files
link_directories(${CMAKE_SOURCE_DIR}/lib)
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/")
# Package management
# OpenGL
set(OpenGL_GL_PREFERENCE GLVND)
find_package(OpenGL REQUIRED)
# GLFW
find_package(GLFW3 REQUIRED)
message(STATUS "Found GLFW3 in ${GLFW3_INCLUDE_DIR}")
# GLAD
add_library(GLAD "thirdparty/glad.c")
# CUDA
find_package(CUDA REQUIRED)
include_directories("${CUDA_INCLUDE_DIRS}")
# netcdf
find_package(netCDF REQUIRED)
message(STATUS "Found netcdf in ${GLFW3_INCLUDE_DIR}")
execute_process(
COMMAND nc-config --includedir
OUTPUT_VARIABLE NETCDF_INCLUDE_DIR
OUTPUT_STRIP_TRAILING_WHITESPACE
)
execute_process(
COMMAND ncxx4-config --libdir
OUTPUT_VARIABLE NETCDFCXX_LIB_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)
set(LIBS ${GLFW3_LIBRARY} ${OPENGL_LIBRARY} GLAD ${CMAKE_DL_LIBS} ${CUDA_LIBRARIES} ${NETCDF_LIB})
include_directories(
"${CMAKE_SOURCE_DIR}/src"
"${CMAKE_SOURCE_DIR}/include"
"${CMAKE_SOURCE_DIR}/include/glad"
)
target_link_libraries(${PROJECT_NAME} ${LIBS})
function(CUDA_CONVERT_FLAGS EXISTING_TARGET)
get_property(old_flags TARGET ${EXISTING_TARGET} PROPERTY INTERFACE_COMPILE_OPTIONS)
if(NOT "${old_flags}" STREQUAL "")
string(REPLACE ";" "," CUDA_flags "${old_flags}")
set_property(TARGET ${EXISTING_TARGET} PROPERTY INTERFACE_COMPILE_OPTIONS
"$<$<BUILD_INTERFACE:$<COMPILE_LANGUAGE:CXX>>:${old_flags}>$<$<BUILD_INTERFACE:$<COMPILE_LANGUAGE:CUDA>>:-Xcompiler=${CUDA_flags}>"
)
endif()
endfunction()
CUDA_CONVERT_FLAGS(${PROJECT_NAME})
# check for cache variable set in cmakepresets called IS_CUDA_DEBUG
if(IS_CUDA_DEBUG)
target_compile_options(${PROJECT_NAME} PRIVATE "$<$<AND:$<CONFIG:Debug>,$<COMPILE_LANGUAGE:CUDA>>:-G>")
message(STATUS "CUDA_DEBUG is ON")
endif()

5912
GLFW/glfw3.h Normal file
View File

File diff suppressed because it is too large Load Diff

628
GLFW/glfw3native.h Normal file
View File

@ -0,0 +1,628 @@
/*************************************************************************
* GLFW 3.3 - www.glfw.org
* A library for OpenGL, window and input
*------------------------------------------------------------------------
* Copyright (c) 2002-2006 Marcus Geelnard
* Copyright (c) 2006-2018 Camilla Löwy <elmindreda@glfw.org>
*
* This software is provided 'as-is', without any express or implied
* warranty. In no event will the authors be held liable for any damages
* arising from the use of this software.
*
* Permission is granted to anyone to use this software for any purpose,
* including commercial applications, and to alter it and redistribute it
* freely, subject to the following restrictions:
*
* 1. The origin of this software must not be misrepresented; you must not
* claim that you wrote the original software. If you use this software
* in a product, an acknowledgment in the product documentation would
* be appreciated but is not required.
*
* 2. Altered source versions must be plainly marked as such, and must not
* be misrepresented as being the original software.
*
* 3. This notice may not be removed or altered from any source
* distribution.
*
*************************************************************************/
#ifndef _glfw3_native_h_
#define _glfw3_native_h_
#ifdef __cplusplus
extern "C" {
#endif
/*************************************************************************
* Doxygen documentation
*************************************************************************/
/*! @file glfw3native.h
* @brief The header of the native access functions.
*
* This is the header file of the native access functions. See @ref native for
* more information.
*/
/*! @defgroup native Native access
* @brief Functions related to accessing native handles.
*
* **By using the native access functions you assert that you know what you're
* doing and how to fix problems caused by using them. If you don't, you
* shouldn't be using them.**
*
* Before the inclusion of @ref glfw3native.h, you may define zero or more
* window system API macro and zero or more context creation API macros.
*
* The chosen backends must match those the library was compiled for. Failure
* to do this will cause a link-time error.
*
* The available window API macros are:
* * `GLFW_EXPOSE_NATIVE_WIN32`
* * `GLFW_EXPOSE_NATIVE_COCOA`
* * `GLFW_EXPOSE_NATIVE_X11`
* * `GLFW_EXPOSE_NATIVE_WAYLAND`
*
* The available context API macros are:
* * `GLFW_EXPOSE_NATIVE_WGL`
* * `GLFW_EXPOSE_NATIVE_NSGL`
* * `GLFW_EXPOSE_NATIVE_GLX`
* * `GLFW_EXPOSE_NATIVE_EGL`
* * `GLFW_EXPOSE_NATIVE_OSMESA`
*
* These macros select which of the native access functions that are declared
* and which platform-specific headers to include. It is then up your (by
* definition platform-specific) code to handle which of these should be
* defined.
*
* If you do not want the platform-specific headers to be included, define
* `GLFW_NATIVE_INCLUDE_NONE` before including the @ref glfw3native.h header.
*
* @code
* #define GLFW_EXPOSE_NATIVE_WIN32
* #define GLFW_EXPOSE_NATIVE_WGL
* #define GLFW_NATIVE_INCLUDE_NONE
* #include <GLFW/glfw3native.h>
* @endcode
*/
/*************************************************************************
* System headers and types
*************************************************************************/
#if !defined(GLFW_NATIVE_INCLUDE_NONE)
#if defined(GLFW_EXPOSE_NATIVE_WIN32) || defined(GLFW_EXPOSE_NATIVE_WGL)
/* This is a workaround for the fact that glfw3.h needs to export APIENTRY (for
* example to allow applications to correctly declare a GL_KHR_debug callback)
* but windows.h assumes no one will define APIENTRY before it does
*/
#if defined(GLFW_APIENTRY_DEFINED)
#undef APIENTRY
#undef GLFW_APIENTRY_DEFINED
#endif
#include <windows.h>
#elif defined(GLFW_EXPOSE_NATIVE_COCOA) || defined(GLFW_EXPOSE_NATIVE_NSGL)
#if defined(__OBJC__)
#import <Cocoa/Cocoa.h>
#else
#include <ApplicationServices/ApplicationServices.h>
#include <objc/objc.h>
#endif
#elif defined(GLFW_EXPOSE_NATIVE_X11) || defined(GLFW_EXPOSE_NATIVE_GLX)
#include <X11/Xlib.h>
#include <X11/extensions/Xrandr.h>
#elif defined(GLFW_EXPOSE_NATIVE_WAYLAND)
#include <wayland-client.h>
#endif
#if defined(GLFW_EXPOSE_NATIVE_WGL)
/* WGL is declared by windows.h */
#endif
#if defined(GLFW_EXPOSE_NATIVE_NSGL)
/* NSGL is declared by Cocoa.h */
#endif
#if defined(GLFW_EXPOSE_NATIVE_GLX)
/* This is a workaround for the fact that glfw3.h defines GLAPIENTRY because by
* default it also acts as an OpenGL header
* However, glx.h will include gl.h, which will define it unconditionally
*/
#if defined(GLFW_GLAPIENTRY_DEFINED)
#undef GLAPIENTRY
#undef GLFW_GLAPIENTRY_DEFINED
#endif
#include <GL/glx.h>
#endif
#if defined(GLFW_EXPOSE_NATIVE_EGL)
#include <EGL/egl.h>
#endif
#if defined(GLFW_EXPOSE_NATIVE_OSMESA)
/* This is a workaround for the fact that glfw3.h defines GLAPIENTRY because by
* default it also acts as an OpenGL header
* However, osmesa.h will include gl.h, which will define it unconditionally
*/
#if defined(GLFW_GLAPIENTRY_DEFINED)
#undef GLAPIENTRY
#undef GLFW_GLAPIENTRY_DEFINED
#endif
#include <GL/osmesa.h>
#endif
#endif /*GLFW_NATIVE_INCLUDE_NONE*/
/*************************************************************************
* Functions
*************************************************************************/
#if defined(GLFW_EXPOSE_NATIVE_WIN32)
/*! @brief Returns the adapter device name of the specified monitor.
*
* @return The UTF-8 encoded adapter device name (for example `\\.\DISPLAY1`)
* of the specified monitor, or `NULL` if an [error](@ref error_handling)
* occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.1.
*
* @ingroup native
*/
GLFWAPI const char* glfwGetWin32Adapter(GLFWmonitor* monitor);
/*! @brief Returns the display device name of the specified monitor.
*
* @return The UTF-8 encoded display device name (for example
* `\\.\DISPLAY1\Monitor0`) of the specified monitor, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.1.
*
* @ingroup native
*/
GLFWAPI const char* glfwGetWin32Monitor(GLFWmonitor* monitor);
/*! @brief Returns the `HWND` of the specified window.
*
* @return The `HWND` of the specified window, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @remark The `HDC` associated with the window can be queried with the
* [GetDC](https://docs.microsoft.com/en-us/windows/win32/api/winuser/nf-winuser-getdc)
* function.
* @code
* HDC dc = GetDC(glfwGetWin32Window(window));
* @endcode
* This DC is private and does not need to be released.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI HWND glfwGetWin32Window(GLFWwindow* window);
#endif
#if defined(GLFW_EXPOSE_NATIVE_WGL)
/*! @brief Returns the `HGLRC` of the specified window.
*
* @return The `HGLRC` of the specified window, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @remark The `HDC` associated with the window can be queried with the
* [GetDC](https://docs.microsoft.com/en-us/windows/win32/api/winuser/nf-winuser-getdc)
* function.
* @code
* HDC dc = GetDC(glfwGetWin32Window(window));
* @endcode
* This DC is private and does not need to be released.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI HGLRC glfwGetWGLContext(GLFWwindow* window);
#endif
#if defined(GLFW_EXPOSE_NATIVE_COCOA)
/*! @brief Returns the `CGDirectDisplayID` of the specified monitor.
*
* @return The `CGDirectDisplayID` of the specified monitor, or
* `kCGNullDirectDisplay` if an [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.1.
*
* @ingroup native
*/
GLFWAPI CGDirectDisplayID glfwGetCocoaMonitor(GLFWmonitor* monitor);
/*! @brief Returns the `NSWindow` of the specified window.
*
* @return The `NSWindow` of the specified window, or `nil` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI id glfwGetCocoaWindow(GLFWwindow* window);
#endif
#if defined(GLFW_EXPOSE_NATIVE_NSGL)
/*! @brief Returns the `NSOpenGLContext` of the specified window.
*
* @return The `NSOpenGLContext` of the specified window, or `nil` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI id glfwGetNSGLContext(GLFWwindow* window);
#endif
#if defined(GLFW_EXPOSE_NATIVE_X11)
/*! @brief Returns the `Display` used by GLFW.
*
* @return The `Display` used by GLFW, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI Display* glfwGetX11Display(void);
/*! @brief Returns the `RRCrtc` of the specified monitor.
*
* @return The `RRCrtc` of the specified monitor, or `None` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.1.
*
* @ingroup native
*/
GLFWAPI RRCrtc glfwGetX11Adapter(GLFWmonitor* monitor);
/*! @brief Returns the `RROutput` of the specified monitor.
*
* @return The `RROutput` of the specified monitor, or `None` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.1.
*
* @ingroup native
*/
GLFWAPI RROutput glfwGetX11Monitor(GLFWmonitor* monitor);
/*! @brief Returns the `Window` of the specified window.
*
* @return The `Window` of the specified window, or `None` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI Window glfwGetX11Window(GLFWwindow* window);
/*! @brief Sets the current primary selection to the specified string.
*
* @param[in] string A UTF-8 encoded string.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED and @ref
* GLFW_PLATFORM_ERROR.
*
* @pointer_lifetime The specified string is copied before this function
* returns.
*
* @thread_safety This function must only be called from the main thread.
*
* @sa @ref clipboard
* @sa glfwGetX11SelectionString
* @sa glfwSetClipboardString
*
* @since Added in version 3.3.
*
* @ingroup native
*/
GLFWAPI void glfwSetX11SelectionString(const char* string);
/*! @brief Returns the contents of the current primary selection as a string.
*
* If the selection is empty or if its contents cannot be converted, `NULL`
* is returned and a @ref GLFW_FORMAT_UNAVAILABLE error is generated.
*
* @return The contents of the selection as a UTF-8 encoded string, or `NULL`
* if an [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED and @ref
* GLFW_PLATFORM_ERROR.
*
* @pointer_lifetime The returned string is allocated and freed by GLFW. You
* should not free it yourself. It is valid until the next call to @ref
* glfwGetX11SelectionString or @ref glfwSetX11SelectionString, or until the
* library is terminated.
*
* @thread_safety This function must only be called from the main thread.
*
* @sa @ref clipboard
* @sa glfwSetX11SelectionString
* @sa glfwGetClipboardString
*
* @since Added in version 3.3.
*
* @ingroup native
*/
GLFWAPI const char* glfwGetX11SelectionString(void);
#endif
#if defined(GLFW_EXPOSE_NATIVE_GLX)
/*! @brief Returns the `GLXContext` of the specified window.
*
* @return The `GLXContext` of the specified window, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI GLXContext glfwGetGLXContext(GLFWwindow* window);
/*! @brief Returns the `GLXWindow` of the specified window.
*
* @return The `GLXWindow` of the specified window, or `None` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.2.
*
* @ingroup native
*/
GLFWAPI GLXWindow glfwGetGLXWindow(GLFWwindow* window);
#endif
#if defined(GLFW_EXPOSE_NATIVE_WAYLAND)
/*! @brief Returns the `struct wl_display*` used by GLFW.
*
* @return The `struct wl_display*` used by GLFW, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.2.
*
* @ingroup native
*/
GLFWAPI struct wl_display* glfwGetWaylandDisplay(void);
/*! @brief Returns the `struct wl_output*` of the specified monitor.
*
* @return The `struct wl_output*` of the specified monitor, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.2.
*
* @ingroup native
*/
GLFWAPI struct wl_output* glfwGetWaylandMonitor(GLFWmonitor* monitor);
/*! @brief Returns the main `struct wl_surface*` of the specified window.
*
* @return The main `struct wl_surface*` of the specified window, or `NULL` if
* an [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.2.
*
* @ingroup native
*/
GLFWAPI struct wl_surface* glfwGetWaylandWindow(GLFWwindow* window);
#endif
#if defined(GLFW_EXPOSE_NATIVE_EGL)
/*! @brief Returns the `EGLDisplay` used by GLFW.
*
* @return The `EGLDisplay` used by GLFW, or `EGL_NO_DISPLAY` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NOT_INITIALIZED.
*
* @remark Because EGL is initialized on demand, this function will return
* `EGL_NO_DISPLAY` until the first context has been created via EGL.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI EGLDisplay glfwGetEGLDisplay(void);
/*! @brief Returns the `EGLContext` of the specified window.
*
* @return The `EGLContext` of the specified window, or `EGL_NO_CONTEXT` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI EGLContext glfwGetEGLContext(GLFWwindow* window);
/*! @brief Returns the `EGLSurface` of the specified window.
*
* @return The `EGLSurface` of the specified window, or `EGL_NO_SURFACE` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.0.
*
* @ingroup native
*/
GLFWAPI EGLSurface glfwGetEGLSurface(GLFWwindow* window);
#endif
#if defined(GLFW_EXPOSE_NATIVE_OSMESA)
/*! @brief Retrieves the color buffer associated with the specified window.
*
* @param[in] window The window whose color buffer to retrieve.
* @param[out] width Where to store the width of the color buffer, or `NULL`.
* @param[out] height Where to store the height of the color buffer, or `NULL`.
* @param[out] format Where to store the OSMesa pixel format of the color
* buffer, or `NULL`.
* @param[out] buffer Where to store the address of the color buffer, or
* `NULL`.
* @return `GLFW_TRUE` if successful, or `GLFW_FALSE` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.3.
*
* @ingroup native
*/
GLFWAPI int glfwGetOSMesaColorBuffer(GLFWwindow* window, int* width, int* height, int* format, void** buffer);
/*! @brief Retrieves the depth buffer associated with the specified window.
*
* @param[in] window The window whose depth buffer to retrieve.
* @param[out] width Where to store the width of the depth buffer, or `NULL`.
* @param[out] height Where to store the height of the depth buffer, or `NULL`.
* @param[out] bytesPerValue Where to store the number of bytes per depth
* buffer element, or `NULL`.
* @param[out] buffer Where to store the address of the depth buffer, or
* `NULL`.
* @return `GLFW_TRUE` if successful, or `GLFW_FALSE` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.3.
*
* @ingroup native
*/
GLFWAPI int glfwGetOSMesaDepthBuffer(GLFWwindow* window, int* width, int* height, int* bytesPerValue, void** buffer);
/*! @brief Returns the `OSMesaContext` of the specified window.
*
* @return The `OSMesaContext` of the specified window, or `NULL` if an
* [error](@ref error_handling) occurred.
*
* @errors Possible errors include @ref GLFW_NO_WINDOW_CONTEXT and @ref
* GLFW_NOT_INITIALIZED.
*
* @thread_safety This function may be called from any thread. Access is not
* synchronized.
*
* @since Added in version 3.3.
*
* @ingroup native
*/
GLFWAPI OSMesaContext glfwGetOSMesaContext(GLFWwindow* window);
#endif
#ifdef __cplusplus
}
#endif
#endif /* _glfw3_native_h_ */

22
README.md
View File

@ -1,22 +1,12 @@
# cuda-based-raytrace # cuda-based-raytrace
## How to run ## How to run
If necessary clean previous build:
```bash
make clean
```
Then build (currently always debug build): Gui branch was a bit easier to compile using Cmake so to run this branch do:
```bash ```bash
make all mkdir build
cd build
cmake ..
make
./cuda-raytracer
``` ```
Run:
```bash
./build/main
```
To open the resulting image you can use (Ubuntu):
```bash
xdg-open output.ppm
```

49
cmake/FindGLFW3.cmake Normal file
View File

@ -0,0 +1,49 @@
# Locate the glfw3 library
#
# This module defines the following variables:
#
# GLFW3_LIBRARY the name of the library;
# GLFW3_INCLUDE_DIR where to find glfw include files.
# GLFW3_FOUND true if both the GLFW3_LIBRARY and GLFW3_INCLUDE_DIR have been found.
#
# To help locate the library and include file, you can define a
# variable called GLFW3_ROOT which points to the root of the glfw library
# installation.
#
# default search dirs
#
# Cmake file from: https://github.com/daw42/glslcookbook
set( _glfw3_HEADER_SEARCH_DIRS
"/usr/include"
"/usr/local/include"
"${CMAKE_SOURCE_DIR}/include"
"C:/Program Files (x86)/glfw/include" )
set( _glfw3_LIB_SEARCH_DIRS
"/usr/lib"
"/usr/local/lib"
"${CMAKE_SOURCE_DIR}/lib"
"C:/Program Files (x86)/glfw/lib-msvc110" )
# Check environment for root search directory
set( _glfw3_ENV_ROOT $ENV{GLFW3_ROOT} )
if( NOT GLFW3_ROOT AND _glfw3_ENV_ROOT )
set(GLFW3_ROOT ${_glfw3_ENV_ROOT} )
endif()
# Put user specified location at beginning of search
if( GLFW3_ROOT )
list( INSERT _glfw3_HEADER_SEARCH_DIRS 0 "${GLFW3_ROOT}/include" )
list( INSERT _glfw3_LIB_SEARCH_DIRS 0 "${GLFW3_ROOT}/lib" )
endif()
# Search for the header
FIND_PATH(GLFW3_INCLUDE_DIR "GLFW/glfw3.h"
PATHS ${_glfw3_HEADER_SEARCH_DIRS} )
# Search for the library
FIND_LIBRARY(GLFW3_LIBRARY NAMES glfw3 glfw
PATHS ${_glfw3_LIB_SEARCH_DIRS} )
INCLUDE(FindPackageHandleStandardArgs)
FIND_PACKAGE_HANDLE_STANDARD_ARGS(GLFW3 DEFAULT_MSG
GLFW3_LIBRARY GLFW3_INCLUDE_DIR)

3694
include/glad/glad.h Normal file
View File

File diff suppressed because it is too large Load Diff

BIN
lib/glfw3.lib Normal file
View File

Binary file not shown.

4
makefile
View File

@ -1,6 +1,6 @@
# Compiler and flags # Compiler and flags
NVCC = nvcc NVCC = nvcc
CXXFLAGS = -I./src -std=c++17 $(shell nc-config --cxx4flags) $(shell nc-config --cxx4libs) -g -G CXXFLAGS = -I./src -std=c++17 $(shell ncxx4-config --cflags) $(shell ncxx4-config --libs) -g -G
COMPILE_OBJ_FLAGS = --device-c COMPILE_OBJ_FLAGS = --device-c
# Directories # Directories
@ -36,4 +36,4 @@ clean:
$(BUILD_DIR): $(BUILD_DIR):
mkdir -p $(BUILD_DIR) mkdir -p $(BUILD_DIR)
.PHONY: all clean debug .PHONY: all clean debug

2
src/consts.cu
View File

@ -15,4 +15,4 @@ void copyConstantsToDevice() {
cudaMemcpyToSymbol(d_cameraDir, &h_cameraDir, sizeof(Vec3)); cudaMemcpyToSymbol(d_cameraDir, &h_cameraDir, sizeof(Vec3));
cudaMemcpyToSymbol(d_cameraUp, &h_cameraUp, sizeof(Vec3)); cudaMemcpyToSymbol(d_cameraUp, &h_cameraUp, sizeof(Vec3));
cudaMemcpyToSymbol(d_lightPos, &h_lightPos, sizeof(Point3)); cudaMemcpyToSymbol(d_lightPos, &h_lightPos, sizeof(Point3));
} }

8
src/consts.h
View File

@ -9,11 +9,11 @@ const int VOLUME_WIDTH = 49;
const int VOLUME_HEIGHT = 51; const int VOLUME_HEIGHT = 51;
const int VOLUME_DEPTH = 42; const int VOLUME_DEPTH = 42;
const int IMAGE_WIDTH = 2560; const int IMAGE_WIDTH = 800;
const int IMAGE_HEIGHT = 1440; const int IMAGE_HEIGHT = 600;
const double epsilon = 1e-10f; const double epsilon = 1e-10f;
const double infty = 1e15f; // This vlalue is used to represent missing values in data const double infty = 1e15f; // This value is used to represent missing values in data
// --------------------------- Raycasting Constants --------------------------- // --------------------------- Raycasting Constants ---------------------------
@ -41,4 +41,4 @@ extern __device__ Point3 d_lightPos;
// --------------------------- Functions for handling external constants --------------------------- // --------------------------- Functions for handling external constants ---------------------------
void copyConstantsToDevice(); void copyConstantsToDevice();
#endif // CONSTS_H #endif // CONSTS_H

13
src/cuda_error.cpp Normal file
View File

@ -0,0 +1,13 @@
#include "cuda_error.h"
#include "cuda_runtime.h"
#include <iostream>
void check_cuda(cudaError_t res, char const* const func, const char* const file, int const line) {
if (res) {
std::cout << "CUDA encountered an error: " << cudaGetErrorString(res) << " in " << file << ":" << line << std::endl;
cudaDeviceReset();
exit(1);
}
}

10
src/cuda_error.h Normal file
View File

@ -0,0 +1,10 @@
#ifndef CUDA_ERROR_H
#define CUDA_ERROR_H
#include "cuda_runtime.h"
#define check_cuda_errors(val) check_cuda( (val), #val, __FILE__, __LINE__ )
void check_cuda(cudaError_t res, char const* const func, const char* const file, int const line);
#endif // CUDA_ERROR_H

125
src/gui/MainWindow.cpp Normal file
View File

@ -0,0 +1,125 @@
#include "MainWindow.h"
#include "cuda_runtime.h"
#include <iostream>
#include <memory>
#include "Shader.h"
Window::Window(unsigned int w, unsigned int h) {
this->w = w;
this->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<Window*>(glfwGetWindowUserPointer(window));
newWin->resize(w, h);
}
int Window::init(float* data) {
// 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 // TODO: switch from this style input to something more resembling an actual gui
glfwSetInputMode(this->window, GLFW_CURSOR, GLFW_CURSOR_DISABLED);
glfwSetWindowUserPointer(this->window, reinterpret_cast<void*>(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(data)) return -1;
this->last_frame = std::chrono::steady_clock::now();
while (!glfwWindowShouldClose(window)) {
Window::tick();
}
Window::free(data);
return 0;
}
int Window::init_quad(float* data) {
this->quad = std::make_unique<Quad>(this->w, this->h);
this->quad->cuda_init(data);
this->quad->make_fbo();
this->shader = std::make_unique<Shader>("./shaders/vertshader.glsl", "./shaders/fragshader.glsl");
this->shader->use();
glUniform1i(glGetUniformLocation(this->shader->ID, "currentFrameTex"), 0);
return 0;
}
void Window::free(float* data) {
// 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->quad = nullptr;
cudaFree(data);
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<std::chrono::milliseconds>(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);
this->quad->render();
this->shader->use();
glBindVertexArray(this->quad->VAO);
glBindTexture(GL_TEXTURE_2D, this->quad->tex);
glDrawArrays(GL_TRIANGLES, 0, 6); // draw current frame to texture
// check for events
glfwSwapBuffers(this->window);
glfwPollEvents();
}
void Window::resize(unsigned int w, unsigned int h) {
this->w = w;
this->h = h;
this->quad->resize(w, h);
}

34
src/gui/MainWindow.h Normal file
View File

@ -0,0 +1,34 @@
#ifndef MAINWINDOW_H
#define MAINWINDOW_H
#include "Quad.h"
#include "Shader.h"
#include <glad/glad.h>
#include <GLFW/glfw3.h>
#include <chrono>
class Window {
public:
unsigned int w;
unsigned int h;
float* data;
Window(unsigned int w, unsigned int h);
int init(float* data);
void free(float* data);
void resize(unsigned int w, unsigned int h);
private:
GLFWwindow* window;
std::unique_ptr<Quad> quad;
std::chrono::steady_clock::time_point last_frame;
void tick();
int init_quad(float* data);
std::unique_ptr<Shader> shader;
};
#endif // MAINWINDOW_H

117
src/gui/Quad.cpp Normal file
View File

@ -0,0 +1,117 @@
#include "Quad.h"
#include "cuda_error.h"
#include <glad/glad.h>
#include <GLFW/glfw3.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda_gl_interop.h>
#include <iostream>
Quad::Quad(unsigned int w, unsigned int h) {
this->w = w;
this->h = h;
std::vector<float> 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);
};
void Quad::make_fbo(){
glGenFramebuffers(1, &fb);
glBindFramebuffer(GL_FRAMEBUFFER, fb);
glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, tex, 0);
glBindFramebuffer(GL_FRAMEBUFFER, 0);
}
Quad::~Quad() {
check_cuda_errors(cudaGraphicsUnregisterResource(CGR));
};
void Quad::cuda_init(float* data) {
check_cuda_errors(cudaGraphicsGLRegisterBuffer(&this->CGR, this->PBO, cudaGraphicsRegisterFlagsNone));
this->renderer = std::make_unique<Raycaster>(this->CGR, this->w, this->h, data);
};
void Quad::render() {
glBindTexture(GL_TEXTURE_2D, 0);
this->renderer->render();
glBindTexture(GL_TEXTURE_2D, this->tex);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, this->PBO);
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) {
check_cuda_errors(cudaGraphicsUnregisterResource(CGR));
check_cuda_errors(cudaGraphicsGLRegisterBuffer(&this->CGR, this->PBO, cudaGraphicsRegisterFlagsNone));
this->renderer->resources = this->CGR;
this->renderer->resize(w, h);
}
};

37
src/gui/Quad.h Normal file
View File

@ -0,0 +1,37 @@
#ifndef QUAD_H
#define QUAD_H
#include <glad/glad.h>
#include <GLFW/glfw3.h>
#include <cuda_runtime.h>
#include "illumination/Raycaster.h"
#include <vector>
#include <memory>
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<Raycaster> renderer;
Quad(unsigned int w, unsigned int h);
~Quad();
void render();
void resize(unsigned int w, unsigned int h);
void cuda_init(float* data);
void make_fbo();
};
#endif // QUAD_H

126
src/gui/Shader.h Normal file
View File

@ -0,0 +1,126 @@
// Shader class nicked from learnOpengl here: https://learnopengl.com/Getting-started/Shaders
// which is published under the CC BY-NC 4.0 license (and thus fine to use here)
#ifndef SHADER_H
#define SHADER_H
#include <glad/glad.h>
#include <string>
#include <fstream>
#include <sstream>
#include <iostream>
class Shader
{
public:
unsigned int ID;
// constructor generates the shader on the fly
// ------------------------------------------------------------------------
Shader(const char* vertexPath, const char* fragmentPath)
{
// 1. retrieve the vertex/fragment source code from filePath
std::string vertexCode;
std::string fragmentCode;
std::ifstream vShaderFile;
std::ifstream fShaderFile;
// ensure ifstream objects can throw exceptions:
vShaderFile.exceptions (std::ifstream::failbit | std::ifstream::badbit);
fShaderFile.exceptions (std::ifstream::failbit | std::ifstream::badbit);
try
{
// open files
vShaderFile.open(vertexPath);
fShaderFile.open(fragmentPath);
std::stringstream vShaderStream, fShaderStream;
// read file's buffer contents into streams
vShaderStream << vShaderFile.rdbuf();
fShaderStream << fShaderFile.rdbuf();
// close file handlers
vShaderFile.close();
fShaderFile.close();
// convert stream into string
vertexCode = vShaderStream.str();
fragmentCode = fShaderStream.str();
}
catch (std::ifstream::failure& e)
{
std::cout << "ERROR::SHADER::FILE_NOT_SUCCESSFULLY_READ: " << e.what() << std::endl;
}
const char* vShaderCode = vertexCode.c_str();
const char * fShaderCode = fragmentCode.c_str();
// 2. compile shaders
unsigned int vertex, fragment;
// vertex shader
vertex = glCreateShader(GL_VERTEX_SHADER);
glShaderSource(vertex, 1, &vShaderCode, NULL);
glCompileShader(vertex);
checkCompileErrors(vertex, "VERTEX");
// fragment Shader
fragment = glCreateShader(GL_FRAGMENT_SHADER);
glShaderSource(fragment, 1, &fShaderCode, NULL);
glCompileShader(fragment);
checkCompileErrors(fragment, "FRAGMENT");
// shader Program
ID = glCreateProgram();
glAttachShader(ID, vertex);
glAttachShader(ID, fragment);
glLinkProgram(ID);
checkCompileErrors(ID, "PROGRAM");
// delete the shaders as they're linked into our program now and no longer necessary
glDeleteShader(vertex);
glDeleteShader(fragment);
}
// activate the shader
// ------------------------------------------------------------------------
void use()
{
glUseProgram(ID);
}
// utility uniform functions
// ------------------------------------------------------------------------
void setBool(const std::string &name, bool value) const
{
glUniform1i(glGetUniformLocation(ID, name.c_str()), (int)value);
}
// ------------------------------------------------------------------------
void setInt(const std::string &name, int value) const
{
glUniform1i(glGetUniformLocation(ID, name.c_str()), value);
}
// ------------------------------------------------------------------------
void setFloat(const std::string &name, float value) const
{
glUniform1f(glGetUniformLocation(ID, name.c_str()), value);
}
private:
// utility function for checking shader compilation/linking errors.
// ------------------------------------------------------------------------
void checkCompileErrors(unsigned int shader, std::string type)
{
int success;
char infoLog[1024];
if (type != "PROGRAM")
{
glGetShaderiv(shader, GL_COMPILE_STATUS, &success);
if (!success)
{
glGetShaderInfoLog(shader, 1024, NULL, infoLog);
std::cout << "ERROR::SHADER_COMPILATION_ERROR of type: " << type << "\n" << infoLog << "\n -- --------------------------------------------------- -- " << std::endl;
}
}
else
{
glGetProgramiv(shader, GL_LINK_STATUS, &success);
if (!success)
{
glGetProgramInfoLog(shader, 1024, NULL, infoLog);
std::cout << "ERROR::PROGRAM_LINKING_ERROR of type: " << type << "\n" << infoLog << "\n -- --------------------------------------------------- -- " << std::endl;
}
}
}
};
#endif

12
src/gui/shaders/fragshader.glsl Normal file
View File

@ -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);
}

11
src/gui/shaders/vertshader.glsl Normal file
View File

@ -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);
}

12
src/illumination/FrameBuffer.cu Normal file
View File

@ -0,0 +1,12 @@
#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);
}

20
src/illumination/FrameBuffer.h Normal file
View File

@ -0,0 +1,20 @@
#ifndef FRAMEBUFFER_H
#define FRAMEBUFFER_H
#include "cuda_runtime.h"
#include "linalg/linalg.h"
#include <cstdint>
class FrameBuffer {
public:
unsigned int* 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

78
src/illumination/raycaster.h → src/illumination/Raycaster.cu
View File

@ -1,14 +1,18 @@
#ifndef RAYCASTER_H #include "Raycaster.h"
#define RAYCASTER_H
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda_runtime.h>
#include "linalg/linalg.h" #include "linalg/linalg.h"
#include "consts.h" #include "consts.h"
#include "cuda_error.h"
#include "shading.h" #include "shading.h"
#include <iostream>
#include "objs/sphere.h"
// Raycast + phong, TODO: Consider wrapping in a class // TODO: instead of IMAGEWIDTH and IMAGEHEIGHT this should reflect the windowSize;
__global__ void raycastKernel(float* volumeData, unsigned char* framebuffer) { __global__ void raycastKernel(float* volumeData, FrameBuffer framebuffer) {
int px = blockIdx.x * blockDim.x + threadIdx.x; int px = blockIdx.x * blockDim.x + threadIdx.x;
int py = blockIdx.y * blockDim.y + threadIdx.y; int py = blockIdx.y * blockDim.y + threadIdx.y;
if (px >= IMAGE_WIDTH || py >= IMAGE_HEIGHT) return; if (px >= IMAGE_WIDTH || py >= IMAGE_HEIGHT) return;
@ -39,7 +43,7 @@ __global__ void raycastKernel(float* volumeData, unsigned char* framebuffer) {
auto intersectAxis = [&](float start, float dirVal) { auto intersectAxis = [&](float start, float dirVal) {
if (fabsf(dirVal) < epsilon) { if (fabsf(dirVal) < epsilon) {
if (start < 0.f || start > 1.f) { if (start < 0.f || start > 1.f) {
tNear = 1e9f; tNear = 1e9f;
tFar = -1e9f; tFar = -1e9f;
} }
} else { } else {
@ -117,10 +121,62 @@ __global__ void raycastKernel(float* volumeData, unsigned char* framebuffer) {
accumB /= (float)SAMPLES_PER_PIXEL; accumB /= (float)SAMPLES_PER_PIXEL;
// Final colour // Final colour
int fbIndex = (py * IMAGE_WIDTH + px) * 3; framebuffer.writePixel(px, py, accumR, accumG, accumB);
framebuffer[fbIndex + 0] = (unsigned char)(fminf(accumR, 1.f) * 255); // int fbIndex = (py * IMAGE_WIDTH + px) * 3;
framebuffer[fbIndex + 1] = (unsigned char)(fminf(accumG, 1.f) * 255); // framebuffer[fbIndex + 0] = (unsigned char)(fminf(accumR, 1.f) * 255);
framebuffer[fbIndex + 2] = (unsigned char)(fminf(accumB, 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
Raycaster::Raycaster(cudaGraphicsResource_t resources, int w, int h, float* data) {
this->resources = resources;
this->w = w;
this->h = h;
this->fb = new FrameBuffer(w, h);
this->data = data;
// 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<Camera*>();
check_cuda_errors(cudaDeviceSynchronize());
}
void Raycaster::render() {
check_cuda_errors(cudaGraphicsMapResources(1, &this->resources));
check_cuda_errors(cudaGraphicsResourceGetMappedPointer((void**)&(this->fb->buffer), &(this->fb->buffer_size), resources));
// FIXME: might not be the best parallelization configuraiton
int tx = 16;
int ty = 16;
dim3 threadSize(this->w / tx + 1, this->h / ty + 1);
dim3 blockSize(tx, ty);
// TODO: pass camera info at some point
// frame buffer is implicitly copied to the device each frame
raycastKernel<<<threadSize, blockSize>>> (this->data, *this->fb);
check_cuda_errors(cudaGetLastError());
check_cuda_errors(cudaDeviceSynchronize());
check_cuda_errors(cudaGraphicsUnmapResources(1, &this->resources));
}
void Raycaster::resize(int w, int h) {
this->w = w;
this->h = h;
delete this->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);
check_cuda_errors(cudaDeviceSynchronize());
}

33
src/illumination/Raycaster.h Normal file
View File

@ -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 <thrust/device_ptr.h>
__global__ void raycastKernel(float* volumeData, unsigned char* framebuffer);
struct Raycaster {
// thrust::device_ptr<Camera*> d_camera;
// CameraInfo camera_info;
cudaGraphicsResource_t resources;
FrameBuffer* fb;
float* data;
int w;
int h;
Raycaster(cudaGraphicsResource_t resources, int nx, int ny, float* data);
// ~Raycaster();
void set_camera(Vec3 position, Vec3 forward, Vec3 up);
void render();
void resize(int nx, int ny);
};
#endif // RAYCASTER_H

4
src/illumination/illumination.h
View File

@ -1,7 +1,7 @@
#ifndef ILLUMINATION_H #ifndef ILLUMINATION_H
#define ILLUMINATION_H #define ILLUMINATION_H
#include "raycaster.h" #include "Raycaster.h"
#include "shading.h" #include "shading.h"
#endif // ILLUMINATION_H #endif // ILLUMINATION_H

14
src/illumination/shading.cu Normal file
View File

@ -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;
};

14
src/illumination/shading.h
View File

@ -4,17 +4,7 @@
#include "linalg/linalg.h" #include "linalg/linalg.h"
#include "consts.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);
__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; #endif // SHADING_H
}
#endif // SHADING_H

46
src/linalg/mat.cu Normal file
View File

@ -0,0 +1,46 @@
#include "mat.h"
#include <vector>
#include <algorithm>
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;
float len = sqrtf(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?
u.in[0] = round(r/len * 255.0f);
u.in[1] = round(g/len * 255.0f);
u.in[2] = round(b/len * 255.0f);
u.in[3] = round(a/len * 255.0f);
return u.out;
}

26
src/linalg/mat.h
View File

@ -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) { #include "vec.h"
// 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); __device__ Vec3 computeGradient(float* volumeData, const int volW, const int volH, const int volD, int x, int y, int z);
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 __device__ unsigned int packUnorm4x8(float r, float g, float b, float a);
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); #endif // MAT_H
}

2
src/linalg/vec.h
View File

@ -27,4 +27,4 @@ struct Vec3 { // TODO: Maybe make this into a class ... maybe
}; };
typedef Vec3 Point3; typedef Vec3 Point3;
typedef Vec3 Color3; typedef Vec3 Color3;

139
src/main.cu
View File

@ -1,20 +1,23 @@
#include <iostream>
#include <fstream>
#include <cmath>
#include <cuda_runtime.h>
#include <vector>
#include <algorithm> #include <algorithm>
#include <cmath>
#include "hurricanedata/datareader.h"
#include "linalg/linalg.h"
#include "objs/sphere.h"
#include "img/handler.h"
#include "consts.h" #include "consts.h"
#include <cuda_runtime.h>
#include <fstream>
#include "gui/MainWindow.h"
#include "hurricanedata/datareader.h"
#include "illumination/illumination.h" #include "illumination/illumination.h"
#include "img/handler.h"
#include <iostream>
#include "linalg/linalg.h"
#include <vector>
static float* d_volume = nullptr; static float* d_volume = nullptr;
// TODO: general
// * pass camera_info to the raycasting function - updated according to glfw.
// * on that note, code for handling input (mouse movement certainly, possibly free input / 4 pre-coded views, q/esc to quit, space for pause (would be were the 'simple' render idea would come in))
// * very similarly - actual code for loading new data as the simulation progresses - right now its effectively a static image loader
void getTemperature(std::vector<float>& temperatureData, int idx = 0) { void getTemperature(std::vector<float>& temperatureData, int idx = 0) {
std::string path = "data/trimmed"; std::string path = "data/trimmed";
@ -46,67 +49,75 @@ void getSpeed(std::vector<float>& speedData, int idx = 0) {
} }
} }
int main(int argc, char** argv) {
std::vector<float> data; int main() {
// getTemperature(data); std::vector<float> data;
getSpeed(data); // getTemperature(data);
getSpeed(data);
// TODO: Eveontually remove debug below (i.e., eliminate for-loop etc.) // TODO: Eveontually remove debug below (i.e., eliminate for-loop etc.)
// Generate debug volume data // Generate debug volume data
float* hostVolume = new float[VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH]; float* hostVolume = new float[VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH];
// generateVolume(hostVolume, VOLUME_WIDTH, VOLUME_HEIGHT, VOLUME_DEPTH); // generateVolume(hostVolume, VOLUME_WIDTH, VOLUME_HEIGHT, VOLUME_DEPTH);
for (int i = 0; i < VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH; i++) { // TODO: This is technically an unnecessary artifact of the old code taking in a float* instead of a std::vector for (int i = 0; i < VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH; i++) { // TODO: This is technically an unnecessary artifact of the old code taking in a float* instead of a std::vector
// Discard temperatures above a small star (supposedly, missing temperature values) // Discard temperatures above a small star (supposedly, missing temperature values)
hostVolume[i] = data[i]; hostVolume[i] = data[i];
if (data[i] + epsilon >= infty) hostVolume[i] = 0.0f; if (data[i] + epsilon >= infty) hostVolume[i] = 0.0f;
} }
// Min-max normalization // Min-max normalization
float minVal = *std::min_element(hostVolume, hostVolume + VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH); float minVal = *std::min_element(hostVolume, hostVolume + VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH);
float maxVal = *std::max_element(hostVolume, hostVolume + VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH); float maxVal = *std::max_element(hostVolume, hostVolume + VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH);
for (int i = 0; i < VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH; i++) { for (int i = 0; i < VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH; i++) {
hostVolume[i] = (hostVolume[i] - minVal) / (maxVal - minVal); hostVolume[i] = (hostVolume[i] - minVal) / (maxVal - minVal);
} }
// Allocate + copy data to GPU // Allocate + copy data to GPU
size_t volumeSize = sizeof(float) * VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH; size_t volumeSize = sizeof(float) * VOLUME_WIDTH * VOLUME_HEIGHT * VOLUME_DEPTH;
cudaMalloc((void**)&d_volume, volumeSize); cudaMalloc((void**)&d_volume, volumeSize);
cudaMemcpy(d_volume, hostVolume, volumeSize, cudaMemcpyHostToDevice); cudaMemcpy(d_volume, hostVolume, volumeSize, cudaMemcpyHostToDevice);
// Allocate framebuffer // Allocate framebuffer
unsigned char* d_framebuffer; // unsigned char* d_framebuffer;
size_t fbSize = IMAGE_WIDTH * IMAGE_HEIGHT * 3 * sizeof(unsigned char); // size_t fbSize = IMAGE_WIDTH * IMAGE_HEIGHT * 3 * sizeof(unsigned char);
cudaMalloc((void**)&d_framebuffer, fbSize); // cudaMalloc((void**)&d_framebuffer, fbSize);
cudaMemset(d_framebuffer, 0, fbSize); // cudaMemset(d_framebuffer, 0, fbSize);
// Copy external constants from consts.h to cuda // Copy external constants from consts.h to cuda
copyConstantsToDevice(); copyConstantsToDevice();
// Launch kernel // NOTE: this is done within the rayTracer class
dim3 blockSize(16, 16); // TODO: Figure out a good size for parallelization // // Launch kernel
dim3 gridSize((IMAGE_WIDTH + blockSize.x - 1)/blockSize.x, // dim3 blockSize(16, 16);
(IMAGE_HEIGHT + blockSize.y - 1)/blockSize.y); // dim3 gridSize((IMAGE_WIDTH + blockSize.x - 1)/blockSize.x,
// (IMAGE_HEIGHT + blockSize.y - 1)/blockSize.y);
//
// raycastKernel<<<gridSize, blockSize>>>(
// d_volume,
// d_framebuffer
// );
// cudaDeviceSynchronize();
raycastKernel<<<gridSize, blockSize>>>( Window window(IMAGE_WIDTH, IMAGE_HEIGHT);
d_volume, int out = window.init(d_volume);
d_framebuffer
);
cudaDeviceSynchronize();
// Copy framebuffer back to CPU cudaFree(d_volume);
unsigned char* hostFramebuffer = new unsigned char[IMAGE_WIDTH * IMAGE_HEIGHT * 3];
cudaMemcpy(hostFramebuffer, d_framebuffer, fbSize, cudaMemcpyDeviceToHost);
// Export image // // Copy framebuffer back to CPU
saveImage("output.ppm", hostFramebuffer, IMAGE_WIDTH, IMAGE_HEIGHT); // unsigned char* hostFramebuffer = new unsigned char[IMAGE_WIDTH * IMAGE_HEIGHT * 3];
// cudaMemcpy(hostFramebuffer, d_framebuffer, fbSize, cudaMemcpyDeviceToHost);
// Cleanup //
delete[] hostVolume; // // Export image
delete[] hostFramebuffer; // saveImage("output.ppm", hostFramebuffer, IMAGE_WIDTH, IMAGE_HEIGHT);
cudaFree(d_volume); //
cudaFree(d_framebuffer); // // Cleanup //TODO: cleanup properly
delete[] hostVolume;
std::cout << "Phong-DVR rendering done. Image saved to output.ppm" << std::endl; // delete[] hostFramebuffer;
return 0; // cudaFree(d_volume);
} // cudaFree(d_framebuffer);
//
// std::cout << "Phong-DVR rendering done. Image saved to output.ppm" << std::endl;
// return 0;
return out;
}

1833
thirdparty/glad.c vendored Normal file
View File

File diff suppressed because it is too large Load Diff