Skip to content
Snippets Groups Projects
Commit 3e46bc0c authored by Thomas Müller's avatar Thomas Müller
Browse files

Merge branch 'master' into patch-1 

# Conflicts:
#	docs/nerf_dataset_tips.md
parents 624605f6 e53d2ca0
No related branches found
No related tags found
No related merge requests found
Expand all Hide whitespace changes
Inline Side-by-side
Showing
with 627 additions and 340 deletions
.editorconfig
+ 1
− 1
View file @ 3e46bc0c
...@@ -11,5 +11,5 @@ trim_trailing_whitespace = true ...@@ -11,5 +11,5 @@ trim_trailing_whitespace = true
trim_trailing_whitespace = false trim_trailing_whitespace = false
[*.yml] [*.yml]
indent_style = spaces indent_style = space
indent_size = 2 indent_size = 2
.github/workflows/main.yml
+ 2
− 1
View file @ 3e46bc0c
...@@ -147,11 +147,12 @@ jobs: ...@@ -147,11 +147,12 @@ jobs:
docs/assets_readme/ docs/assets_readme/
data/ data/
scripts/flip/* scripts/flip/*
scripts/category2id.json
scripts/colmap2nerf.py
scripts/common.py scripts/common.py
scripts/convert_image.py scripts/convert_image.py
scripts/download_colmap.bat scripts/download_colmap.bat
scripts/download_ffmpeg.bat scripts/download_ffmpeg.bat
scripts/colmap2nerf.py
scripts/nsvf2nerf.py scripts/nsvf2nerf.py
scripts/record3d2nerf.py scripts/record3d2nerf.py
......
.gitignore
+ 1
− 0
View file @ 3e46bc0c
...@@ -12,6 +12,7 @@ ...@@ -12,6 +12,7 @@
/*.dll /*.dll
/*.exe /*.exe
/*.json /*.json
*.ingp
*.msgpack *.msgpack
*.training *.training
__pycache__ __pycache__
......
.gitmodules
+ 9
− 0
View file @ 3e46bc0c
...@@ -22,3 +22,12 @@ ...@@ -22,3 +22,12 @@
[submodule "dependencies/dlss"] [submodule "dependencies/dlss"]
path = dependencies/dlss path = dependencies/dlss
url = https://github.com/NVIDIA/DLSS url = https://github.com/NVIDIA/DLSS
[submodule "dependencies/zstr"]
path = dependencies/zstr
url = https://github.com/Tom94/zstr
[submodule "dependencies/zlib"]
path = dependencies/zlib
url = https://github.com/Tom94/zlib
[submodule "dependencies/OpenXR-SDK"]
path = dependencies/OpenXR-SDK
url = https://github.com/KhronosGroup/OpenXR-SDK.git
CMakeLists.txt
+ 66
− 24
View file @ 3e46bc0c
...@@ -71,7 +71,6 @@ set(CMAKE_CUDA_RUNTIME_LIBRARY Shared) ...@@ -71,7 +71,6 @@ set(CMAKE_CUDA_RUNTIME_LIBRARY Shared)
if (MSVC) if (MSVC)
list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=/bigobj") list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=/bigobj")
else() else()
list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=-mf16c")
list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=-Wno-float-conversion") list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=-Wno-float-conversion")
list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=-fno-strict-aliasing") list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=-fno-strict-aliasing")
list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=-fPIC") list(APPEND CUDA_NVCC_FLAGS "-Xcompiler=-fPIC")
...@@ -104,11 +103,7 @@ if (NGP_BUILD_WITH_GUI) ...@@ -104,11 +103,7 @@ if (NGP_BUILD_WITH_GUI)
list(APPEND NGP_INCLUDE_DIRECTORIES "dependencies/dlss/include") list(APPEND NGP_INCLUDE_DIRECTORIES "dependencies/dlss/include")
if (MSVC) if (MSVC)
list(APPEND NGP_LINK_DIRECTORIES "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Windows_x86_64/x86_64") list(APPEND NGP_LINK_DIRECTORIES "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Windows_x86_64/x86_64")
if (CMAKE_BUILD_TYPE STREQUAL "Debug") list(APPEND NGP_LIBRARIES "$<IF:$<CONFIG:Debug>,nvsdk_ngx_d_dbg,nvsdk_ngx_d>")
list(APPEND NGP_LIBRARIES nvsdk_ngx_d_dbg)
else()
list(APPEND NGP_LIBRARIES nvsdk_ngx_d)
endif()
else() else()
list(APPEND NGP_LINK_DIRECTORIES "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Linux_x86_64") list(APPEND NGP_LINK_DIRECTORIES "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Linux_x86_64")
list(APPEND NGP_LIBRARIES nvsdk_ngx) list(APPEND NGP_LIBRARIES nvsdk_ngx)
...@@ -123,6 +118,38 @@ if (NGP_BUILD_WITH_GUI) ...@@ -123,6 +118,38 @@ if (NGP_BUILD_WITH_GUI)
endif() endif()
endif() endif()
# OpenXR
if (WIN32)
list(APPEND NGP_DEFINITIONS -DXR_USE_PLATFORM_WIN32 -DGLFW_EXPOSE_NATIVE_WGL)
elseif (UNIX AND NOT APPLE)
list(APPEND NGP_DEFINITIONS -DGLFW_EXPOSE_NATIVE_GLX)
if (JK_USE_WAYLAND)
set(PRESENTATION_BACKEND wayland CACHE STRING " " FORCE)
set(BUILD_WITH_XLIB_HEADERS OFF CACHE BOOL " " FORCE)
set(BUILD_WITH_XCB_HEADERS OFF CACHE BOOL " " FORCE)
set(BUILD_WITH_WAYLAND_HEADERS ON CACHE BOOL " " FORCE)
list(APPEND NGP_DEFINITIONS -DGLFW_EXPOSE_NATIVE_WAYLAND -DXR_USE_PLATFORM_WAYLAND)
else()
set(PRESENTATION_BACKEND xlib CACHE STRING " " FORCE)
set(BUILD_WITH_XLIB_HEADERS ON CACHE BOOL " " FORCE)
set(BUILD_WITH_XCB_HEADERS OFF CACHE BOOL " " FORCE)
set(BUILD_WITH_WAYLAND_HEADERS OFF CACHE BOOL " " FORCE)
list(APPEND NGP_DEFINITIONS -DGLFW_EXPOSE_NATIVE_X11 -DXR_USE_PLATFORM_XLIB)
endif()
else()
message(FATAL_ERROR "No OpenXR platform set for this OS")
endif()
add_subdirectory(dependencies/OpenXR-SDK)
list(APPEND NGP_INCLUDE_DIRECTORIES "dependencies/OpenXR-SDK/include" "dependencies/OpenXR-SDK/src/common")
list(APPEND NGP_LIBRARIES openxr_loader)
list(APPEND GUI_SOURCES src/openxr_hmd.cu)
# OpenGL
find_package(OpenGL REQUIRED)
# GLFW
set(GLFW_BUILD_EXAMPLES OFF CACHE BOOL " " FORCE) set(GLFW_BUILD_EXAMPLES OFF CACHE BOOL " " FORCE)
set(GLFW_BUILD_TESTS OFF CACHE BOOL " " FORCE) set(GLFW_BUILD_TESTS OFF CACHE BOOL " " FORCE)
set(GLFW_BUILD_DOCS OFF CACHE BOOL " " FORCE) set(GLFW_BUILD_DOCS OFF CACHE BOOL " " FORCE)
...@@ -193,6 +220,26 @@ if (Python_FOUND) ...@@ -193,6 +220,26 @@ if (Python_FOUND)
add_subdirectory("dependencies/pybind11") add_subdirectory("dependencies/pybind11")
endif() endif()
# Compile zlib (only on Windows)
if (WIN32)
set(ZLIB_USE_STATIC_LIBS ON CACHE BOOL " " FORCE)
set(ZLIB_BUILD_STATIC_LIBS ON CACHE BOOL " " FORCE)
set(ZLIB_BUILD_SHARED_LIBS OFF CACHE BOOL " " FORCE)
set(SKIP_INSTALL_ALL ON CACHE BOOL " " FORCE)
add_subdirectory("dependencies/zlib")
set_property(TARGET zlibstatic PROPERTY FOLDER "dependencies")
set(ZLIB_INCLUDE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/zlib" CACHE PATH " " FORCE)
set(ZLIB_LIBRARY zlibstatic)
include_directories(${ZLIB_INCLUDE_DIR} "${CMAKE_CURRENT_BINARY_DIR}/dependencies/zlib")
list(APPEND NGP_LIBRARIES zlibstatic)
endif()
add_subdirectory("dependencies/zstr")
list(APPEND NGP_LIBRARIES zstr::zstr)
############################################################################### ###############################################################################
# Program # Program
...@@ -286,24 +333,17 @@ target_link_libraries(instant-ngp PRIVATE ngp) ...@@ -286,24 +333,17 @@ target_link_libraries(instant-ngp PRIVATE ngp)
# Copy DLSS shared libraries # Copy DLSS shared libraries
if (NGP_VULKAN) if (NGP_VULKAN)
if (CMAKE_BUILD_TYPE STREQUAL "Debug") set(NGX_BUILD_DIR "$<IF:$<CONFIG:Debug>,dev,rel>")
set(NGX_BUILD_DIR "dev")
else()
set(NGX_BUILD_DIR "rel")
endif()
if (MSVC) if (MSVC)
add_custom_command(TARGET instant-ngp POST_BUILD set(NGX_SHARED_LIB "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Windows_x86_64/${NGX_BUILD_DIR}/nvngx_dlss.dll")
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Windows_x86_64/${NGX_BUILD_DIR}/nvngx_dlss.dll" $<TARGET_FILE_DIR:instant-ngp>
COMMAND_EXPAND_LISTS
)
else() else()
file(GLOB DLSS_SOS "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Linux_x86_64/${NGX_BUILD_DIR}/libnvidia-ngx-dlss.so.*") set(NGX_SHARED_LIB "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Linux_x86_64/${NGX_BUILD_DIR}/libnvidia-ngx-dlss.so.*")
add_custom_command(TARGET instant-ngp POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy ${DLSS_SOS} $<TARGET_FILE_DIR:instant-ngp>
COMMAND_EXPAND_LISTS
)
endif() endif()
add_custom_command(TARGET instant-ngp POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy "${NGX_SHARED_LIB}" $<TARGET_FILE_DIR:instant-ngp>
COMMAND_EXPAND_LISTS
)
endif() endif()
if (MSVC) if (MSVC)
...@@ -329,9 +369,11 @@ endif() ...@@ -329,9 +369,11 @@ endif()
# Link the executable to the project directory and copy over DLLs such that instant-ngp can be invoked without going into the build folder. # Link the executable to the project directory and copy over DLLs such that instant-ngp can be invoked without going into the build folder.
set(NGP_BINARY_FILE "\"${CMAKE_CURRENT_SOURCE_DIR}/$<TARGET_FILE_NAME:instant-ngp>\"") set(NGP_BINARY_FILE "\"${CMAKE_CURRENT_SOURCE_DIR}/$<TARGET_FILE_NAME:instant-ngp>\"")
if (MSVC) if (MSVC)
file(TO_NATIVE_PATH ${NGP_BINARY_FILE} NGP_BINARY_FILE) add_custom_command(TARGET instant-ngp POST_BUILD COMMAND ${CMAKE_COMMAND} -E copy $<TARGET_FILE:instant-ngp> ${CMAKE_CURRENT_SOURCE_DIR})
add_custom_command(TARGET instant-ngp POST_BUILD COMMAND cmd /C del /F /Q "${NGP_BINARY_FILE}" && cmd /C mklink /H "${NGP_BINARY_FILE}" "\"$<TARGET_FILE:instant-ngp>\"") file(GLOB NGP_REQUIRED_DLLS "${CUDA_COMPILER_BIN}/cudart64*.dll")
file(GLOB NGP_REQUIRED_DLLS "${CUDA_COMPILER_BIN}/cudart64*.dll" "${CMAKE_CURRENT_SOURCE_DIR}/dependencies/dlss/lib/Windows_x86_64/${NGX_BUILD_DIR}/nvngx_dlss.dll") if (NGP_VULKAN)
list(APPEND NGP_REQUIRED_DLLS "${NGX_SHARED_LIB}")
endif()
if (NGP_REQUIRED_DLLS) if (NGP_REQUIRED_DLLS)
add_custom_command(TARGET instant-ngp POST_BUILD add_custom_command(TARGET instant-ngp POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy ${NGP_REQUIRED_DLLS} ${CMAKE_CURRENT_SOURCE_DIR} COMMAND ${CMAKE_COMMAND} -E copy ${NGP_REQUIRED_DLLS} ${CMAKE_CURRENT_SOURCE_DIR}
......
LICENSE.txt
+ 1
− 1
View file @ 3e46bc0c
Copyright (c) 2022, NVIDIA Corporation & affiliates. All rights reserved. Copyright (c) 2022-2023, NVIDIA Corporation & affiliates. All rights reserved.
NVIDIA Source Code License for instant neural graphics primitives NVIDIA Source Code License for instant neural graphics primitives
......
README.md
+ 143
− 131
View file @ 3e46bc0c
This diff is collapsed.
OpenXR-SDK @ e2da9ce8
Subproject commit e2da9ce83a4388c9622da328bf48548471261290
dependencies/filesystem/filesystem/directory.h
+ 4
− 4
View file @ 3e46bc0c
...@@ -39,8 +39,8 @@ public: ...@@ -39,8 +39,8 @@ public:
iterator(const directory& dir) { iterator(const directory& dir) {
m_dir = dir; m_dir = dir;
#if defined(_WIN32) #if defined(_WIN32)
std::string search_path(dir.make_absolute().str() + "/*.*"); std::wstring search_path(dir.make_absolute().wstr() + L"/*.*");
m_handle = _findfirst(search_path.c_str(), &m_data); m_handle = _wfindfirst(search_path.c_str(), &m_data);
if (is_valid_handler()) if (is_valid_handler())
{ {
m_result = m_dir / m_data.name; m_result = m_dir / m_data.name;
...@@ -72,7 +72,7 @@ public: ...@@ -72,7 +72,7 @@ public:
if (is_valid_handler()) if (is_valid_handler())
{ {
#if defined(_WIN32) #if defined(_WIN32)
if (_findnext(m_handle, &m_data)) if (_wfindnext(m_handle, &m_data))
{ {
if (ENOENT == errno) /* reaching the end */ if (ENOENT == errno) /* reaching the end */
{ {
...@@ -133,7 +133,7 @@ public: ...@@ -133,7 +133,7 @@ public:
path m_result; path m_result;
#if defined(_WIN32) #if defined(_WIN32)
intptr_t m_handle; intptr_t m_handle;
_finddata_t m_data; _wfinddata_t m_data;
#else #else
DIR* m_handle; DIR* m_handle;
struct dirent* m_data; struct dirent* m_data;
......
dependencies/filesystem/filesystem/path.h
+ 2
− 0
View file @ 3e46bc0c
...@@ -296,6 +296,8 @@ public: ...@@ -296,6 +296,8 @@ public:
m_path = tokenize(str, "/"); m_path = tokenize(str, "/");
m_absolute = !str.empty() && str[0] == '/'; m_absolute = !str.empty() && str[0] == '/';
} }
m_path.erase(std::remove(std::begin(m_path), std::end(m_path), ""), std::end(m_path));
} }
path &operator=(const path &path) { path &operator=(const path &path) {
......
tiny-cuda-nn @ b242f6e4
Compare 14053e9a...b242f6e4
Subproject commit 14053e9a87ebf449d32bda335c0363dd4f5667a4 Subproject commit b242f6e4ed6a09f39adc7fa4f98ac305ed987fde
zlib @ fafd714d
Subproject commit fafd714d8b3ce85a6decf566e3b5b611352eadea
zstr @ 007f9760
Subproject commit 007f97607865934fb52e08248f152712e428688a
docs/nerf_dataset_tips.md
+ 12
− 7
View file @ 3e46bc0c
...@@ -43,7 +43,7 @@ If you have an existing dataset in `transforms.json` format, it should be center ...@@ -43,7 +43,7 @@ If you have an existing dataset in `transforms.json` format, it should be center
You can set any of the following parameters, where the listed values are the default. You can set any of the following parameters, where the listed values are the default.
```json ```json
{ {
"aabb_scale": 16, "aabb_scale": 64,
"scale": 0.33, "scale": 0.33,
"offset": [0.5, 0.5, 0.5], "offset": [0.5, 0.5, 0.5],
... ...
...@@ -60,6 +60,8 @@ To train on self-captured data, one has to process the data into an existing for ...@@ -60,6 +60,8 @@ To train on self-captured data, one has to process the data into an existing for
All require [Python](https://www.python.org/) 3.7 or higher to be installed and available in your PATH. All require [Python](https://www.python.org/) 3.7 or higher to be installed and available in your PATH.
On Windows you can [download an installer from here](https://www.python.org/downloads/). During installation, make sure to check "add python.exe to PATH".
If you are using Debian based Linux distribution, install Python with If you are using Debian based Linux distribution, install Python with
```sh ```sh
sudo apt-get install python3-dev python3-pip sudo apt-get install python3-dev python3-pip
...@@ -70,9 +72,7 @@ Alternatively, if you are using Arch or Arch derivatives, install Python with ...@@ -70,9 +72,7 @@ Alternatively, if you are using Arch or Arch derivatives, install Python with
sudo pacman -S python python-pip sudo pacman -S python python-pip
``` ```
On Windows you can also install Python from the Microsoft Store, which will add Python to your PATH automatically. For all operating systems, after having installed Python, you need to install the required Python packages by opening a Windows Command Prompt / Linux terminal and calling
Then you need to install the required Python packages for running this software by to do so opening a Linux terminal / Windows Command Prompt and calling
```sh ```sh
pip install -r requirements.txt pip install -r requirements.txt
``` ```
...@@ -88,7 +88,7 @@ If you use Windows, you do not need to install anything. COLMAP and FFmpeg will ...@@ -88,7 +88,7 @@ If you use Windows, you do not need to install anything. COLMAP and FFmpeg will
If you are training from a video file, run the [scripts/colmap2nerf.py](/scripts/colmap2nerf.py) script from the folder containing the video, with the following recommended parameters: If you are training from a video file, run the [scripts/colmap2nerf.py](/scripts/colmap2nerf.py) script from the folder containing the video, with the following recommended parameters:
```sh ```sh
data-folder$ python [path-to-instant-ngp]/scripts/colmap2nerf.py --video_in <filename of video> --video_fps 2 --run_colmap --aabb_scale 16 data-folder$ python [path-to-instant-ngp]/scripts/colmap2nerf.py --video_in <filename of video> --video_fps 2 --run_colmap --aabb_scale 64
``` ```
The above assumes a single video file as input, which then has frames extracted at the specified framerate (2). It is recommended to choose a frame rate that leads to around 50-150 images. So for a one minute video, `--video_fps 2` is ideal. The above assumes a single video file as input, which then has frames extracted at the specified framerate (2). It is recommended to choose a frame rate that leads to around 50-150 images. So for a one minute video, `--video_fps 2` is ideal.
...@@ -96,7 +96,7 @@ The above assumes a single video file as input, which then has frames extracted ...@@ -96,7 +96,7 @@ The above assumes a single video file as input, which then has frames extracted
For training from images, place them in a subfolder called `images` and then use suitable options such as the ones below: For training from images, place them in a subfolder called `images` and then use suitable options such as the ones below:
```sh ```sh
data-folder$ python [path-to-instant-ngp]/scripts/colmap2nerf.py --colmap_matcher exhaustive --run_colmap --aabb_scale 16 data-folder$ python [path-to-instant-ngp]/scripts/colmap2nerf.py --colmap_matcher exhaustive --run_colmap --aabb_scale 64
``` ```
The script will run (and install, if you use Windows) FFmpeg and COLMAP as needed, followed by a conversion step to the required `transforms.json` format, which will be written in the current directory. The script will run (and install, if you use Windows) FFmpeg and COLMAP as needed, followed by a conversion step to the required `transforms.json` format, which will be written in the current directory.
...@@ -104,7 +104,12 @@ The script will run (and install, if you use Windows) FFmpeg and COLMAP as neede ...@@ -104,7 +104,12 @@ The script will run (and install, if you use Windows) FFmpeg and COLMAP as neede
By default, the script invokes colmap with the "sequential matcher", which is suitable for images taken from a smoothly changing camera path, as in a video. The exhaustive matcher is more appropriate if the images are in no particular order, as shown in the image example above. By default, the script invokes colmap with the "sequential matcher", which is suitable for images taken from a smoothly changing camera path, as in a video. The exhaustive matcher is more appropriate if the images are in no particular order, as shown in the image example above.
For more options, you can run the script with `--help`. For more advanced uses of COLMAP or for challenging scenes, please see the [COLMAP documentation](https://colmap.github.io/cli.html); you may need to modify the [scripts/colmap2nerf.py](/scripts/colmap2nerf.py) script itself. For more options, you can run the script with `--help`. For more advanced uses of COLMAP or for challenging scenes, please see the [COLMAP documentation](https://colmap.github.io/cli.html); you may need to modify the [scripts/colmap2nerf.py](/scripts/colmap2nerf.py) script itself.
The `aabb_scale` parameter is the most important __instant-ngp__ specific parameter. It specifies the extent of the scene, defaulting to 1; that is, the scene is scaled such that the camera positions are at an average distance of 1 unit from the origin. For small synthetic scenes such as the original NeRF dataset, the default `aabb_scale` of 1 is ideal and leads to fastest training. The NeRF model makes the assumption that the training images can entirely be explained by a scene contained within this bounding box. However, for natural scenes where there is a background that extends beyond this bounding box, the NeRF model will struggle and may hallucinate "floaters" at the boundaries of the box. By setting `aabb_scale` to a larger power of 2 (up to a maximum of 16), the NeRF model will extend rays to a much larger bounding box. Note that this can impact training speed slightly. If in doubt, for natural scenes, start with an `aabb_scale` of 128, and subsequently reduce it if possible. The value can be directly edited in the `transforms.json` output file, without re-running the [scripts/colmap2nerf.py](/scripts/colmap2nerf.py) script. The `aabb_scale` parameter is the most important __instant-ngp__ specific parameter. It specifies the extent of the scene, defaulting to 1; that is, the scene is scaled such that the camera positions are at an average distance of 1 unit from the origin. For small synthetic scenes such as the original NeRF dataset, the default `aabb_scale` of 1 is ideal and leads to fastest training. The NeRF model makes the assumption that the training images can entirely be explained by a scene contained within this bounding box. However, for natural scenes where there is a background that extends beyond this bounding box, the NeRF model will struggle and may hallucinate "floaters" at the boundaries of the box. By setting `aabb_scale` to a larger power of 2 (up to a maximum of 128), the NeRF model will extend rays to a much larger bounding box. Note that this can impact training speed slightly. If in doubt, for natural scenes, start with an `aabb_scale` of 128, and subsequently reduce it if possible. The value can be directly edited in the `transforms.json` output file, without re-running the [scripts/colmap2nerf.py](/scripts/colmap2nerf.py) script.
You can optionally pass in object categories (e.g. `--mask_categories person car`) which runs [Detectron2](https://github.com/facebookresearch/detectron2) to generate masks automatically.
__instant-ngp__ will not use the masked pixels for training.
This utility is helpful for users who wish to ignore moving or sensitive objects such as people, cars, or bikes.
See [scripts/category2id.json](/scripts/category2id.json) for a list of categories.
Assuming success, you can now train your NeRF model as follows, starting in the __instant-ngp__ folder: Assuming success, you can now train your NeRF model as follows, starting in the __instant-ngp__ folder:
......
include/neural-graphics-primitives/camera_path.h
+ 4
− 4
View file @ 3e46bc0c
...@@ -126,14 +126,14 @@ struct CameraPath { ...@@ -126,14 +126,14 @@ struct CameraPath {
return spline(t-floorf(t), get_keyframe(t1-1), get_keyframe(t1), get_keyframe(t1+1), get_keyframe(t1+2)); return spline(t-floorf(t), get_keyframe(t1-1), get_keyframe(t1), get_keyframe(t1+1), get_keyframe(t1+2));
} }
void save(const std::string& filepath_string); void save(const fs::path& path);
void load(const std::string& filepath_string, const Eigen::Matrix<float, 3, 4> &first_xform); void load(const fs::path& path, const Eigen::Matrix<float, 3, 4> &first_xform);
#ifdef NGP_GUI #ifdef NGP_GUI
ImGuizmo::MODE m_gizmo_mode = ImGuizmo::LOCAL; ImGuizmo::MODE m_gizmo_mode = ImGuizmo::LOCAL;
ImGuizmo::OPERATION m_gizmo_op = ImGuizmo::TRANSLATE; ImGuizmo::OPERATION m_gizmo_op = ImGuizmo::TRANSLATE;
int imgui(char path_filename_buf[128], float frame_milliseconds, Eigen::Matrix<float, 3, 4>& camera, float slice_plane_z, float scale, float fov, float aperture_size, float bounding_radius, const Eigen::Matrix<float, 3, 4>& first_xform, int glow_mode, float glow_y_cutoff); int imgui(char path_filename_buf[1024], float frame_milliseconds, Eigen::Matrix<float, 3, 4>& camera, float slice_plane_z, float scale, float fov, float aperture_size, float bounding_radius, const Eigen::Matrix<float, 3, 4>& first_xform, int glow_mode, float glow_y_cutoff);
bool imgui_viz(ImDrawList* list, Eigen::Matrix<float, 4, 4>& view2proj, Eigen::Matrix<float, 4, 4>& world2proj, Eigen::Matrix<float, 4, 4>& world2view, Eigen::Vector2f focal, float aspect); bool imgui_viz(ImDrawList* list, Eigen::Matrix<float, 4, 4>& view2proj, Eigen::Matrix<float, 4, 4>& world2proj, Eigen::Matrix<float, 4, 4>& world2view, Eigen::Vector2f focal, float aspect, float znear, float zfar);
#endif #endif
}; };
......
include/neural-graphics-primitives/common.h
+ 67
− 3
View file @ 3e46bc0c
...@@ -63,10 +63,20 @@ ...@@ -63,10 +63,20 @@
NGP_NAMESPACE_BEGIN NGP_NAMESPACE_BEGIN
namespace fs = filesystem;
bool is_wsl(); bool is_wsl();
filesystem::path get_executable_dir(); fs::path get_executable_dir();
filesystem::path get_root_dir(); fs::path get_root_dir();
#ifdef _WIN32
std::string utf16_to_utf8(const std::wstring& utf16);
std::wstring utf8_to_utf16(const std::string& utf16);
std::wstring native_string(const fs::path& path);
#else
std::string native_string(const fs::path& path);
#endif
bool ends_with(const std::string& str, const std::string& ending); bool ends_with(const std::string& str, const std::string& ending);
bool ends_with_case_insensitive(const std::string& str, const std::string& ending); bool ends_with_case_insensitive(const std::string& str, const std::string& ending);
...@@ -222,8 +232,13 @@ enum class ELensMode : int { ...@@ -222,8 +232,13 @@ enum class ELensMode : int {
FTheta, FTheta,
LatLong, LatLong,
OpenCVFisheye, OpenCVFisheye,
Equirectangular,
}; };
static constexpr const char* LensModeStr = "Perspective\0OpenCV\0F-Theta\0LatLong\0OpenCV Fisheye\0\0"; static constexpr const char* LensModeStr = "Perspective\0OpenCV\0F-Theta\0LatLong\0OpenCV Fisheye\0Equirectangular\0\0";
inline bool supports_dlss(ELensMode mode) {
return mode == ELensMode::Perspective || mode == ELensMode::OpenCV || mode == ELensMode::OpenCVFisheye;
}
struct Lens { struct Lens {
ELensMode mode = ELensMode::Perspective; ELensMode mode = ELensMode::Perspective;
...@@ -333,4 +348,53 @@ private: ...@@ -333,4 +348,53 @@ private:
std::chrono::time_point<std::chrono::steady_clock> m_creation_time; std::chrono::time_point<std::chrono::steady_clock> m_creation_time;
}; };
template <typename T>
struct Buffer2DView {
T* data = nullptr;
Eigen::Vector2i resolution = Eigen::Vector2i::Zero();
// Lookup via integer pixel position (no bounds checking)
NGP_HOST_DEVICE T at(const Eigen::Vector2i& xy) const {
return data[xy.x() + xy.y() * resolution.x()];
}
// Lookup via UV coordinates in [0,1]^2
NGP_HOST_DEVICE T at(const Eigen::Vector2f& uv) const {
Eigen::Vector2i xy = resolution.cast<float>().cwiseProduct(uv).cast<int>().cwiseMax(0).cwiseMin(resolution - Eigen::Vector2i::Ones());
return at(xy);
}
// Lookup via UV coordinates in [0,1]^2 and LERP the nearest texels
NGP_HOST_DEVICE T at_lerp(const Eigen::Vector2f& uv) const {
const Eigen::Vector2f xy_float = resolution.cast<float>().cwiseProduct(uv);
const Eigen::Vector2i xy = xy_float.cast<int>();
const Eigen::Vector2f weight = xy_float - xy.cast<float>();
auto read_val = [&](Eigen::Vector2i pos) {
pos = pos.cwiseMax(0).cwiseMin(resolution - Eigen::Vector2i::Ones());
return at(pos);
};
return (
(1 - weight.x()) * (1 - weight.y()) * read_val({xy.x(), xy.y()}) +
(weight.x()) * (1 - weight.y()) * read_val({xy.x()+1, xy.y()}) +
(1 - weight.x()) * (weight.y()) * read_val({xy.x(), xy.y()+1}) +
(weight.x()) * (weight.y()) * read_val({xy.x()+1, xy.y()+1})
);
}
NGP_HOST_DEVICE operator bool() const {
return data;
}
};
uint8_t* load_stbi(const fs::path& path, int* width, int* height, int* comp, int req_comp);
float* load_stbi_float(const fs::path& path, int* width, int* height, int* comp, int req_comp);
uint16_t* load_stbi_16(const fs::path& path, int* width, int* height, int* comp, int req_comp);
bool is_hdr_stbi(const fs::path& path);
int write_stbi(const fs::path& path, int width, int height, int comp, const uint8_t* pixels, int quality = 100);
FILE* native_fopen(const fs::path& path, const char* mode);
NGP_NAMESPACE_END NGP_NAMESPACE_END
include/neural-graphics-primitives/common_device.cuh
+ 291
− 138
View file @ 3e46bc0c
...@@ -28,6 +28,52 @@ NGP_NAMESPACE_BEGIN ...@@ -28,6 +28,52 @@ NGP_NAMESPACE_BEGIN
using precision_t = tcnn::network_precision_t; using precision_t = tcnn::network_precision_t;
// The maximum depth that can be produced when rendering a frame.
// Chosen somewhat low (rather than std::numeric_limits<float>::infinity())
// to permit numerically stable reprojection and DLSS operation,
// even when rendering the infinitely distant horizon.
inline constexpr __device__ float MAX_DEPTH() { return 16384.0f; }
template <typename T>
class Buffer2D {
public:
Buffer2D() = default;
Buffer2D(const Eigen::Vector2i& resolution) {
resize(resolution);
}
T* data() const {
return m_data.data();
}
size_t bytes() const {
return m_data.bytes();
}
void resize(const Eigen::Vector2i& resolution) {
m_data.resize(resolution.prod());
m_resolution = resolution;
}
const Eigen::Vector2i& resolution() const {
return m_resolution;
}
Buffer2DView<T> view() const {
// Row major for now.
return {data(), m_resolution};
}
Buffer2DView<const T> const_view() const {
// Row major for now.
return {data(), m_resolution};
}
private:
tcnn::GPUMemory<T> m_data;
Eigen::Vector2i m_resolution;
};
inline NGP_HOST_DEVICE float srgb_to_linear(float srgb) { inline NGP_HOST_DEVICE float srgb_to_linear(float srgb) {
if (srgb <= 0.04045f) { if (srgb <= 0.04045f) {
return srgb / 12.92f; return srgb / 12.92f;
...@@ -76,42 +122,9 @@ inline NGP_HOST_DEVICE Eigen::Array3f linear_to_srgb_derivative(const Eigen::Arr ...@@ -76,42 +122,9 @@ inline NGP_HOST_DEVICE Eigen::Array3f linear_to_srgb_derivative(const Eigen::Arr
return {linear_to_srgb_derivative(x.x()), linear_to_srgb_derivative(x.y()), (linear_to_srgb_derivative(x.z()))}; return {linear_to_srgb_derivative(x.x()), linear_to_srgb_derivative(x.y()), (linear_to_srgb_derivative(x.z()))};
} }
template <uint32_t N_DIMS, typename T>
NGP_HOST_DEVICE Eigen::Matrix<float, N_DIMS, 1> read_image(const T* __restrict__ data, const Eigen::Vector2i& resolution, const Eigen::Vector2f& pos) {
const Eigen::Vector2f pos_float = Eigen::Vector2f{pos.x() * (float)(resolution.x()-1), pos.y() * (float)(resolution.y()-1)};
const Eigen::Vector2i texel = pos_float.cast<int>();
const Eigen::Vector2f weight = pos_float - texel.cast<float>();
auto read_val = [&](Eigen::Vector2i pos) {
pos.x() = std::max(std::min(pos.x(), resolution.x()-1), 0);
pos.y() = std::max(std::min(pos.y(), resolution.y()-1), 0);
Eigen::Matrix<float, N_DIMS, 1> result;
if (std::is_same<T, float>::value) {
result = *(Eigen::Matrix<T, N_DIMS, 1>*)&data[(pos.x() + pos.y() * resolution.x()) * N_DIMS];
} else {
auto val = *(tcnn::vector_t<T, N_DIMS>*)&data[(pos.x() + pos.y() * resolution.x()) * N_DIMS];
NGP_PRAGMA_UNROLL
for (uint32_t i = 0; i < N_DIMS; ++i) {
result[i] = (float)val[i];
}
}
return result;
};
return (
(1 - weight.x()) * (1 - weight.y()) * read_val({texel.x(), texel.y()}) +
(weight.x()) * (1 - weight.y()) * read_val({texel.x()+1, texel.y()}) +
(1 - weight.x()) * (weight.y()) * read_val({texel.x(), texel.y()+1}) +
(weight.x()) * (weight.y()) * read_val({texel.x()+1, texel.y()+1})
);
}
template <uint32_t N_DIMS, typename T> template <uint32_t N_DIMS, typename T>
__device__ void deposit_image_gradient(const Eigen::Matrix<float, N_DIMS, 1>& value, T* __restrict__ gradient, T* __restrict__ gradient_weight, const Eigen::Vector2i& resolution, const Eigen::Vector2f& pos) { __device__ void deposit_image_gradient(const Eigen::Matrix<float, N_DIMS, 1>& value, T* __restrict__ gradient, T* __restrict__ gradient_weight, const Eigen::Vector2i& resolution, const Eigen::Vector2f& pos) {
const Eigen::Vector2f pos_float = Eigen::Vector2f{pos.x() * (resolution.x()-1), pos.y() * (resolution.y()-1)}; const Eigen::Vector2f pos_float = resolution.cast<float>().cwiseProduct(pos);
const Eigen::Vector2i texel = pos_float.cast<int>(); const Eigen::Vector2i texel = pos_float.cast<int>();
const Eigen::Vector2f weight = pos_float - texel.cast<float>(); const Eigen::Vector2f weight = pos_float - texel.cast<float>();
...@@ -142,6 +155,138 @@ __device__ void deposit_image_gradient(const Eigen::Matrix<float, N_DIMS, 1>& va ...@@ -142,6 +155,138 @@ __device__ void deposit_image_gradient(const Eigen::Matrix<float, N_DIMS, 1>& va
deposit_val(value, (weight.x()) * (weight.y()), {texel.x()+1, texel.y()+1}); deposit_val(value, (weight.x()) * (weight.y()), {texel.x()+1, texel.y()+1});
} }
struct FoveationPiecewiseQuadratic {
NGP_HOST_DEVICE FoveationPiecewiseQuadratic() = default;
FoveationPiecewiseQuadratic(float center_pixel_steepness, float center_inverse_piecewise_y, float center_radius) {
float center_inverse_radius = center_radius * center_pixel_steepness;
float left_inverse_piecewise_switch = center_inverse_piecewise_y - center_inverse_radius;
float right_inverse_piecewise_switch = center_inverse_piecewise_y + center_inverse_radius;
if (left_inverse_piecewise_switch < 0) {
left_inverse_piecewise_switch = 0.0f;
}
if (right_inverse_piecewise_switch > 1) {
right_inverse_piecewise_switch = 1.0f;
}
float am = center_pixel_steepness;
float d = (right_inverse_piecewise_switch - left_inverse_piecewise_switch) / center_pixel_steepness / 2;
// binary search for l,r,bm since analytical is very complex
float bm;
float m_min = 0.0f;
float m_max = 1.0f;
for (uint32_t i = 0; i < 20; i++) {
float m = (m_min + m_max) / 2.0f;
float l = m - d;
float r = m + d;
bm = -((am - 1) * l * l) / (r * r - 2 * r + l * l + 1);
float l_actual = (left_inverse_piecewise_switch - bm) / am;
float r_actual = (right_inverse_piecewise_switch - bm) / am;
float m_actual = (l_actual + r_actual) / 2;
if (m_actual > m) {
m_min = m;
} else {
m_max = m;
}
}
float l = (left_inverse_piecewise_switch - bm) / am;
float r = (right_inverse_piecewise_switch - bm) / am;
// Full linear case. Default construction covers this.
if ((l == 0.0f && r == 1.0f) || (am == 1.0f)) {
return;
}
// write out solution
switch_left = l;
switch_right = r;
this->am = am;
al = (am - 1) / (r * r - 2 * r + l * l + 1);
bl = (am * (r * r - 2 * r + 1) + am * l * l + (2 - 2 * am) * l) / (r * r - 2 * r + l * l + 1);
cl = 0;
this->bm = bm = -((am - 1) * l * l) / (r * r - 2 * r + l * l + 1);
ar = -(am - 1) / (r * r - 2 * r + l * l + 1);
br = (am * (r * r + 1) - 2 * r + am * l * l) / (r * r - 2 * r + l * l + 1);
cr = -(am * r * r - r * r + (am - 1) * l * l) / (r * r - 2 * r + l * l + 1);
inv_switch_left = am * switch_left + bm;
inv_switch_right = am * switch_right + bm;
}
// left parabola: al * x^2 + bl * x + cl
float al = 0.0f, bl = 0.0f, cl = 0.0f;
// middle linear piece: am * x + bm. am should give 1:1 pixel mapping between warped size and full size.
float am = 1.0f, bm = 0.0f;
// right parabola: al * x^2 + bl * x + cl
float ar = 0.0f, br = 0.0f, cr = 0.0f;
// points where left and right switch over from quadratic to linear
float switch_left = 0.0f, switch_right = 1.0f;
// same, in inverted space
float inv_switch_left = 0.0f, inv_switch_right = 1.0f;
NGP_HOST_DEVICE float warp(float x) const {
x = tcnn::clamp(x, 0.0f, 1.0f);
if (x < switch_left) {
return al * x * x + bl * x + cl;
} else if (x > switch_right) {
return ar * x * x + br * x + cr;
} else {
return am * x + bm;
}
}
NGP_HOST_DEVICE float unwarp(float y) const {
y = tcnn::clamp(y, 0.0f, 1.0f);
if (y < inv_switch_left) {
return (std::sqrt(-4 * al * cl + 4 * al * y + bl * bl) - bl) / (2 * al);
} else if (y > inv_switch_right) {
return (std::sqrt(-4 * ar * cr + 4 * ar * y + br * br) - br) / (2 * ar);
} else {
return (y - bm) / am;
}
}
NGP_HOST_DEVICE float density(float x) const {
x = tcnn::clamp(x, 0.0f, 1.0f);
if (x < switch_left) {
return 2 * al * x + bl;
} else if (x > switch_right) {
return 2 * ar * x + br;
} else {
return am;
}
}
};
struct Foveation {
NGP_HOST_DEVICE Foveation() = default;
Foveation(const Eigen::Vector2f& center_pixel_steepness, const Eigen::Vector2f& center_inverse_piecewise_y, const Eigen::Vector2f& center_radius)
: warp_x{center_pixel_steepness.x(), center_inverse_piecewise_y.x(), center_radius.x()}, warp_y{center_pixel_steepness.y(), center_inverse_piecewise_y.y(), center_radius.y()} {}
FoveationPiecewiseQuadratic warp_x, warp_y;
NGP_HOST_DEVICE Eigen::Vector2f warp(const Eigen::Vector2f& x) const {
return {warp_x.warp(x.x()), warp_y.warp(x.y())};
}
NGP_HOST_DEVICE Eigen::Vector2f unwarp(const Eigen::Vector2f& y) const {
return {warp_x.unwarp(y.x()), warp_y.unwarp(y.y())};
}
NGP_HOST_DEVICE float density(const Eigen::Vector2f& x) const {
return warp_x.density(x.x()) * warp_y.density(x.y());
}
};
template <typename T> template <typename T>
NGP_HOST_DEVICE inline void opencv_lens_distortion_delta(const T* extra_params, const T u, const T v, T* du, T* dv) { NGP_HOST_DEVICE inline void opencv_lens_distortion_delta(const T* extra_params, const T u, const T v, T* du, T* dv) {
const T k1 = extra_params[0]; const T k1 = extra_params[0];
...@@ -292,37 +437,53 @@ inline NGP_HOST_DEVICE Eigen::Vector3f latlong_to_dir(const Eigen::Vector2f& uv) ...@@ -292,37 +437,53 @@ inline NGP_HOST_DEVICE Eigen::Vector3f latlong_to_dir(const Eigen::Vector2f& uv)
return {sp * ct, st, cp * ct}; return {sp * ct, st, cp * ct};
} }
inline NGP_HOST_DEVICE Ray pixel_to_ray( inline NGP_HOST_DEVICE Eigen::Vector3f equirectangular_to_dir(const Eigen::Vector2f& uv) {
float ct = (uv.y() - 0.5f) * 2.0f;
float st = std::sqrt(std::max(1.0f - ct * ct, 0.0f));
float phi = (uv.x() - 0.5f) * PI() * 2.0f;
float sp, cp;
sincosf(phi, &sp, &cp);
return {sp * st, ct, cp * st};
}
inline NGP_HOST_DEVICE Ray uv_to_ray(
uint32_t spp, uint32_t spp,
const Eigen::Vector2i& pixel, const Eigen::Vector2f& uv,
const Eigen::Vector2i& resolution, const Eigen::Vector2i& resolution,
const Eigen::Vector2f& focal_length, const Eigen::Vector2f& focal_length,
const Eigen::Matrix<float, 3, 4>& camera_matrix, const Eigen::Matrix<float, 3, 4>& camera_matrix,
const Eigen::Vector2f& screen_center, const Eigen::Vector2f& screen_center,
const Eigen::Vector3f& parallax_shift, const Eigen::Vector3f& parallax_shift = Eigen::Vector3f::Zero(),
bool snap_to_pixel_centers = false,
float near_distance = 0.0f, float near_distance = 0.0f,
float focus_z = 1.0f, float focus_z = 1.0f,
float aperture_size = 0.0f, float aperture_size = 0.0f,
const Foveation& foveation = {},
Buffer2DView<const uint8_t> hidden_area_mask = {},
const Lens& lens = {}, const Lens& lens = {},
const float* __restrict__ distortion_grid = nullptr, Buffer2DView<const Eigen::Vector2f> distortion = {}
const Eigen::Vector2i distortion_grid_resolution = Eigen::Vector2i::Zero()
) { ) {
Eigen::Vector2f offset = ld_random_pixel_offset(snap_to_pixel_centers ? 0 : spp); Eigen::Vector2f warped_uv = foveation.warp(uv);
Eigen::Vector2f uv = (pixel.cast<float>() + offset).cwiseQuotient(resolution.cast<float>());
// Check the hidden area mask _after_ applying foveation, because foveation will be undone
// before blitting to the framebuffer to which the hidden area mask corresponds.
if (hidden_area_mask && !hidden_area_mask.at(warped_uv)) {
return Ray::invalid();
}
Eigen::Vector3f dir; Eigen::Vector3f dir;
if (lens.mode == ELensMode::FTheta) { if (lens.mode == ELensMode::FTheta) {
dir = f_theta_undistortion(uv - screen_center, lens.params, {1000.f, 0.f, 0.f}); dir = f_theta_undistortion(warped_uv - screen_center, lens.params, {0.f, 0.f, 0.f});
if (dir.x() == 1000.f) { if (dir == Eigen::Vector3f::Zero()) {
return {{1000.f, 0.f, 0.f}, {0.f, 0.f, 1.f}}; // return a point outside the aabb so the pixel is not rendered return Ray::invalid();
} }
} else if (lens.mode == ELensMode::LatLong) { } else if (lens.mode == ELensMode::LatLong) {
dir = latlong_to_dir(uv); dir = latlong_to_dir(warped_uv);
} else if (lens.mode == ELensMode::Equirectangular) {
dir = equirectangular_to_dir(warped_uv);
} else { } else {
dir = { dir = {
(uv.x() - screen_center.x()) * (float)resolution.x() / focal_length.x(), (warped_uv.x() - screen_center.x()) * (float)resolution.x() / focal_length.x(),
(uv.y() - screen_center.y()) * (float)resolution.y() / focal_length.y(), (warped_uv.y() - screen_center.y()) * (float)resolution.y() / focal_length.y(),
1.0f 1.0f
}; };
...@@ -332,8 +493,9 @@ inline NGP_HOST_DEVICE Ray pixel_to_ray( ...@@ -332,8 +493,9 @@ inline NGP_HOST_DEVICE Ray pixel_to_ray(
iterative_opencv_fisheye_lens_undistortion(lens.params, &dir.x(), &dir.y()); iterative_opencv_fisheye_lens_undistortion(lens.params, &dir.x(), &dir.y());
} }
} }
if (distortion_grid) {
dir.head<2>() += read_image<2>(distortion_grid, distortion_grid_resolution, uv); if (distortion) {
dir.head<2>() += distortion.at_lerp(warped_uv);
} }
Eigen::Vector3f head_pos = {parallax_shift.x(), parallax_shift.y(), 0.f}; Eigen::Vector3f head_pos = {parallax_shift.x(), parallax_shift.y(), 0.f};
...@@ -341,26 +503,60 @@ inline NGP_HOST_DEVICE Ray pixel_to_ray( ...@@ -341,26 +503,60 @@ inline NGP_HOST_DEVICE Ray pixel_to_ray(
dir = camera_matrix.block<3, 3>(0, 0) * dir; dir = camera_matrix.block<3, 3>(0, 0) * dir;
Eigen::Vector3f origin = camera_matrix.block<3, 3>(0, 0) * head_pos + camera_matrix.col(3); Eigen::Vector3f origin = camera_matrix.block<3, 3>(0, 0) * head_pos + camera_matrix.col(3);
if (aperture_size != 0.0f) {
if (aperture_size > 0.0f) {
Eigen::Vector3f lookat = origin + dir * focus_z; Eigen::Vector3f lookat = origin + dir * focus_z;
Eigen::Vector2f blur = aperture_size * square2disk_shirley(ld_random_val_2d(spp, (uint32_t)pixel.x() * 19349663 + (uint32_t)pixel.y() * 96925573) * 2.0f - Eigen::Vector2f::Ones()); Eigen::Vector2f blur = aperture_size * square2disk_shirley(ld_random_val_2d(spp, uv.cwiseProduct(resolution.cast<float>()).cast<int>().dot(Eigen::Vector2i{19349663, 96925573})) * 2.0f - Eigen::Vector2f::Ones());
origin += camera_matrix.block<3, 2>(0, 0) * blur; origin += camera_matrix.block<3, 2>(0, 0) * blur;
dir = (lookat - origin) / focus_z; dir = (lookat - origin) / focus_z;
} }
origin += dir * near_distance;
origin += dir * near_distance;
return {origin, dir}; return {origin, dir};
} }
inline NGP_HOST_DEVICE Eigen::Vector2f pos_to_pixel( inline NGP_HOST_DEVICE Ray pixel_to_ray(
uint32_t spp,
const Eigen::Vector2i& pixel,
const Eigen::Vector2i& resolution,
const Eigen::Vector2f& focal_length,
const Eigen::Matrix<float, 3, 4>& camera_matrix,
const Eigen::Vector2f& screen_center,
const Eigen::Vector3f& parallax_shift = Eigen::Vector3f::Zero(),
bool snap_to_pixel_centers = false,
float near_distance = 0.0f,
float focus_z = 1.0f,
float aperture_size = 0.0f,
const Foveation& foveation = {},
Buffer2DView<const uint8_t> hidden_area_mask = {},
const Lens& lens = {},
Buffer2DView<const Eigen::Vector2f> distortion = {}
) {
return uv_to_ray(
spp,
(pixel.cast<float>() + ld_random_pixel_offset(snap_to_pixel_centers ? 0 : spp)).cwiseQuotient(resolution.cast<float>()),
resolution,
focal_length,
camera_matrix,
screen_center,
parallax_shift,
near_distance,
focus_z,
aperture_size,
foveation,
hidden_area_mask,
lens,
distortion
);
}
inline NGP_HOST_DEVICE Eigen::Vector2f pos_to_uv(
const Eigen::Vector3f& pos, const Eigen::Vector3f& pos,
const Eigen::Vector2i& resolution, const Eigen::Vector2i& resolution,
const Eigen::Vector2f& focal_length, const Eigen::Vector2f& focal_length,
const Eigen::Matrix<float, 3, 4>& camera_matrix, const Eigen::Matrix<float, 3, 4>& camera_matrix,
const Eigen::Vector2f& screen_center, const Eigen::Vector2f& screen_center,
const Eigen::Vector3f& parallax_shift, const Eigen::Vector3f& parallax_shift,
const Foveation& foveation = {},
const Lens& lens = {} const Lens& lens = {}
) { ) {
// Express ray in terms of camera frame // Express ray in terms of camera frame
...@@ -386,10 +582,23 @@ inline NGP_HOST_DEVICE Eigen::Vector2f pos_to_pixel( ...@@ -386,10 +582,23 @@ inline NGP_HOST_DEVICE Eigen::Vector2f pos_to_pixel(
dir.y() += dv; dir.y() += dv;
Eigen::Vector2f uv = Eigen::Vector2f{dir.x(), dir.y()}.cwiseProduct(focal_length).cwiseQuotient(resolution.cast<float>()) + screen_center; Eigen::Vector2f uv = Eigen::Vector2f{dir.x(), dir.y()}.cwiseProduct(focal_length).cwiseQuotient(resolution.cast<float>()) + screen_center;
return uv.cwiseProduct(resolution.cast<float>()); return foveation.unwarp(uv);
} }
inline NGP_HOST_DEVICE Eigen::Vector2f motion_vector_3d( inline NGP_HOST_DEVICE Eigen::Vector2f pos_to_pixel(
const Eigen::Vector3f& pos,
const Eigen::Vector2i& resolution,
const Eigen::Vector2f& focal_length,
const Eigen::Matrix<float, 3, 4>& camera_matrix,
const Eigen::Vector2f& screen_center,
const Eigen::Vector3f& parallax_shift,
const Foveation& foveation = {},
const Lens& lens = {}
) {
return pos_to_uv(pos, resolution, focal_length, camera_matrix, screen_center, parallax_shift, foveation, lens).cwiseProduct(resolution.cast<float>());
}
inline NGP_HOST_DEVICE Eigen::Vector2f motion_vector(
const uint32_t sample_index, const uint32_t sample_index,
const Eigen::Vector2i& pixel, const Eigen::Vector2i& pixel,
const Eigen::Vector2i& resolution, const Eigen::Vector2i& resolution,
...@@ -400,6 +609,8 @@ inline NGP_HOST_DEVICE Eigen::Vector2f motion_vector_3d( ...@@ -400,6 +609,8 @@ inline NGP_HOST_DEVICE Eigen::Vector2f motion_vector_3d(
const Eigen::Vector3f& parallax_shift, const Eigen::Vector3f& parallax_shift,
const bool snap_to_pixel_centers, const bool snap_to_pixel_centers,
const float depth, const float depth,
const Foveation& foveation = {},
const Foveation& prev_foveation = {},
const Lens& lens = {} const Lens& lens = {}
) { ) {
Ray ray = pixel_to_ray( Ray ray = pixel_to_ray(
...@@ -414,98 +625,39 @@ inline NGP_HOST_DEVICE Eigen::Vector2f motion_vector_3d( ...@@ -414,98 +625,39 @@ inline NGP_HOST_DEVICE Eigen::Vector2f motion_vector_3d(
0.0f, 0.0f,
1.0f, 1.0f,
0.0f, 0.0f,
lens, foveation,
nullptr, {}, // No hidden area mask
Eigen::Vector2i::Zero() lens
); );
Eigen::Vector2f prev_pixel = pos_to_pixel( Eigen::Vector2f prev_pixel = pos_to_pixel(
ray.o + ray.d * depth, ray(depth),
resolution, resolution,
focal_length, focal_length,
prev_camera, prev_camera,
screen_center, screen_center,
parallax_shift, parallax_shift,
prev_foveation,
lens lens
); );
return prev_pixel - (pixel.cast<float>() + ld_random_pixel_offset(sample_index)); return prev_pixel - (pixel.cast<float>() + ld_random_pixel_offset(sample_index));
} }
inline NGP_HOST_DEVICE Eigen::Vector2f pixel_to_image_uv( // Maps view-space depth (physical units) in the range [znear, zfar] hyperbolically to
const uint32_t sample_index, // the interval [1, 0]. This is the reverse-z-component of "normalized device coordinates",
const Eigen::Vector2i& pixel, // which are commonly used in rasterization, where linear interpolation in screen space
const Eigen::Vector2i& resolution, // has to be equivalent to linear interpolation in real space (which, in turn, is
const Eigen::Vector2i& image_resolution, // guaranteed by the hyperbolic mapping of depth). This format is commonly found in
const Eigen::Vector2f& screen_center, // z-buffers, and hence expected by downstream image processing functions, such as DLSS
const float view_dist, // and VR reprojection.
const Eigen::Vector2f& image_pos, inline NGP_HOST_DEVICE float to_ndc_depth(float z, float n, float f) {
const bool snap_to_pixel_centers // View depth outside of the view frustum leads to output outside of [0, 1]
) { z = tcnn::clamp(z, n, f);
Eigen::Vector2f jit = ld_random_pixel_offset(snap_to_pixel_centers ? 0 : sample_index);
Eigen::Vector2f offset = screen_center.cwiseProduct(resolution.cast<float>()) + jit;
float y_scale = view_dist;
float x_scale = y_scale * resolution.x() / resolution.y();
return {
((x_scale * (pixel.x() + offset.x())) / resolution.x() - view_dist * image_pos.x()) / image_resolution.x() * image_resolution.y(),
(y_scale * (pixel.y() + offset.y())) / resolution.y() - view_dist * image_pos.y()
};
}
inline NGP_HOST_DEVICE Eigen::Vector2f image_uv_to_pixel(
const Eigen::Vector2f& uv,
const Eigen::Vector2i& resolution,
const Eigen::Vector2i& image_resolution,
const Eigen::Vector2f& screen_center,
const float view_dist,
const Eigen::Vector2f& image_pos
) {
Eigen::Vector2f offset = screen_center.cwiseProduct(resolution.cast<float>());
float y_scale = view_dist;
float x_scale = y_scale * resolution.x() / resolution.y();
return { float scale = n / (n - f);
((uv.x() / image_resolution.y() * image_resolution.x()) + view_dist * image_pos.x()) * resolution.x() / x_scale - offset.x(), float bias = -f * scale;
(uv.y() + view_dist * image_pos.y()) * resolution.y() / y_scale - offset.y() return tcnn::clamp((z * scale + bias) / z, 0.0f, 1.0f);
};
}
inline NGP_HOST_DEVICE Eigen::Vector2f motion_vector_2d(
const uint32_t sample_index,
const Eigen::Vector2i& pixel,
const Eigen::Vector2i& resolution,
const Eigen::Vector2i& image_resolution,
const Eigen::Vector2f& screen_center,
const float view_dist,
const float prev_view_dist,
const Eigen::Vector2f& image_pos,
const Eigen::Vector2f& prev_image_pos,
const bool snap_to_pixel_centers
) {
Eigen::Vector2f uv = pixel_to_image_uv(
sample_index,
pixel,
resolution,
image_resolution,
screen_center,
view_dist,
image_pos,
snap_to_pixel_centers
);
Eigen::Vector2f prev_pixel = image_uv_to_pixel(
uv,
resolution,
image_resolution,
screen_center,
prev_view_dist,
prev_image_pos
);
return prev_pixel - (pixel.cast<float>() + ld_random_pixel_offset(sample_index));
} }
inline NGP_HOST_DEVICE float fov_to_focal_length(int resolution, float degrees) { inline NGP_HOST_DEVICE float fov_to_focal_length(int resolution, float degrees) {
...@@ -587,7 +739,8 @@ inline NGP_HOST_DEVICE void apply_quilting(uint32_t* x, uint32_t* y, const Eigen ...@@ -587,7 +739,8 @@ inline NGP_HOST_DEVICE void apply_quilting(uint32_t* x, uint32_t* y, const Eigen
if (quilting_dims == Eigen::Vector2i{2, 1}) { if (quilting_dims == Eigen::Vector2i{2, 1}) {
// Likely VR: parallax_shift.x() is the IPD in this case. The following code centers the camera matrix between both eyes. // Likely VR: parallax_shift.x() is the IPD in this case. The following code centers the camera matrix between both eyes.
parallax_shift.x() = idx ? (-0.5f * parallax_shift.x()) : (0.5f * parallax_shift.x()); // idx == 0 -> left eye -> -1/2 x
parallax_shift.x() = (idx == 0) ? (-0.5f * parallax_shift.x()) : (0.5f * parallax_shift.x());
} else { } else {
// Likely HoloPlay lenticular display: in this case, `parallax_shift.z()` is the inverse height of the head above the display. // Likely HoloPlay lenticular display: in this case, `parallax_shift.z()` is the inverse height of the head above the display.
// The following code computes the x-offset of views as a function of this. // The following code computes the x-offset of views as a function of this.
...@@ -756,7 +909,7 @@ inline NGP_HOST_DEVICE float read_depth(Eigen::Vector2f pos, const Eigen::Vector ...@@ -756,7 +909,7 @@ inline NGP_HOST_DEVICE float read_depth(Eigen::Vector2f pos, const Eigen::Vector
Eigen::Matrix<float, 3, 4> log_space_lerp(const Eigen::Matrix<float, 3, 4>& begin, const Eigen::Matrix<float, 3, 4>& end, float t); Eigen::Matrix<float, 3, 4> log_space_lerp(const Eigen::Matrix<float, 3, 4>& begin, const Eigen::Matrix<float, 3, 4>& end, float t);
tcnn::GPUMemory<float> load_exr(const std::string& filename, int& width, int& height); tcnn::GPUMemory<float> load_exr_gpu(const fs::path& path, int* width, int* height);
tcnn::GPUMemory<float> load_stbi(const std::string& filename, int& width, int& height); tcnn::GPUMemory<float> load_stbi_gpu(const fs::path& path, int* width, int* height);
NGP_NAMESPACE_END NGP_NAMESPACE_END
include/neural-graphics-primitives/dlss.h
+ 9
− 9
View file @ 3e46bc0c
...@@ -54,16 +54,16 @@ public: ...@@ -54,16 +54,16 @@ public:
virtual EDlssQuality quality() const = 0; virtual EDlssQuality quality() const = 0;
}; };
#ifdef NGP_VULKAN class IDlssProvider {
std::shared_ptr<IDlss> dlss_init(const Eigen::Vector2i& out_resolution); public:
virtual ~IDlssProvider() {}
void vulkan_and_ngx_init(); virtual size_t allocated_bytes() const = 0;
size_t dlss_allocated_bytes(); virtual std::unique_ptr<IDlss> init_dlss(const Eigen::Vector2i& out_resolution) = 0;
void vulkan_and_ngx_destroy(); };
#else
inline size_t dlss_allocated_bytes() { #ifdef NGP_VULKAN
return 0; std::shared_ptr<IDlssProvider> init_vulkan_and_ngx();
}
#endif #endif
NGP_NAMESPACE_END NGP_NAMESPACE_END
include/neural-graphics-primitives/envmap.cuh
+ 7
− 16
View file @ 3e46bc0c
...@@ -26,31 +26,22 @@ ...@@ -26,31 +26,22 @@
NGP_NAMESPACE_BEGIN NGP_NAMESPACE_BEGIN
template <typename T> inline __device__ Eigen::Array4f read_envmap(const Buffer2DView<const Eigen::Array4f>& envmap, const Eigen::Vector3f& dir) {
__device__ Eigen::Array4f read_envmap(const T* __restrict__ envmap_data, const Eigen::Vector2i envmap_resolution, const Eigen::Vector3f& dir) {
auto dir_cyl = dir_to_spherical_unorm({dir.z(), -dir.x(), dir.y()}); auto dir_cyl = dir_to_spherical_unorm({dir.z(), -dir.x(), dir.y()});
auto envmap_float = Eigen::Vector2f{dir_cyl.y() * (envmap_resolution.x()-1), dir_cyl.x() * (envmap_resolution.y()-1)}; auto envmap_float = Eigen::Vector2f{dir_cyl.y() * (envmap.resolution.x()-1), dir_cyl.x() * (envmap.resolution.y()-1)};
Eigen::Vector2i envmap_texel = envmap_float.cast<int>(); Eigen::Vector2i envmap_texel = envmap_float.cast<int>();
auto weight = envmap_float - envmap_texel.cast<float>(); auto weight = envmap_float - envmap_texel.cast<float>();
auto read_val = [&](Eigen::Vector2i pos) { auto read_val = [&](Eigen::Vector2i pos) {
if (pos.x() < 0) { if (pos.x() < 0) {
pos.x() += envmap_resolution.x(); pos.x() += envmap.resolution.x();
} else if (pos.x() >= envmap_resolution.x()) { } else if (pos.x() >= envmap.resolution.x()) {
pos.x() -= envmap_resolution.x(); pos.x() -= envmap.resolution.x();
}
pos.y() = std::max(std::min(pos.y(), envmap_resolution.y()-1), 0);
Eigen::Array4f result;
if (std::is_same<T, float>::value) {
result = *(Eigen::Array4f*)&envmap_data[(pos.x() + pos.y() * envmap_resolution.x()) * 4];
} else {
auto val = *(tcnn::vector_t<T, 4>*)&envmap_data[(pos.x() + pos.y() * envmap_resolution.x()) * 4];
result = {(float)val[0], (float)val[1], (float)val[2], (float)val[3]};
} }
return result; pos.y() = std::max(std::min(pos.y(), envmap.resolution.y()-1), 0);
return envmap.at(pos);
}; };
auto result = ( auto result = (
......
include/neural-graphics-primitives/json_binding.h
+ 4
− 0
View file @ 3e46bc0c
...@@ -166,6 +166,7 @@ inline void to_json(nlohmann::json& j, const NerfDataset& dataset) { ...@@ -166,6 +166,7 @@ inline void to_json(nlohmann::json& j, const NerfDataset& dataset) {
j["from_mitsuba"] = dataset.from_mitsuba; j["from_mitsuba"] = dataset.from_mitsuba;
j["is_hdr"] = dataset.is_hdr; j["is_hdr"] = dataset.is_hdr;
j["wants_importance_sampling"] = dataset.wants_importance_sampling; j["wants_importance_sampling"] = dataset.wants_importance_sampling;
j["n_extra_learnable_dims"] = dataset.n_extra_learnable_dims;
} }
inline void from_json(const nlohmann::json& j, NerfDataset& dataset) { inline void from_json(const nlohmann::json& j, NerfDataset& dataset) {
...@@ -209,11 +210,14 @@ inline void from_json(const nlohmann::json& j, NerfDataset& dataset) { ...@@ -209,11 +210,14 @@ inline void from_json(const nlohmann::json& j, NerfDataset& dataset) {
dataset.aabb_scale = j.at("aabb_scale"); dataset.aabb_scale = j.at("aabb_scale");
dataset.from_mitsuba = j.at("from_mitsuba"); dataset.from_mitsuba = j.at("from_mitsuba");
dataset.is_hdr = j.value("is_hdr", false); dataset.is_hdr = j.value("is_hdr", false);
if (j.contains("wants_importance_sampling")) { if (j.contains("wants_importance_sampling")) {
dataset.wants_importance_sampling = j.at("wants_importance_sampling"); dataset.wants_importance_sampling = j.at("wants_importance_sampling");
} else { } else {
dataset.wants_importance_sampling = true; dataset.wants_importance_sampling = true;
} }
dataset.n_extra_learnable_dims = j.value("n_extra_learnable_dims", 0);
} }
NGP_NAMESPACE_END NGP_NAMESPACE_END
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment