nvidia/gvdb-voxels

Sparse volume compute and rendering on NVIDIA GPUs

NVIDIA® GVDB VOXELS Release 1.1.1

OVERVIEW

NVIDIA® GVDB Voxels is a new library and SDK for simulation, compute, and rendering of sparse volumetric data. Details on the GVDB technology can be found at http://developer.nvidia.com/gvdb.

RELEASE NOTES

1/29/2020, GVDB Voxels 1.1.1

• Significantly streamlined CMake build system
• gImportVDB Linux support
• Uses surface and texture objects instead of surface and texture references
• Remove dependency on CUDPP, now using Thrust for radix sorting and reduction
• Watertight voxelization robustness improvements, using Woop, Benthin, and Wald's watertight rasterization technique and D3D10-style top-left edge rules
• Remove limitation of 65,535 bricks in UpdateApron and UpdateApronFaces
• Voxel Size has been removed (i.e. GVDB's world space is now equal to its index space; please use SetTransform to apply your own arbitrary affine transformations)
• Many bug and formatting fixes (including public contributions)

3/25/2018, GVDB Voxels 1.1

• Dynamic topology on GPU
• Multiple contexts and GVDB objects
• Multiple samples with OptiX support
• Improved compute and raytrace performance
• Render from any channel
• Resampling functions
• Grid transforms for rendering
• All limitations in Programming Guide 1.0 addressed

9/19/2017, GVDB Voxels, Incremental fix

• Fixed Depth map integrated raycasting
• Fixed Resample function
• New gDepthMap sample, shows how to render OpenGL and GVDB volumes with proper depth integration.
• New gResample sample, shows how to load a dense RAW file into GVDB as sparse data.
• RAW data courtesy of the Visible Human Project, National Library of Medicine, National Institutes of Health, USA

5/1/2017, GVDB Voxels, Release 1.0 Created by Rama Hoetzlein, 2017

• First public release
• Open source (BSD 3-clause; changed on 2/1/2021 to Apache 2.0)
• Samples: g3DPrint, gFluidSurface, gInteractiveGL, gInteractiveOptix, gJetsonTX, gRenderKernel, gRenderToFile, gSprayDeposit
• Builds on Windows and Linux
• Runs on Quadro, GeForce, JetsonTX1/2, and Tegra/GRID

REQUIREMENTS

• NVIDIA Kepler generation or later GPU
• CUDA Toolkit 6.5 or higher (Toolkit 10.2 recommended)
• CMake 3.10 or later

GVDB is released as a library with samples. The library and each sample can be built separately, using CMake.

WHAT'S IN THE PACKAGE?

• The GVDB Voxels Library
• Code Samples
  • g3DPrint - Demonstrates generating cross-section slices for 3D printing from a polygonal model.
  • gDepthMap - Renders a polygonal mesh and an emissive volume at the same time using OpenGL and CUDA.
  • gFluidSurface - A dynamic fluid simulation with ray traced surface rendering using GVDB and OptiX.
  • gImportVDB - Loads and ray traces a volume stored in OpenVDB's .vdb format.
  • gInteractiveGL - Interactive rendering of a volume using GVDB and OpenGL.
  • gInteractiveOptix - Interactive ray tracing of a volume and a mesh, with light bouncing between the two.
  • gNanoVDB - Exports a GVDB volume to a NanoVDB volume, then renders it using NanoVDB.
  • gPointFusion - Fuses points from a moving camera into a full 3D volume.
  • gRenderKernel - Rendering without OpenGL using a custom GVDB kernel.
  • gRenderToFile - Rendering a semitransparent object to a file without OpenGL using GVDB.
  • gResample - Imaging of a sparse volume generated from dense data.
  • gSprayDeposit - Demonstrates simulated spray deposition onto a 3D park.
  • gJetsonTX - Simple 3D printing driver for the Jetson TX1/2 with volume slicing on a Tegra chip.
• The GVDB VBX File Specification
• GVDB Sample Descriptions
• The GVDB Programming Guide

SAMPLE USAGE

All interactive samples use the following user interface:

• Camera rotation -> move mouse

• Change orientation -> left mouse click

• Zoom -> right mouse click

• Panning -> hold middle button

A few samples have on-screen GUIs with features that can be toggled by clicking on them.

QUICK BUILD PROCESS (Windows and Linux)

Install dependencies

1. Install CMake 3.10 or later.
2. Install CUDA Toolkit 10.2 or later.

To build the samples and the library at the same time:

3. In CMake, configure and generate the build system using gvdb-voxels/CMakeLists.txt, then build the BUILD_ALL target in your IDE (such as Visual Studio).

That's it!

You can also build a specific sample or the GVDB library this way by building its target in this project. Additionally, you can collect a build of GVDB and its samples into a redistributable package by building the INSTALL target.

(Wondering what the GVDB_BUILD_OPTIX_SAMPLES, GVDB_BUILD_OPENVDB, GVDB_BUILD_OPENVDB_SAMPLES, GVDB_BUILD_NANOVDB, and GVDB_BUILD_NANOVDB checkboxes in the CMake GUI do? See "To build the OptiX samples", "To build GVDB with OpenVDB", and "To build the NanoVDB sample" below.)

To build the GVDB Library by itself:

3. In CMake, configure and generate the build system using gvdb-voxels/source/gvdb_library/CMakeLists.txt, then build the gvdb target.

As above, you can create a redistributable build of the NVIDIA GVDB Voxels library by building the INSTALL target.

To build a sample by itself:

3. In CMake, configure and generate the build system using gvdb-voxels/source/[your sample name here]/CMakeLists.txt, then build your sample's target.

That's it! In Visual Studio, you can also run a sample by right-clicking it inside Visual Studio, selecting "Set as StartUp Project", and then pressing F5 or clicking the green triangle in the toolbar.

For some samples, you'll see targets named things like gSample, gSampleApp, and gSamplePTX. In this case, you'll want to build and run gSample; the other two targets compile the application and its PTX files, while the gSample target collects everything together.

Building a sample will also automatically build GVDB, so you no longer need to build and install GVDB before building a sample.

To build the OptiX samples:

1. Install OptiX 6.5 from https://developer.nvidia.com/designworks/optix/download.
2. On Windows, check GVDB_BUILD_OPTIX_SAMPLES in the CMake GUI, or add -DGVDB_BUILD_OPTIX_SAMPLES=ON to the CMake command line.
3. On Linux, check GVDB_BUILD_OPTIX_SAMPLES in the CMake GUI, and also add an entry, OPTIX_ROOT_DIR, pointing to the path to the OptiX SDK (the folder containing OptiX's lib64 directory). Or if you're using the CMake command line, add -DGVDB_BUILD_OPTIX_SAMPLES=ON -DOPTIX_ROOT_DIR=<path to OptiX SDK>, replacing <path to OptiX SDK> with the correct path.
4. Finally, generate and build the CMake project.

To build GVDB with OpenVDB:

Windows:

1. Install OpenVDB: On Windows, one of the easiest ways to install OpenVDB is to use Microsoft's vcpkg; install vcpkg, then run vcpkg install openvdb[tools]:x64-windows. Make sure vcpkg is using the same compiler you'll use to compile GVDB!
2. Work around a temporary issue in vcpkg: If you plan to build GVDB in release mode, go to your vcpkg/installed/x64-windows-debug/bin folder and copy openvdb_d.dll to openvdb.dll. This works around an issue where a build system copies debug-mode openvdb.lib and openvdb.dll to openvdb_d.lib and openvdb_d.dll respectively, but doesn't update the DLL reference in openvdb_d.lib.
3. Configure CMake:
   1. If you're using the CMake GUI, delete the cache, then click the "Configure" button, specify your generator and platform, check "Specify toolchain file for cross-compiling", and click "Next". Then specify the path to vcpkg/scripts/buildsystems/vcpkg.cmake, and click Finish. Then check GVDB_BUILD_OPENVDB (and if you'd like to build the gImportVDB sample as well, check GVDB_BUILD_OPENVDB_SAMPLES) and click "Configure" again.
   2. If you're using the CMake command line, you can also do this by specifying -DCMAKE_TOOLCHAIN_FILE=<path to vcpkg.cmake> -DGVDB_BUILD_OPENVDB=ON -DGVDB_BUILD_OPENVDB_SAMPLES=ON.
   3. Alternatively, if you're not using vcpkg, you can also specify GVDB_OPENVDB_INCLUDE_DIR, GVDB_OPENVDB_LIB_RELEASE_DIR, GVDB_OPENVDB_LIB_DEBUG_DIR, and GVDB_OPENVDB_DLL_RELEASE_DIR and copy in OpenVDB's DLLs using any method - see gvdb_library/CMakeLists.txt for more information.
4. Finally, generate and build the CMake project. Now you can run GVDB with OpenVDB!

Linux:

1. Install OpenVDB 6.1+: On Linux, we recommend building OpenVDB from source using the instructions on OpenVDB's Developer Quick Start, unless OpenVDB 6.1+ is available through your distro's package manager (6.1 introduced a new CMake build system in OpenVDB that we rely upon). Note that you may have to add -DCMAKE_NO_SYSTEM_FROM_IMPORTED:BOOL=TRUE if you run into OpenVDB Issue 144.
2. Configure CMake:
   1. If you're using the CMake GUI, check GVDB_BUILD_OPENVDB and GVDB_BUILD_OPENVDB_SAMPLES and click "Configure" again.
   2. If you're using the CMake command line, you can also do this by specifying -DGVDB_BUILD_OPENVDB=ON -DGVDB_BUILD_OPENVDB_SAMPLES=ON.
3. Finally, generate and build the CMake project. Now you can run GVDB with OpenVDB!

To build the NanoVDB sample:

1. Download NanoVDB from the OpenVDB repository at https://github.com/AcademySoftwareFoundation/openvdb/tree/feature/nanovdb/nanovdb. Since NanoVDB is a header-only library, there's no need to build OpenVDB.
2. In CMake, set GVDB_NANOVDB_INCLUDE_DIR to the path to NanoVDB (this folder contains a nanovdb folder which contains NanoVDB.h), set GVDB_BUILD_NANOVDB to ON, and set GVDB_BUILD_NANOVDB_SAMPLES to ON.
3. Finally, generate and build the CMake project. Now you can run the gNanoVDB sample!

Dependencies

GVDB uses LodePNG and The OpenGL Extension Wranger Library (GLEW). Their licenses can be found in source/sample_utils/lodepng.h and source/gvdb_library/glew/include/GL/glew.h, respectively.

License

==========================

Apache 2.0. Please refer to LICENSE.txt.

Before commit dfcaf6 on February 1st, 2021, GVDB used 3-clause BSD. See #111 for more details.

© 2017-2022 NVIDIA Corporation

gvdb-voxels's Issues

Samples failed to build on Linux and GCC

Several samples (e.g. gPointCloud and gFluidSurface) failed to compile on Linux using GCC (gcc version 5.4.0 20160609 (Ubuntu 5.4.0-6ubuntu1~16.04.10)) because of these errors:

Invalid initialization of non-const reference

error: invalid initialization of non-const reference of type ‘nvdb::DataPtr&’ from an rvalue of type ‘nvdb::DataPtr’

This error might be caused because GCC prohibits passing an rvalue to a non-const reference.

Workaround: in the erroring C++ sample code, find any function call similar to functionName(..., DataPtr(), ...), then replace it with something similar to:

DataPtr dptr1 = DataPtr();
functionName(..., dptr1, ...);

Windows consts and functions not declared

Windows consts and functions (e.g. HANDLE, ReadFile, etc.) are not declared and causes linker errors.

Workaround: comment Windows-specific code when building on Linux, e.g. in source/gPointCloud/main_point_cloud.cpp, comment code in the if (m_io_method == WIN_IO) block.

Undefined reference to checkGL(char*)

Linker error: undefined reference to 'checkGL(char*)' when not running in debug mode. Caused by checkGL(char*) function in source/sample_utils/main_x11.cpp only defined in debug mode.

Workaround: in the source/sample_utils/main_x11.cpp, add the following code after the checkGL(char*) function definition and before the #endif (around line 177):

#else
void checkGL( char* msg ) {}

Win10 x64 compile library error

Hi,
I have an errors when I compile library.

Thanks.

Build Suggestions (Linux)

A few important suggestions on building for linux:

1. It is strongly recommended that you use Cmake to build GVDB. If you do not, you are inviting a lot of pain and trouble. We made the CmakeLists to do the work for you.

2. The suggested CMAKE_INSTALL_PREFIX for gvdb_library is /usr/local/gvdb. You can install to other paths, but there may be side effects we didn't check that you need to work out.

3. Don't forget that you need to do "sudo make install" after you perform "make" on the gvdb_library. This is required in order to place all the necessary files into /usr/local/gvdb.

When done correctly for GVDB 1.1, your /usr/local/gdb should contain the following:
/usr/local/gvdb/include/cuda_gvdb_*.cuh (9 files)
/usr/local/gvdb/include/gvdb_*.h (12 files)
/usr/local/gvdb/include/loader_*.h (3 files)
/usr/local/gvdb/include/radixsort_kernel.cuh and string_helper.h

/usr/local/gvdb/lib/cuda_gvdb_copydata.ptx
/usr/local/gvdb/lib/cuda_gvdb_module.ptx
/usr/local/gvdb/lib/libgvdb.so
/usr/local/gvdb/lib/libcudpp_hash_.so (created when you built cudpp for your platform)
/usr/local/gvdb/lib/libcudpp_.so (created when you built cudpp for your platform)
/usr/local/gvdb/lib/*.glsl (5 files)

4. When building a sample on linux, the CMAKE_INSTALL_PREFIX should be "/home/{user}/build/g3DPrint" or something similar. It should not be /usr/local something, because Cmake may not have sufficient privileges to perform all the necessary steps to that path. This is also true on Windows, you cannot use "C:\Program Files\sample" as the CMAKE_INSTALL_PREFIX due to restricted permissions on this path.

Compile g3DPrint

I'm using, Linux, and I follow your tutorial. I got a problem, for compile all the g3DPrint, gDepthMap... because, when I use make, I got an error that : "Unknown Platform", this error come when we call the platform.h (l. 184) define in sample_utils. I would know, how can I fix it?
So the program recognize that I use GNUC as a compiler but he don't find a "signature" (arm). I put you a copy of my terminal.

[ 14%] Building CXX object CMakeFiles/gDepthMap.dir/main_depthmap.cpp.o
In file included from /usr/local/gvdb/source/sample_utils/main.h:27:0,
from /usr/local/gvdb/source/gDepthMap/main_depthmap.cpp:9:
/usr/local/gvdb/source/sample_utils/platform.h:184:6: error: #error "Unknown platform"

error "Unknown platform"

Thanks for your answer!
Colin

memory leak

I try to contruct and destroy a nvdb::VolumeGVDB multiple times. I monitor a memory leak. This simple code illustrates the leak.
while (1) { nvdb::VolumeGVDB gvdb; gvdb.SetCudaDevice(GVDB_DEV_FIRST); gvdb.Initialize(); }
Is there a function to free the VolumeGVDB or is this pattern not permitted in gvdb ?

Windows build issue: CUDPP

First of all thanks for updating gvdb, really appreciate the support for multiple contexts!
Just wanted to point out a small issue during the build process. The readme states:

"Install CUDA Toolkit 8.0"

However, all the provided CUDPP binaries are compiled for CUDA 9 and are not compatible during a CUDA 8 build.

It was a simple matter of compiling CUDPP again for CUDA 8 and renaming the binaries using the convention:
cudpp_1900x64.lib -> cudpp_1900cu8x64.lib
etc.

wrong topology computation in example g3DPrint

I'm trying to use this project to implement some material simulations. But before to go ahead, I'd like to know if the output of the g3Dprint example is correct. Assuming that It is defined a GVDB configuration such that <3,3,3,3,5>, the number of visualized level should be 5, am I right?

If I'm right, it's means that some thing is not working on. So any clue to resolve this?

Samples running trouble

I have install the gvdb library but I can not make the samples to run. I can compile the samples .cpp and when the programs are running a CUDA error appear "CUDA ERROR: file not found (func: cuda_gvdb_module.ptx, caller: ModuleLoad)".

My guess is that the gvdb library trying to find the file but is not in the supposed direction. Since the installation is made on the system the address of the file is not properly defined. I can copy the file to the corresponding address, but I don't know where to put it the system, in order for the gvdb should look for.

class Allocator without GVDB_API

GVDB_API dll macro is not added to class Allocator. This cause linker errors, e.g. using gvdb.mPool->getAtlasRes(chan) in the user code.

FindOptix.cmake in example fails with Windows10, CMake 3.7.2

CMake fails on _FIND_FILE macro.

As a remedy, I replaced with the FindOptix.cmake in the Optix SDK and CMake correctly works.

SolidVoxelize violation access with GLFW

I'm trying to use GLFW in place of NVPWindow because i want to run a voxelization in background without a window. I've made a naive initialization with a GLFWwindow, but i experienced a violation access at SolidVoxelize(). Do i've missed something ?

VolumeGVDB gvdb;
glewInit();
glfwInit();

glfwWindowHint(GLFW_CONTEXT_VERSION_MAJOR, 3);
glfwWindowHint(GLFW_CONTEXT_VERSION_MINOR, 2);
glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE);
glfwWindowHint(GLFW_OPENGL_FORWARD_COMPAT, GL_TRUE);
glfwWindowHint(GLFW_VISIBLE, GLFW_FALSE);

GLFWwindow* window = glfwCreateWindow(800, 600, "OpenGL", nullptr, nullptr); // Windowed
glfwMakeContextCurrent(window);

// Initialize GVDB
gvdb.SetDebug(true);
gvdb.SetVerbose(true);
gvdb.SetProfile(false, false);
gvdb.SetCudaDevice(GVDB_DEV_FIRST);
gvdb.Initialize();
gvdb.AddPath("../../source/shared_assets/");
gvdb.AddPath("../shared_assets/");
gvdb.AddPath(ASSET_PATH);
gvdb.StartRasterGL();

gvdb.getScene()->SetSteps(0.25f, 16.f, 0.25f); // Set raycasting steps
gvdb.getScene()->SetVolumeRange(0.25f, 0.0f, 1.0f); // Set volume value range
gvdb.getScene()->SetExtinct(-1.0f, 1.1f, 0.f); // Set volume extinction
gvdb.getScene()->SetCutoff(0.005f, 0.005f, 0.f);
gvdb.getScene()->SetShadowParams(0, 0, 0);

int model_id = gvdb.getScene()->AddModel("lucy.obj", 1.0, 0, 0, 0);
gvdb.CommitGeometry(model_id);

Model* m = gvdb.getScene()->getModel(model_id);
Matrix4F xform;
Vector3DF voxelsize(0.2f, 0.2f, 0.2f); // Voxel size (mm)
gvdb.SetVoxelSize(voxelsize.x, voxelsize.y, voxelsize.z);
gvdb.SolidVoxelize(0, m, &xform, 1.0, 0.5);

no kernel image is available for execution on the device

Does any one know what the issue is. I'm running the pointCloud example
Cuda 9.2
Optix 5.11
laptop gpu GTX680m
VS2017

gFluidSurface fails on GTX 970 on Win10..

testing new SDK 1.1 reléase on Windows 10 GTX 970 and 391.35 driver..
all samples work ok except: gFluidSurface.exe
it's due to 4gb GPU memory not enough?
thanks..
error log:

FLUID SYSTEM, CUDA ERROR:
Launch status: unspecified launch failure
Kernel status: no error
Caller: FluidSystem::TransferFromCUDA
Call: cuMemcpyDtoH
Args: FPOS
Error. Application will exit.

full log below:

gFluidSurface.exe
Reading TGA: arial.tga
Creating OptiX context.
Setting Ray Generation program.
Loading: optix_trace_primary.ptx, trace_primary
Loading: optix_trace_primary.ptx, exception
Setting Miss program.
Loading: optix_trace_miss.ptx, miss
Creating random number buffer.
Creating Mesh intersection programs.
Loading: optix_mesh_intersect.ptx, mesh_intersect
Loading: optix_mesh_intersect.ptx, mesh_bounds
Creating Volume intersection programs.
Loading: optix_vol_intersect.ptx, vol_intersect
Loading: optix_vol_intersect.ptx, vol_levelset
Loading: optix_vol_intersect.ptx, vol_deep
Loading: optix_vol_intersect.ptx, vol_bounds
Creating OptiX Main Group and BVH.
PERF_INIT: No CPU markers.
PERF_INIT: Disabling GPU markers. Did not find nvToolsExt64_1.dll.
PERF_INIT: Disabling perf. No CPU or GPU markers.
Starting GVDB Voxels. ver 1.11
Creating Allocator..
Creating Scene..
Creating CUDPP..
Starting Fluid System.
Add Particles. Density: 400.000000, Spacing: 1.393426, PDist: 0.008008
CUDA Config:
Pnts: 1500000, t:2930x512=1500160, Size:1500160
Grid: 1203052, t:2350x512=1203200, bufGrid:0, Res: 134x67x134
Creating screen buffer. 1024 x 768
Loading: optix_trace_surface.ptx, trace_surface
Loading: optix_trace_surface.ptx, trace_shadow
Adding GVDB Volume to OptiX graph.
Validating OptiX.
Update GVDB Volume.
Rebuild Optix.. Done.
Running..
OptiX buffer size: 1024, 768
FLUID SYSTEM, CUDA ERROR:
Launch status: unspecified launch failure
Kernel status: no error
Caller: FluidSystem::TransferFromCUDA
Call: cuMemcpyDtoH
Args: FPOS
Error. Application will exit.

build error in interactive_optix

VS2017
GVDB 1.1
CUDA 10
optix_extra_math.cuh(34): fatal error C1083: Cannot open include file: 'optix.h': No such file or directory

I doubled checked in visual studio optix.h was added to the includes.

All other optix related examples are all working.

GVDB-Vulkan Interop, and other feature?

I understand that GVDB is indeed a render. Either Vulkan or OpenGL is merely a image shower. But it still would be nice if you guys can update a vulkan interop for the fans. Yep, I'm a big fan. Thank you guys for the great work.

It looks like it has been a while since last update under this repo. Do owners intent to maintain this repo in the future?

potential build dependency issue

Hi! Thanks for the great talk at GTC Rama!

Just attempted building, did following:

git clone https://github.com/NVIDIA/gvdb-voxels.git
cd gvdb-voxels
mkdir build; cd build;

Tried building gvdb library first

cmake ../source/gvdb_library/

Got:

....
CMake Error: The following variables are used in this project, but they are set to NOTFOUND.
Please set them or make sure they are set and tested correctly in the CMake files:
X11_Xcursor_LIB (ADVANCED)
linked by target "gvdb" in directory /home/jimmy/dev/gvdb-voxels/source/gvdb_library
X11_Xi_LIB (ADVANCED)
linked by target "gvdb" in directory /home/jimmy/dev/gvdb-voxels/source/gvdb_library
X11_Xinerama_LIB (ADVANCED)
linked by target "gvdb" in directory /home/jimmy/dev/gvdb-voxels/source/gvdb_library
X11_Xrandr_LIB (ADVANCED)
linked by target "gvdb" in directory /home/jimmy/dev/gvdb-voxels/source/gvdb_library
....

So possibly I am missing some X11 dependency?

THanks!

Does GVDB need more customization?

Hi,

I am working on a simulation and we started using GVDB to handle self collisions.
So far, we have got it working nicely, but I needed to make a bunch of modifications in the source and would like to know if they were really necessary and/or are planned to be included in future releases.

As far as I am aware, in the InsertPointsSubcell, launching cuFunc[FUNC_INSERT_SUBCELL] copies the point positions in the slot of each sub-cell overlapping the AABB of the points. Optionally, it copies velocities and colors.

For our application, we needed different info. We could have reused the color attribute, since we are not using it anyway, but the algorithm is not finished yet and we will probably need to add more attributes...
Since I couldn't find any way to launch a custom "myInsertPointsSubcellKernel", we recompiled the SDK with a InsertPointsSubcell with an extra argument of type CUfunction.

In the "myInsertPointsSubcellKernel" instead of positions, we store just the point ID. Further, we are also using a custom "gather" function where we read that ID and actually compute the needed info into custom channels.

Since we will be using GVDB with very large point clouds, copying only the ID consumes less memory and gives us more freedom to compute custom channels.

So, it was not a big change in the core, just extra methods.
A "InsertPointsSubcell" and a "gather" that accept a user CUfunction.

thank you.

LoadVBX bug (not a SaveVBX bug)

In SaveVBX function, at line 1814
if (major >= 2) {
fwrite( &use_bitmask, sizeof(uchar), 1, fp ); // bitmask usage (GVDB 2.0 or higher)
}

It 's not writed bitmask sign, but in LoadVBX, read this sign.

Modified:
if ((major == 1 && minor >= 1) || major > 1) {
fwrite( &use_bitmask, sizeof(uchar), 1, fp ); // bitmask usage (GVDB 2.0 or higher)
}

Missing CUDA_TOOLKIT_VERSION creates libcudpp_.so

I'm building CUDPP on CentOS 7 with CUDA installed in a non-standard location (but findable by cmake 3.9.6). My commands are:

cd gvdb-voxels
mkdir build_cudpp
cd build_cudpp
cmake -DCUDA_SDK_ROOT_DIR=/mnt/modules/cuda/cuda-9.0/samples -DCMAKE_BUILD_TYPE=Release ../source/shared_cudpp/
make

This builds the libs just fine, but the names of the .so files are wrong:

gvdb-voxels/build_cudpp $ ls -l lib/
total 35568
-rw-rw-r-- 1 mstock n-apps      559 Jun 12 10:19 cudpp-config.cmake
-rw-rw-r-- 1 mstock n-apps      396 Jun 12 10:19 cudpp-config.cmake.install
-rw-rw-r-- 1 mstock n-apps      295 Jun 12 10:19 cudpp-config-version.cmake
-rw-rw-r-- 1 mstock n-apps     2843 Jun 12 10:19 cudpp-targets.cmake
-rwxrwxr-x 1 mstock n-apps 35582064 Jun 12 10:22 libcudpp_.so
-rwxrwxr-x 1 mstock n-apps   815232 Jun 12 10:22 libcudpp_hash_.so

Naturally, this causes the gvdb cmake command to fail.
I can try to track down why cmake was able to find CUDA but didn't set CUDA_TOOLKIT_VERSION, but I think that source/shared_cudpp/CMakeLists.txt should throw an error if that variable is not defined.

Compute total volume (reduction): Difference between atlas_cnt and atlas_res, layout of bricks in atlas

I've created a simple kernel function to walk over the bricks in the atlas space with the objective of get statistics like mean, min and max value of each bricks. The kernel is launch by ComputeKernel method. This kernel bounds CUDA threads using atlas_cnt attribute of VDBInfo. To get node information of the current brick: first I generate bitmask_pos by getBitPos; and then, I get the node according las bitmask_pos and lev = 0. Using three concatenated loops I visit all voxels in the bricks (from atlas_apron to brick_res - 2*altas_apron). Each voxel access to the texture by tex(volIn, p.x + o.x, p.y+o.y, p.z+o.z) where o = node->mValue.

I'm not totally sure that this is working properly because I'm not sure about how are stored bricks in the atlas. I mean if all bricks in the atlas are contiguous then this strategy is fine because the kernel is bounded to the altas_cnt value which is the number of bricks stored (or what is the same, it's pointing to the last birck stored). Otherwise I should use atlas_res which actually is the resolution of the atlas (but is this in the voxel level or brick level?)

I was looking for an answer in the programing guide, but there is not a clear one.

Strange behaviour of SetModule(CUModule)

I have a custom module which use macros like SCN_WIDTH and VBOX so to be able to use this information in my kernels I set up the new module. The problem come here, the processing of data in these custom kernel is good but now the visualizer provided with GVDB starts to do strange thinks, like the atlas information is not transforming according with the camera transformation. I've attached an image to better understanding. As can you see, the topology has change according with the camera transformation, but the data which came from the atlas is frozen.

I guess that the problem is because when you set a new module the scene information is just being updating in the custom module but not in the default one. Am I right? Or any other suggestion?

StartRasterGL() Error

Hello, I am following the initialization process of using OpenGL

I am using SDL2 library and created a OpenGL context using SDL_GL_CreateContext() then called glewInit()

However, afterwards when I called StartRasterGL(), I got access violation

Exception thrown at 0x0000000000000000 in sim.exe: 0xC0000005: Access violation executing location 0x0000000000000000.

suggesting I am de-referencing an uninitialized pointer . I guess it is the OpenGL context pointer that VolumeGVDB is missing. How could I link them?

My code

Window window;
Window::Window(const unsigned int width, const unsigned int height, const std::string & title)
{
	_width = width;
	_height = height;
	_title = title;
	SDL_Init(SDL_INIT_EVERYTHING);	// Initilize everything using SDL

	SDL_GL_SetAttribute(SDL_GL_RED_SIZE, 8);
	SDL_GL_SetAttribute(SDL_GL_GREEN_SIZE, 8);
	SDL_GL_SetAttribute(SDL_GL_BLUE_SIZE, 8);
	SDL_GL_SetAttribute(SDL_GL_ALPHA_SIZE, 8);
	SDL_GL_SetAttribute(SDL_GL_BUFFER_SIZE, 32);
	SDL_GL_SetAttribute(SDL_GL_DEPTH_SIZE, 16);
	SDL_GL_SetAttribute(SDL_GL_DOUBLEBUFFER, 1);	// Set attributes for SDL window

	_window = SDL_CreateWindow(title.c_str(), SDL_WINDOWPOS_CENTERED, SDL_WINDOWPOS_CENTERED, width, height, SDL_WINDOW_OPENGL);	// Create OpenGL window using SDL
	_gl_context = SDL_GL_CreateContext(_window);	// Create OpenGL context

	GLenum res = glewInit();

	if (res != GLEW_OK)
	{
		std::cerr << "Error:" << glewGetErrorString(res) << std::endl;	// Monitor if glew initilization succeed
	}

	glEnable(GL_DEPTH_TEST);	// Enable GL depth test
	glEnable(GL_CULL_FACE);		// Enable GL cull face
	glCullFace(GL_BACK);		// Only one face will be rendered
	glEnable(GL_PROGRAM_POINT_SIZE);
	glEnable(GL_POINT_SPRITE);
}

then

Simulation sim;
sim.Init();
	void Simulation::Init()
{
	Scene* scn = _gvdb.getScene();
	// Initialize GVDB
	_gvdb.SetDebug(true);
	_gvdb.SetVerbose(true);		// enable/disable console output from gvdb
	_gvdb.SetProfile(true, true);
	_gvdb.SetCudaDevice(GVDB_DEV_FIRST);
	_gvdb.Initialize();
	_gvdb.StartRasterGL();
	_gvdb.AddPath("./Assets/");
	_gvdb.Configure(3, 3, 3, 3, 5); // May use overload function to define N-level tree
	_gvdb.AddChannel(0, T_FLOAT, 1);
	scn->AddModel("torus.obj", 1.0, 0.0, 0.0, 0.0);
	_gvdb.CommitGeometry(0);

}

Voxel index restriction due to CUDPP Radix Sort

gvdb-voxels/source/gvdb_library/kernels/cuda_gvdb_particles.cuh

Lines 366 to 368 in 96781dd

	extern "C" __global__ void gvdbCalcBrickId ( int num_pnts, int lev_depth, int* range_res,
	char* ppos, int pos_off, int pos_stride,
	float3 orig, unsigned short* pout)

The function argument for gvdbCalcBrickId requires the level-index list to be unsigned short, which I believe is due to the CUDPP radix sort limit of taking CUDPP_LONGLONG

gvdb-voxels/source/gvdb_library/src/gvdb_volume_gvdb.cpp

Lines 2559 to 2564 in 96781dd

	CUDPPConfiguration config_sort;
	config_sort.algorithm = CUDPP_SORT_RADIX;
	config_sort.datatype = CUDPP_LONGLONG;
	config_sort.options = CUDPP_OPTION_KEYS_ONLY;
	result = cudppPlan(mCudpp, &mPlan_sort, config_sort, maxNum, 1, 0);

This limits the maximum voxel index to 0~2^(16+brick_log_res). Given brick_log_res is recommended to be 5, the final limit is 0~2^21, or -2^20~2^20. Could there be any other places that a tighter restriction exists?

Color Expansion Kernel

I've rewritten the color expansion kernel to use the accumulated alpha channel as a weight, instead of dividing by the highest RGB channel, which I don't understand the motivation for. I believe this is more suitable in most circumstances, assuming the alpha channel is used.

inline __host__ __device__ int4 make_int4(uchar4 c) {
	return make_int4(c.x, c.y, c.z, c.w);
}

extern "C" __global__ void gvdbOpClrExpand ( int3 res, uchar chan, float p1, float p2, float p3 )
{
	GVDB_COPY_SMEM_UC4

	int4 cs;
	int4 c;
	c = make_int4( svox[ndx.x][ndx.y][ndx.z] );   cs = c*p1;   
	c = make_int4( svox[ndx.x-1][ndx.y][ndx.z] ); cs += c*p2;
	c = make_int4( svox[ndx.x+1][ndx.y][ndx.z] ); cs += c*p2;
	c = make_int4( svox[ndx.x][ndx.y-1][ndx.z] ); cs += c*p2;
	c = make_int4( svox[ndx.x][ndx.y+1][ndx.z] ); cs += c*p2;
	c = make_int4( svox[ndx.x][ndx.y][ndx.z-1] ); cs += c*p2;
	c = make_int4( svox[ndx.x][ndx.y][ndx.z+1] ); cs += c*p2;
	
	cs = (cs.w > 255) ? make_int4(cs.x * 255 / cs.w, cs.y * 255 / cs.w, cs.z * 255 / cs.w, 255) : cs; 
	
	surf3Dwrite ( make_uchar4(cs.x, cs.y, cs.z, cs.w), volOut[chan], vox.x*sizeof(uchar4), vox.y, vox.z );
}

OptiX CMake version error

The current CMake file for OptiX related sample can only search OptiX version up to 4.0.0. If you are using OptiX 5.0.0 and above, you need to manully select OPTIX_BIN_DIR, OPTIX_INCLUDE_DIR, OPTIX_LIB_DIR and OPTIX_ROOT_DIR

4D visulization

Hello,
I'm evaluating the library to see if it fits my needs.
I just want to display medical volume data, say 512512512*sizeof(float) = 536MB of float data per frame. I need to display at least at 20fps.
The entire volume cannot fit inside the GPU.
Do you think I can copy fast enough the data from CPU to GPU ?

Can I use the concurrent copy engine (streams) with GVDB ?

Do you have advices on the best strategies ?

thank you

André

nvcc fatal : Stray '"' character

On arch linux 4.11.3-1 using cmake 3.8.2, I get the following when following the linux build instructions:

CMake Warning (dev) at cuda_compile_ptx_1_generated_cuda_gvdb_copydata.cu.ptx.Release.cmake:77:
Syntax Warning in cmake code at column 56

Argument not separated from preceding token by whitespace.
This warning is for project developers. Use -Wno-dev to suppress it.

nvcc fatal : Stray '"' character
CMake Error at cuda_compile_ptx_1_generated_cuda_gvdb_copydata.cu.ptx.Release.cmake:214 (message):
Error generating
/tmp/gvdb-voxels-master/build/gvdb_library/cuda_compile_ptx_1_generated_cuda_gvdb_copydata.cu.ptx

make[2]: *** [CMakeFiles/gvdb.dir/build.make:65: cuda_compile_ptx_1_generated_cuda_gvdb_copydata.cu.ptx] Error 1
make[1]: *** [CMakeFiles/Makefile2:67: CMakeFiles/gvdb.dir/all] Error 2
make: *** [Makefile:130: all] Error 2

All precompiled samples fail to run on Titan V (Volta) due to CUDPP/PTXAS memory leak(?)

Hi,
just changed a GTX 970 by a Titan V..
I was expecting all to work ok similar to GTX 970 and even perhaps be able to run fluid surface sample
(#15)..
what I see is no matter if on Windows or Linux using latest drivers (using 398.58 currently but others seem to fail too)..
output is:

gInteractiveGL.exe
  Device List:
   0. TITAN V
   Using Device: 0, TITAN V, Context: 000001FDCE9C1D00
Starting GVDB Voxels. ver 1.11
 Creating Allocator..
 Creating Scene..
 Creating CUDPP..
ptxas fatal   : Memory allocation failure

before "ptxas fatal : Memory allocation failure" I see on process manager process gets over 40GB memory before failing as I have 32GB RAM and virtual memory up to 16GB more..
don't know if compiling CUDPP by myself with CUDA 9.2 SDK and explicit sm_70,compute_70 flags would solve the problem..
hope GVDB gets some Volta fixes soon..
thanks..

Why isn't the 3rd parameter of gvdbBrickFunc2_t passed by reference?

Hi all,

I've been working in a modification of the raysurfaceTrilinearBrick function,

gvdb-voxels/source/gvdb_library/kernels/cuda_gvdb_raycast.cuh

Line 226 in 25ed5f4

__device__ void raySurfaceVoxelBrick ( VDBInfo* gvdb, uchar chan, int nodeid, float3 t, float3 pos, float3 dir, float3& pStep, float3& hit, float3& norm, float4& hclr )

used by the raycast function

gvdb-voxels/source/gvdb_library/kernels/cuda_gvdb_raycast.cuh

Line 504 in 25ed5f4

__device__ void rayCast ( VDBInfo* gvdb, uchar chan, float3 pos, float3 dir, float3& hit, float3& norm, float4& clr, gvdbBrickFunc_t brickFunc )

to render volumes. If the variable t which is used to go forward over the ray in the 3DDDA algorithm, my question is, why isn't the 3rd parameter (float3 t) of gvdbBrickFunc2_t passed by reference when the algorithm is in level 1 of the tree (that is, processing brick level)? Changing this will make much more coherent the implementation regards with the algorithm since t once the method had explored the brick content, t will be pointing to the correct place, otherwise t keeps pointing to the last value in level 2.

The interface definition of these methods is write there,

gvdb-voxels/source/gvdb_library/kernels/cuda_gvdb_raycast.cuh

Line 34 in 25ed5f4

typedef void(*gvdbBrickFunc_t)( VDBInfo*, uchar, int, float3, float3, float3, float3&, float3&, float3&, float4& );

Incorrect depth buffer check in depth buffer test

The current depth buffer test compares the linear depth for a particular pixel with the raybox intersection distance along the ray.

This is incorrect, and should be comparing the linear depth against the distance along the ray projected to the camera z axis (assuming that scene depth buffer is equivalent to a depth image transformed by the projection matrix to normalized 0..1 values)

i.e., instead of

if (t.x > getLinearDepth(SCN_DBUF) )

it should be something along the lines of

if (t.x * dot(scn.dir_vec, dir) > getLinearDepth(SCN_DBUF) )

where for convenience I have exposed the camera z in world frame as scn.dir_vec
Of possible relevance to @drmateo and #44

Unable to run samples - cudaMalloc fills all the system memory

I'm building on Windows 10 (tried on Linux Ubuntu also, same problem).

When I run a sample it blocks in function allocReduceStorage :
gvdb_1.1\shared_cudpp\src\cudpp\app\reduce_app.cu

[...]
void allocReduceStorage(CUDPPReducePlan *plan)
...
    case CUDPP_FLOAT:
-->        cudaMalloc(&plan->m_blockSums, blocks * sizeof(float));
        break;

System :
Version 10.0.17134 Build 17134
Processor i7-8750H
Graphics card: NVIDIA GeForce GTX 1070 with Max-Q Design
RAM: 32GB
GVDB: 1.1
CUDA: 10

---

`

•   plan	0x000001f057310b30 {m_threadsPerBlock=256 m_maxBlocks=64 m_blockSums=0xcdcdcdcdcdcdcdcd }	CUDPPReducePlan *
  

•   CUDPPPlan	{m_config={algorithm=CUDPP_REDUCE (3) op=CUDPP_MAX (3) datatype=CUDPP_FLOAT (6) ...} m_numElements=1000000 ...}	CUDPPPlan
  

•   __vfptr	0x00007ffb96c87928 {cudpp_1915x64d.dll!void(* CUDPPReducePlan::`vftable'[2])()} {0x00007ffb96417662 {cudpp_1915x64d.dll!CUDPPReducePlan::`vector deleting destructor'(unsigned int)}}	void * *
  

•   m_config	{algorithm=CUDPP_REDUCE (3) op=CUDPP_MAX (3) datatype=CUDPP_FLOAT (6) ...}	CUDPPConfiguration
    m_numElements	1000000	unsigned __int64
    m_numRows	1	unsigned __int64
    m_rowPitch	0	unsigned __int64
  

•   m_planManager	0x000001f0571fe540 {m_deviceProps={name=0x000001f0571fe540 "GeForce GTX 1070 with Max-Q Design" uuid=...} }	CUDPPManager *
    m_threadsPerBlock	256	unsigned int
    m_maxBlocks	64	unsigned int
    m_blockSums	0xcdcdcdcdcdcdcdcd	void *
  

•   (*((CUDPPPlan*)plan)).m_planManager->m_deviceProps	{name=0x000001f0571fe540 "GeForce GTX 1070 with Max-Q Design" uuid={bytes=0x000001f0571fe640 "ÀæÙþ?ÆÐPÀÇ_¢ñ9¬3... } ...}	cudaDeviceProp
  

•   name	0x000001f0571fe540 "GeForce GTX 1070 with Max-Q Design"	char[256]
  

•   uuid	{bytes=0x000001f0571fe640 "ÀæÙþ?ÆÐPÀÇ_¢ñ9¬3... }	CUuuid_st
  

•   luid	0x000001f0571fe650 "�û"	char[8]
    luidDeviceNodeMask	1	unsigned int
    totalGlobalMem	8589934592	unsigned __int64
    sharedMemPerBlock	49152	unsigned __int64
    regsPerBlock	65536	int
    warpSize	32	int
    memPitch	2147483647	unsigned __int64
    maxThreadsPerBlock	1024	int
  

•   maxThreadsDim	0x000001f0571fe684 {1024, 1024, 64}	int[3]
  

•   maxGridSize	0x000001f0571fe690 {2147483647, 65535, 65535}	int[3]
    clockRate	1265500	int
    totalConstMem	65536	unsigned __int64
    major	6	int
    minor	1	int
    textureAlignment	512	unsigned __int64
    texturePitchAlignment	32	unsigned __int64
    deviceOverlap	1	int
    multiProcessorCount	16	int
    kernelExecTimeoutEnabled	1	int
    integrated	0	int
    canMapHostMemory	1	int
    computeMode	0	int
    maxTexture1D	131072	int
    maxTexture1DMipmap	16384	int
    maxTexture1DLinear	134217728	int
  

•   maxTexture2D	0x000001f0571fe6e4 {131072, 65536}	int[2]
  

•   maxTexture2DMipmap	0x000001f0571fe6ec {32768, 32768}	int[2]
  

•   maxTexture2DLinear	0x000001f0571fe6f4 {131072, 65000, 2097120}	int[3]
  

•   maxTexture2DGather	0x000001f0571fe700 {32768, 32768}	int[2]
  

•   maxTexture3D	0x000001f0571fe708 {16384, 16384, 16384}	int[3]
  

•   maxTexture3DAlt	0x000001f0571fe714 {8192, 8192, 32768}	int[3]
    maxTextureCubemap	32768	int
  

•   maxTexture1DLayered	0x000001f0571fe724 {32768, 2048}	int[2]
  

•   maxTexture2DLayered	0x000001f0571fe72c {32768, 32768, 2048}	int[3]
  

•   maxTextureCubemapLayered	0x000001f0571fe738 {32768, 2046}	int[2]
    maxSurface1D	32768	int
  

•   maxSurface2D	0x000001f0571fe744 {131072, 65536}	int[2]
  

•   maxSurface3D	0x000001f0571fe74c {16384, 16384, 16384}	int[3]
  

•   maxSurface1DLayered	0x000001f0571fe758 {32768, 2048}	int[2]
  

•   maxSurface2DLayered	0x000001f0571fe760 {32768, 32768, 2048}	int[3]
    maxSurfaceCubemap	32768	int
  

•   maxSurfaceCubemapLayered	0x000001f0571fe770 {32768, 2046}	int[2]
    surfaceAlignment	512	unsigned __int64
    concurrentKernels	1	int
    ECCEnabled	0	int
    pciBusID	1	int
    pciDeviceID	0	int
    pciDomainID	0	int
    tccDriver	0	int
    asyncEngineCount	5	int
    unifiedAddressing	1	int
    memoryClockRate	4004000	int
    memoryBusWidth	256	int
    l2CacheSize	2097152	int
    maxThreadsPerMultiProcessor	2048	int
    streamPrioritiesSupported	1	int
    globalL1CacheSupported	1	int
    localL1CacheSupported	1	int
    sharedMemPerMultiprocessor	98304	unsigned __int64
    regsPerMultiprocessor	65536	int
    managedMemory	1	int
    isMultiGpuBoard	0	int
    multiGpuBoardGroupID	0	int
    hostNativeAtomicSupported	0	int
    singleToDoublePrecisionPerfRatio	32	int
    pageableMemoryAccess	0	int
    concurrentManagedAccess	0	int
    computePreemptionSupported	0	int
    canUseHostPointerForRegisteredMem	0	int
    cooperativeLaunch	0	int
    cooperativeMultiDeviceLaunch	0	int
    sharedMemPerBlockOptin	0	unsigned __int64
    pageableMemoryAccessUsesHostPageTables	0	int
    directManagedMemAccessFromHost	0	int
  

`

Linux compilation error

Hi,
I have an error when I compile examples on my ubuntu computer (windows works fine), the problem comes with the m_screenquad_utex.

main_x11.cpp:825:2: error: ‘m_screenquad_utex’ was not declared in this scope
m_screenquad_utex = glGetUniformLocation ( m_screenquad_prog, "uTex" );

And error in functions declaration : renderScreenQuadGL

error: prototype for ‘void NVPWindow::renderScreenQuadGL(int, float, float, float, float, char, char)’ does not match any in class ‘NVPWindow’

Thanks,

Add export STL option for 3D printing

I know that it wasn't intended to be used this way, but it would be a great if we could

export 3D mesh out of GVDB voxels.

Gaps in voxels:

Hi , I am wondering if this is only a bug on my hardware: GTX 1060m, win10 x64 -- I am getting gaps in the g3DPrint demo :

gInteractiveOptix won't build in Release 1.1

The gInteractiveOptix sample program in Release 1.1 seems to have multiple problems preventing it from building (on Ubuntu Linux, but the problems don't seem to be platform specific).

First, even when OPTIX_ROOT_DIR is specified in the CMake configuration step, the Optix include directory is not passed to nvcc, causing this error:

[ 4%] Building NVCC ptx file cuda_compile_ptx_generated_optix_vol_intersect.cu.ptx
In file included from <omitted>/gvdb-voxels-1.1/source/gInteractiveOptix/optix_vol_intersect.cu:30:0:
<omitted>/gvdb-voxels-1.1/source/gInteractiveOptix/optix_extra_math.cuh:34:19: fatal error: optix.h: No such file or directory

A work-around (probably not the best) is to configure CMake to explicitly set CUDA_NVCC_FLAGS:
-I<Optix include path>

With that work-around, additional compilation errors appear. One is the following:

<omitted>/gvdb-voxels-1.1/source/gInteractiveOptix/optix_vol_intersect.cu(78): error: identifier "raySurfaceBrick" is undefined

In Release 1.0, the include/cuda_gvdb_raycast.cuh file contained a raySurfaceBrick() function that seems to have been removed for Release 1.1.

There are enough other differences between the Release 1.0 and Release 1.1 versions of cuda_gvdb_raycast.cuh that I have not tried resurrecting the 1.0 version of raySurfaceBrick().

Will there be an update to fix these problems? Thanks.

Radius search in point cloud

Hello, I'm here for some helps.
Currently I'm implementing radius search(Specifically, I want to find a set of indices of registered point around the query point within the given radius.) for a given point using GVDB.
I'm implementing this because I expected that GVDB has this kind of function but I couldn't find any.
I've read quite some examples(particle surface, point cloud, point fusion) and guide documents but couldn't find the functionality that I want.

Am I missing somethings or GVDB does not have the kNN search function?

Thank you.

gInteractiveOptix

In cuda_gvdb_raycast.cuh the function getLinearDepth uses the scn.* data structure only found under #define CUDA_PATHWAY without an ifdef guarding it. This means that OPTIX_PATHWAY will not compile if you include cuda_gvdb_raycast.cuh which the gInteractiveOptix sample does.

cmake on Linux generates malformed nvcc commands

I am trying to build gvdb on CentOS with gcc 6.3.0, CUDA 9.0, and cmake 3.9.6. After working around issues building CUDPP, I ran into a problem with the resulting Makefiles from cmake for gvdb. I ran:

cd gvdb-voxels
mkdir build_gvdb
cd build_gvdb
cmake -DCUDA_SDK_ROOT_DIR=/mnt/modules/cuda/cuda-9.0/samples -DCMAKE_BUILD_TYPE=Release -DCUDPP_ROOT_DIR=/home/me/opt/gvdb-voxels/build_cudpp ../source/gvdb_library/
VERBOSE=1 make

And the very first compilation command emits:

...
[  5%] Building NVCC ptx file cuda_compile_ptx_1_generated_cuda_gvdb_module.cu.ptx
...
/home/me/opt/gvdb-voxels/build_gvdb/CMakeFiles/cuda_compile_ptx_1.dir/kernels/cuda_compile_ptx_1_generated_cuda_gvdb_module.cu.ptx.NVCC-depend -ccbin /mnt/modules/gcc/gcc-6.3.0/bin/gcc -m64 -DRESOURCE_DIRECTORY= -D/home/me/opt/gvdb-voxels/source/gvdb_library/resources/\" -DLINUX -DGLEW_STATIC -DGLEW_NO_GLU -DPROJECT_RELDIRECTORY=\"../../../gvdb_library/\" -DPROJECT_ABSDIRECTORY=\"/home/me/opt/gvdb-voxels/source/gvdb_library/\" -DPROJECT_NAME=\"gvdb\" -DUSECUDA -DUSE_CUDPP\" -DNVCC -I/mnt/modules/cuda/cuda-9.0/include -I/home/me/opt/gvdb-voxels/source/gvdb_library/glew/include -I/home/me/opt/gvdb-voxels/source/gvdb_library/src -I/home/me/opt/gvdb-voxels/build_cudpp/include
nvcc fatal   : Stray '"' character

which contains some weirdly-formed macros.

calling gvdb.SetCudaDevice error

calling gvdb.SetCudaDevice error in main_point_cloud.cpp line 463

After enable debug, it shows LoadFunction error.

cudaCheck ( cuModuleLoad ( &cuModule[mid], cptx ), "VolumeGVDB", "LoadFunction", "cuModuleLoad", cptx, mbDebug);

SolidVoxelize takes T_FLOAT channel instead of T_UCHAR

gvdb-voxels/source/gvdb_library/src/gvdb_volume_gvdb.h

Line 483 in e9a2350

void SolidVoxelize ( uchar chan, Model* model, Matrix4F* xform, float val_surf, float val_inside, float vthresh=0.0 );

SolidVoxelize actually takes in a T_FLOAT channel instead of T_UCHAR channel. And it seems this function does not have the ability to assign 'S' for surface and 'I' for interior. Also, there is not function called SurfaceVoxelize

point cloud processing related question

What are the steps of imputing a point cloud and converting it to mesh. Is there a inverse of polygon models for gvdb?

timer for volume rendering incorrect in sample code

in the sample code:

gvdb.TimerStart ();
gvdb.Render ( SHADE_VOLUME, 0, 0 );
float rtime = gvdb.TimerStop();
nvprintf ( "Render volume. %6.3f ms\n", rtime );

It seems the call Render return before the actual ray casting was done.
Resulting measured sub millisecond performance which I was not expecting.
I think something needs to be add to wait for return of the operation on GPU.

Problems using DepthBuf inputs

Hi,

I'm trying to use DepthBuf as inputs for my use case. The code that I've written is similar to the gRenderKernel sample which actually is working correctly in my computer.

The host part looks like the following:

	int w = 1024, h = 768;

  	VolumeGVDB gvdb;

	int devid = -1;
	gvdb.SetCudaDevice ( devid );
	gvdb.Initialize ();
	gvdb.getScene()->SetVolumeRange ( 0.01f, 1.0f, 0.0f );

	Camera3D* cam = new Camera3D;						
	cam->setOrbit ( Vector3DF(20,30,0), Vector3DF(125,160,125), 800, 1.0f );		
	gvdb.getScene()->SetCamera( cam );	
	gvdb.getScene()->SetRes ( w, h );
	
	Light* lgt = new Light;	
	lgt->setOrbit ( Vector3DF(50,65,0), Vector3DF(125,140,125), 200, 1.0f );
	gvdb.getScene()->SetLight ( 0, lgt );		
	
	gvdb.AddDepthBuf ( 3, w, h );

       void* data;
       // do some stuff to get image in data pointer
     
       GLuint gl_depth_raw_tex;
       glGenBuffers(0, &gl_depth_raw_tex);
       glBindTexture(GL_TEXTURE_2D, gl_depth_raw_tex);
       glTexImage2D(GL_TEXTURE_2D, 0, GL_DEPTH_COMPONENT, w, h, 0, GL_DEPTH_COMPONENT, GL_FLOAT, data);
        glBindTexture(GL_TEXTURE_2D, 0);

	cudaCheck ( cuModuleLoad ( &cuCustom, "custom.ptx" ), "cuModuleLoad (render_custom)" );
	cudaCheck ( cuModuleGetFunction ( &cuRaycastKernel, cuCustom, "kernel" ), "cuModuleGetFunction (kernel)" );	

	gvdb.SetModule ( cuCustom );

	gvdb.getScene()->SetSteps ( 0.25f, 16, 0.25f );
	gvdb.getScene()->SetVolumeRange ( 0.1f, 0.0, 1.0f );
	gvdb.RenderKernel ( cuRaycastKernel, 0, 3 );

       gvdb.WriteDepthTexGL(3, gl_depth_raw_tex);

And the kernel part is like this:

extern "C" __global__ void allocate_bricks (VDBInfo* gvdb, uchar chan, uchar4* outBuf )
{
	int u = blockIdx.x * blockDim.x + threadIdx.x;
	int v = blockIdx.y * blockDim.y + threadIdx.y;
	printf("%d %d\n", scn.width, scn.height);
	if ( u >= scn.width || v >= scn.height ) return;

       // do something
}

And the problem that I'm facing is that width and height values of scn (ScnInfo) (can be used through macros SCN_WIDTH and SCN_HEIGHT) are 0, so this kernel is doing nothing. I've been reading the gvdb core to find some clues but now I'm more confused because actually when RenderKernel is called, PrepareRender is executed setting up mScnInfo and then transferring the new ScnInfo to device.

Then, does anyone know where I'm losing?

Macro Redefinition of RGBA2INT and CLR2INT

There is redefinition of macro RGBA2INT and CLR2INT in cuda_gvdb_operators.cuh and cuda_gvdb_particles.cuh. They have the same function, but are written as addition and bit operation respectively.

activate space function

Is it safe to use mValue for nodes at intermediate levels ?

Hi !

I would like to know if I can put arbitrary custom data inside mValue for nodes that are not bricks ?

Thanks in advance !

GVDB_VOX macro and aprons

gvdb-voxels/source/gvdb_library/kernels/cuda_gvdb_operators.cuh

Lines 298 to 300 in e9a2350

	#define GVDB_VOX \
	uint3 vox = blockIdx * make_uint3(blockDim.x, blockDim.y, blockDim.z) + threadIdx + make_uint3(1,1,1); \
	if ( vox.x >= res.x|| vox.y >= res.y || vox.z >= res.z ) return;

Here GVDB_VOX seems to cover the apron voxel in atlas as well. Is it by design?
Firstly the + make_uint3(1,1,1); part seems to indicate that GVDB_VOX only supports one layer apron setup.
Secondly vox is a continuous uint3 which covers apron voxels in atlas as well. Despite slight increase of computation amount, this will typically cause no problem because even the user directly changes value of apron voxel using vox as index entry in a kernel function, UpdateApron() will still overwrite the aprons to the correct value. But what if two successive operations e.g. adding noise then smoothing are manually coded in a kernel without calling UpdateApron() in between, there could be a discrepancy of apron values.

My suggestion is to modify uint3 vox as

uint3 vox= blockIdx * blockDim + threadIdx;
uint3 brick_idx = vox/ (brick_res - apron - apron);
vox += brick_idx * apron * 2 + apron;

so this can skip the apron voxels when going through atlas.

ActivateSpace Crashes on -ve coordinates

calling gvdb.ActivateSpace( v ) with any of v.x , v.y or v.z being negative seems to crash. Maybe its the order that I am calling methods. Can send an example if needed.

Some topology configuration return CUDA ERROR "invalid argument"

I've been trying some different GVDB configurations in g3DPrint example and I get the CUDA ERROR "invalid argument" most of the times.
I think this problem come from some bad kernel param configuration. It will welcome any other guess.

Texture sampling in compute shaders

Hi, I've been working with GVDB for a few months and I've been running into some weird behavior sampling from the atlas data. We're using an encoded UCHAR and a FLOAT channel to store the data for our compute kernels. However, we've run into issues trying to sample neighbors with the UCHAR where the samples are sometimes incorrect.
Looking at the GVDB_COPY_SMEM_F macro I noticed that make_float3(0.5f,0.5f,0.5f) is added to the vox variable. However, with the other macros like GVDB_VOX and GVDB_COPY_SMEM_UC this is not added. I was wondering if there is a reason why (0.5f,0.5f,0.5f) is added before sampling in the GVDB_COPY_SMEM_F macro? Should this always be added when sampling the atlas, or only for floats, and why? It seems like adding (0.5,0.5,0.5) should perform interpolation on the sample, not giving us the exact value.
Thanks.