nvblox
repository·public·Indexed 23 days ago
https://github.com/nvidia-isaac/nvbloxA high-performance library for real-time 3D reconstruction optimized for NVIDIA hardware. It enables building, manipulating, and querying reconstructions on the GPU and provides interfaces for C++, Python (via nvblox_torch), and ROS2. The library includes tools for TSDF voxel grid integration, surface and ESDF evaluation, and visualization modules for meshes, pointclouds, and voxel grids.
What's inside nvblox
- nvblox is a library designed for real-time 3D reconstruction in robotic applications. It allows users to build, manipulate, and query reconstructions directly on the GPU. The library is optimized for high runtime performance using CUDA and NVIDIA hardware.
Understand the limitations of nvblox_torch
publicAs of the first official release,
nvblox_torchhas several functional and performance differences compared to thenvbloxcore library and ROS wrapper:Functional Differences
- Dynamic Scene Elements: Unlike the core library and ROS wrapper,
nvblox_torchdoes not yet support mapping in the presence of moving elements (e.g., people segmentation). - Incremental Visualization:
nvblox_torchdoes not support incremental visualization (streaming only parts of the visualization, such as the mesh, to the pipeline).
Performance and Resource Constraints
- Compute Performance:
nvblox_torchprovides an easy-to-use interface but is not as highly optimized for performance as the core library, though it utilizes zero-copy interfaces. - Memory Usage: Memory consumption scales with the volume of mapped space and increases cubically with resolution. Mapping volumes or resolutions that exceed available GPU memory will cause the library to crash.
- Deep Feature Reconstruction: This feature is particularly memory-intensive because it requires storing long channel length features in 3D voxels.
- Reconstruction Quality: The quality of the reconstruction is strictly dependent on the quality of the input depth maps.
- Dynamic Scene Elements: Unlike the core library and ROS wrapper,
Windowed vs Headless rendering modes
publicThe renderer can be used in two primary modes:
- Windowed mode: Opens a GLFW window with interactive arcball camera controls. Controls include:
- Left-drag: Rotate camera
- Right-drag: Pan camera
- Scroll: Zoom
Rkey: Reset camera
- Headless mode: Renders offscreen without a window. This is ideal for testing, CI, or server-side rendering. In this mode, you must control the viewpoint using
ViewCameramethods directly.
- Windowed mode: Opens a GLFW window with interactive arcball camera controls. Controls include:
Access voxels without copying (GPU-side API)
publicFor high-performance applications, you should access voxels directly within GPU kernels to avoid CPU-GPU synchronization and memory copy overhead.
To do this, you must first obtain a GPU view of the map's hash table using
layer.getGpuLayerView(CudaStreamOwning()). This view provides access to the underlying hash map used to transform 3D coordinates into voxel memory locations.Inside your CUDA kernel, use the
getVoxelAtPosition<T>helper function. This function takes theblock_hash, thequery_location, and theblock_size, and returns a pointer to the voxel if it has been allocated.// 1. Get the GPU view of the layer GPULayerView<TsdfBlock> gpu_layer_view = layer.getGpuLayerView(CudaStreamOwning()); // 2. Use the hash in a kernel __global__ void queryVoxelsKernel( int num_queries, Index3DDeviceHashMapType<TsdfBlock> block_hash, float block_size, const Vector3f* query_locations_ptr, TsdfVoxel* voxels_ptr, bool* success_flags_ptr) { const int idx = threadIdx.x + blockIdx.x * blockDim.x; if (idx >= num_queries) return; const Vector3f query_location = query_locations_ptr[idx]; TsdfVoxel* voxel; // Use getVoxelAtPosition to find the voxel in the hash map if (!getVoxelAtPosition<TsdfVoxel>(block_hash, query_location, block_size, &voxel)) { success_flags_ptr[idx] = false; } else { success_flags_ptr[idx] = true; voxels_ptr[idx] = *voxel; } }How the nvblox hierarchical sparse voxel grid is structured
publicnvblox implements a hierarchical sparse voxel grid to manage memory and performance. The hierarchy consists of the following levels:
- LayerCake: The top-level container. It holds multiple
Layersthat are colocated voxel grids. Each layer contains a different type of mapped quantity (e.g., TSDF or ESDF). - Layer: A sparse voxel grid containing a single type of mapped quantity. A layer is composed of a sparse collection of
(Voxel)Blocks. - (Voxel)Blocks: The grid elements within a layer. Each block defines a small cubical region of space where voxels are densely allocated in an 8x8x8 grid.
- Voxels: The smallest unit of resolution, holding a single value of the mapped quantity (e.g., the TSDF value).
Benefits of this structure:
- Sparsity: Memory is only allocated for mapped voxels, allowing the map to grow and shrink dynamically.
- Data Locality: Each
Blockis stored in a contiguous chunk of memory. This allows GPU kernels to process singleVoxelBlocksusing CUDA ThreadBlocks, benefiting from coalesced global GPU memory access. - Extensibility: Users can define custom
Layersto map new quantities alongside the built-in ones.
- LayerCake: The top-level container. It holds multiple
Understand the NvbloxRenderer lifecycle
publicThe
NvbloxRendererfollows a multi-stage lifecycle. Splitting initialization from construction allows for safer error handling (viaboolreturns) and the ability to re-initialize the object (e.g., to change resolution).Lifecycle Stages:
- Construct: Default construction. No GPU/Vulkan resources are allocated.
- Initialize: Call
init(),initWithWindow(), orinitHeadless(). This allocates Vulkan and GPU resources. Returnsfalseon failure. - Initialize visualizers: Call
initVisualizer(mode)for eachRenderModeyou intend to use. - Render loop: For each frame:
- Push data via
update*methods on a CUDA stream. - Synchronize the stream.
- Call
render(). - Call
pollEvents()(windowed mode only).
- Push data via
- Destroy: The destructor calls
destroy()automatically, but you can call it explicitly to release resources early. You can re-initialize the same object after callingdestroy().
Handle Thread Safety and CUDA-Vulkan Synchronization
publicThread Safety
NvbloxRendererand all visualizers are not thread-safe. All calls (init, destroy, render, update) must originate from the same thread.- Windowed Mode: Must use the program's main thread (required for GLFW).
- Headless Mode: Any single thread is acceptable.
CUDA to Vulkan Synchronization
To ensure data integrity when updating the renderer from CUDA, follow this sequence:
- Call the update methods (
updateDepth,updateColor,updatePointCloud,updateMesh, orupdateMeshTexture) using a CUDA stream. - Call
stream.synchronize()to ensure all CUDA writes are complete. - Call
render()to submit Vulkan commands.
How to use the nvblox_torch Mapper in a RealSense loop
publicTo build a reconstruction loop with
nvblox_torchand RealSense, follow this pattern:- Initialize Parameters: Configure
ProjectiveIntegratorParamsandMapperParams. - Initialize Mapper: Create the
Mapperobject with desired voxel sizes and integrator types. - Track Pose: Use
PyCuVSLAMto get the camera pose (T_W_C_left_infrared) from stereo infrared images. - Add Depth: Pass the depth frame and the tracked pose to
nvblox_mapper.add_depth_frame(). - Add Color: Calculate the color camera pose by applying the extrinsic calibration matrix to the infrared pose (
T_W_C_color = T_W_C_left_infrared @ T_C_left_infrared_C_color), then callnvblox_mapper.add_color_frame(). - Update Mesh: Periodically call
nvblox_mapper.update_color_mesh()andnvblox_mapper.get_color_mesh()to retrieve the reconstructed mesh for visualization.
# Initialize nvblox mapper projective_integrator_params = ProjectiveIntegratorParams() projective_integrator_params.projective_integrator_max_integration_distance_m = args.max_integration_distance_m mapper_params = MapperParams() mapper_params.set_projective_integrator_params(projective_integrator_params) nvblox_mapper = Mapper(voxel_sizes_m=args.voxel_size_m, integrator_types=ProjectiveIntegratorType.TSDF, mapper_parameters=mapper_params) # Processing loop for _, frame in enumerate(realsense_dataloader): # 1. Track pose if frame['left_infrared_image'] is not None and frame['right_infrared_image'] is not None: T_W_C_left_infrared = cuvslam_tracker.track( frame['timestamp'], (frame['left_infrared_image'], frame['right_infrared_image'])) # 2. Add depth if frame['depth'] is not None and T_W_C_left_infrared is not None: nvblox_mapper.add_depth_frame(frame['depth'], T_W_C_left_infrared, depth_intrinsics) # 3. Add color if T_W_C_left_infrared is not None and frame['rgb'] is not None: T_W_C_color = T_W_C_left_infrared @ T_C_left_infrared_C_color nvblox_mapper.add_color_frame(frame['rgb'], T_W_C_color, color_intrinsics) # 4. Visualize nvblox_mapper.update_color_mesh() color_mesh = nvblox_mapper.get_color_mesh()- Initialize Parameters: Configure
Use NvbloxRenderer in windowed mode
publicWhen running in windowed mode, the
render()method handles several edge cases automatically:- Minimized window: Returns
trueimmediately without drawing or erroring. - Window resize: Automatically recreates the swapchain. One frame is skipped after resize, and the
ViewCameraaspect ratio updates automatically from the new framebuffer size.
Note: The background clear color is fixed at dark gray
(0.1, 0.1, 0.1).- Minimized window: Returns
How sparse (direct) voxel access works
publicSparse access allows you to interact with voxels directly in the
TsdfLayerwithout copying them into a dense tensor. This is more memory-efficient and faster for large maps but requires iterating over allocated blocks.Internal Structure:
- The world is divided into a grid of
VoxelBlocks (size 8x8x8). - Memory is only allocated for blocks that have been observed.
nvblox_torchprovidestorchwrappers around these blocks, so accessing them does not trigger a memory copy.
Warning:
nvblox_torchdoes not currently prevent dangling references to deletedVoxelBlocks. If you call functions that might delete blocks (likemapper.clear()ormapper.decay*()), you must re-acquire the block tensors.- The world is divided into a grid of
How dense (copy-based) voxel access works
publicDense access is the most straightforward method but requires copying the sparse internal representation into a dense
torchtensor. This is suitable if speed and memory are not primary concerns.Internally,
nvbloxstores voxels sparsely inVoxelBlocks (8x8x8) to save memory. To use dense access, you must convert theTsdfLayerinto a dense tensor representing the Axis-Aligned Bounding Box (AABB) of the observed voxels.Understand nvblox map representations
publicnvblox primarily uses a Truncated Signed Distance Function (TSDF) stored in a 3D voxel grid for map reconstruction.
Key characteristics of the TSDF approach:
- Surface Extraction: The environment surface is extracted as the zero-level set of the voxelized function.
- Quality: Typically provides higher quality surface reconstructions compared to occupancy grids.
- Planning Utility: Distance fields provide immediate collision checking for potential robot positions, making them highly useful for path planning.
In addition to TSDF, nvblox also supports occupancy grids.