nvblox

repository·public·Indexed 23 days ago

https://github.com/nvidia-isaac/nvblox

A 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.

Tokens
24.9K
Snippets
47
Records
133
Agent score
78%

What's inside nvblox

  1. Overview of nvblox for real-time 3D reconstruction

    public
    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.
  2. Understand the limitations of nvblox_torch

    public

    As of the first official release, nvblox_torch has several functional and performance differences compared to the nvblox core library and ROS wrapper:

    Functional Differences

    • Dynamic Scene Elements: Unlike the core library and ROS wrapper, nvblox_torch does not yet support mapping in the presence of moving elements (e.g., people segmentation).
    • Incremental Visualization: nvblox_torch does not support incremental visualization (streaming only parts of the visualization, such as the mesh, to the pipeline).

    Performance and Resource Constraints

    • Compute Performance: nvblox_torch provides 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.
  3. Windowed vs Headless rendering modes

    public

    The 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
      • R key: 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 ViewCamera methods directly.
  4. Access voxels without copying (GPU-side API)

    public

    For 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 the block_hash, the query_location, and the block_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;
      }
    }
  5. How the nvblox hierarchical sparse voxel grid is structured

    public

    nvblox implements a hierarchical sparse voxel grid to manage memory and performance. The hierarchy consists of the following levels:

    1. LayerCake: The top-level container. It holds multiple Layers that are colocated voxel grids. Each layer contains a different type of mapped quantity (e.g., TSDF or ESDF).
    2. Layer: A sparse voxel grid containing a single type of mapped quantity. A layer is composed of a sparse collection of (Voxel)Blocks.
    3. (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.
    4. 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 Block is stored in a contiguous chunk of memory. This allows GPU kernels to process single VoxelBlocks using CUDA ThreadBlocks, benefiting from coalesced global GPU memory access.
    • Extensibility: Users can define custom Layers to map new quantities alongside the built-in ones.
  6. Understand the NvbloxRenderer lifecycle

    public

    The NvbloxRenderer follows a multi-stage lifecycle. Splitting initialization from construction allows for safer error handling (via bool returns) and the ability to re-initialize the object (e.g., to change resolution).

    Lifecycle Stages:

    1. Construct: Default construction. No GPU/Vulkan resources are allocated.
    2. Initialize: Call init(), initWithWindow(), or initHeadless(). This allocates Vulkan and GPU resources. Returns false on failure.
    3. Initialize visualizers: Call initVisualizer(mode) for each RenderMode you intend to use.
    4. Render loop: For each frame:
      • Push data via update* methods on a CUDA stream.
      • Synchronize the stream.
      • Call render().
      • Call pollEvents() (windowed mode only).
    5. Destroy: The destructor calls destroy() automatically, but you can call it explicitly to release resources early. You can re-initialize the same object after calling destroy().
  7. Handle Thread Safety and CUDA-Vulkan Synchronization

    public

    Thread Safety

    NvbloxRenderer and 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:

    1. Call the update methods (updateDepth, updateColor, updatePointCloud, updateMesh, or updateMeshTexture) using a CUDA stream.
    2. Call stream.synchronize() to ensure all CUDA writes are complete.
    3. Call render() to submit Vulkan commands.
  8. How to use the nvblox_torch Mapper in a RealSense loop

    public

    To build a reconstruction loop with nvblox_torch and RealSense, follow this pattern:

    1. Initialize Parameters: Configure ProjectiveIntegratorParams and MapperParams.
    2. Initialize Mapper: Create the Mapper object with desired voxel sizes and integrator types.
    3. Track Pose: Use PyCuVSLAM to get the camera pose (T_W_C_left_infrared) from stereo infrared images.
    4. Add Depth: Pass the depth frame and the tracked pose to nvblox_mapper.add_depth_frame().
    5. 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 call nvblox_mapper.add_color_frame().
    6. Update Mesh: Periodically call nvblox_mapper.update_color_mesh() and nvblox_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()
  9. Use NvbloxRenderer in windowed mode

    public

    When running in windowed mode, the render() method handles several edge cases automatically:

    • Minimized window: Returns true immediately without drawing or erroring.
    • Window resize: Automatically recreates the swapchain. One frame is skipped after resize, and the ViewCamera aspect ratio updates automatically from the new framebuffer size.

    Note: The background clear color is fixed at dark gray (0.1, 0.1, 0.1).

  10. How sparse (direct) voxel access works

    public

    Sparse access allows you to interact with voxels directly in the TsdfLayer without 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_torch provides torch wrappers around these blocks, so accessing them does not trigger a memory copy.

    Warning: nvblox_torch does not currently prevent dangling references to deleted VoxelBlocks. If you call functions that might delete blocks (like mapper.clear() or mapper.decay*()), you must re-acquire the block tensors.

  11. How dense (copy-based) voxel access works

    public

    Dense access is the most straightforward method but requires copying the sparse internal representation into a dense torch tensor. This is suitable if speed and memory are not primary concerns.

    Internally, nvblox stores voxels sparsely in VoxelBlocks (8x8x8) to save memory. To use dense access, you must convert the TsdfLayer into a dense tensor representing the Axis-Aligned Bounding Box (AABB) of the observed voxels.

  12. Understand nvblox map representations

    public

    nvblox 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.