The example snippets below are pulled directly from compiled example source files via Doxygen's \snippet mechanism, so the code shown here is always in sync with what is built and tested in CI. Enable -DNANOVDB_BUILD_EXAMPLES=ON to build them locally.
Convert an OpenVDB grid to NanoVDB and use accessors
The first example shows how to convert an OpenVDB level set sphere into a NanoVDB level set, use accessors to print out multiple values from both grids, and save the NanoVDB grid to file. Note that this example depends on both OpenVDB and NanoVDB.
#include <nanovdb/tools/CreateNanoGrid.h>
#include <nanovdb/io/IO.h>
int main()
{
try {
auto srcGrid = openvdb::tools::createLevelSetSphere<openvdb::FloatGrid>(100.0f,
openvdb::Vec3f(0.0f), 1.0f);
auto handle = nanovdb::tools::createNanoGrid(*srcGrid);
auto* dstGrid = handle.grid<float>();
if (!dstGrid)
throw std::runtime_error("GridHandle does not contain a grid with value type float");
auto dstAcc = dstGrid->getAccessor();
auto srcAcc = srcGrid->getAccessor();
for (int i = 97; i < 104; ++i) {
printf(
"(%3i,0,0) OpenVDB cpu: % -4.2f, NanoVDB cpu: % -4.2f\n", i, srcAcc.getValue(
openvdb::Coord(i, 0, 0)), dstAcc.getValue(nanovdb::Coord(i, 0, 0)));
}
nanovdb::io::writeGrid("data/sphere.nvdb", handle);
}
catch (const std::exception& e) {
std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl;
}
return 0;
}
Read a NanoVDB grid from a file (host only)
The second example reads a NanoVDB grid from a file (the one saved in the previous example) and prints out multiple values. Note that this example does not depend on OpenVDB (nor CUDA), only NanoVDB.
#include <nanovdb/io/IO.h>
int main()
{
try {
auto handle = nanovdb::io::readGrid("data/sphere.nvdb");
auto* grid = handle.grid<float>();
if (!grid)
throw std::runtime_error("File did not contain a grid with value type float");
auto acc = grid->getAccessor();
for (int i = 97; i < 104; ++i) {
printf("(%3i,0,0) NanoVDB cpu: % -4.2f\n", i, acc.getValue(nanovdb::Coord(i, 0, 0)));
}
}
catch (const std::exception& e) {
std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl;
}
return 0;
}
Read a NanoVDB grid and access it on the CPU and GPU
The third example reads a NanoVDB grid from a file (the one saved in the first example) and prints out multiple values on both the CPU and GPU. Note that this example does NOT depend on OpenVDB, only NanoVDB and CUDA.
#include <nanovdb/io/IO.h> // this is required to read (and write) NanoVDB files on the host
#include <nanovdb/cuda/DeviceBuffer.h> // required for CUDA memory management
#include <nanovdb/GridHandle.h>
extern "C" void launch_kernels(const nanovdb::NanoGrid<float>*,
const nanovdb::NanoGrid<float>*,
cudaStream_t stream);
/// @brief Read a NanoVDB grid from a file and print out multiple values on both the cpu and gpu.
///
/// @note Note This example does NOT depend on OpenVDB, only NanoVDB and CUDA.
int main(int, char**)
{
try {
// returns a GridHandle using CUDA for memory management.
auto handle = nanovdb::io::readGrid<nanovdb::CudaDeviceBuffer>("data/sphere.nvdb");
cudaStream_t stream; // Create a CUDA stream to allow for asynchronous copy of pinned CUDA memory.
cudaStreamCreate(&stream);
handle.deviceUpload(stream, false); // Copy the NanoVDB grid to the GPU asynchronously
auto* cpuGrid = handle.grid<float>(); // get a (raw) pointer to a NanoVDB grid of value type float on the CPU
auto* deviceGrid = handle.deviceGrid<float>(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU
if (!deviceGrid || !cpuGrid)
throw std::runtime_error("GridHandle did not contain a grid with value type float");
launch_kernels(deviceGrid, cpuGrid, stream); // Call a host method to print a grid values on both the CPU and GPU
cudaStreamDestroy(stream); // Destroy the CUDA stream
}
catch (const std::exception& e) {
std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl;
}
return 0;
}
NVCC requires the CUDA kernel to be defined in a separate .cu file:
#include <nanovdb/NanoVDB.h> // this defined the core tree data structure of NanoVDB accessable on both the host and device
#include <stdio.h> // for printf
// This is called by the host only
void cpu_kernel(const nanovdb::NanoGrid<float>* cpuGrid)
{
auto cpuAcc = cpuGrid->getAccessor();
for (int i = 97; i < 104; ++i) {
printf("(%3i,0,0) NanoVDB cpu: % -4.2f\n", i, cpuAcc.getValue(nanovdb::Coord(i, 0, 0)));
}
}
// This is called by the device only
__global__ void gpu_kernel(const nanovdb::NanoGrid<float>* deviceGrid)
{
if (threadIdx.x > 6)
return;
int i = 97 + threadIdx.x;
auto gpuAcc = deviceGrid->getAccessor();
printf("(%3i,0,0) NanoVDB gpu: % -4.2f\n", i, gpuAcc.getValue(nanovdb::Coord(i, 0, 0)));
}
// This is called by the client code on the host
extern "C" void launch_kernels(const nanovdb::NanoGrid<float>* deviceGrid,
const nanovdb::NanoGrid<float>* cpuGrid,
cudaStream_t stream)
{
// Launch the device kernel asynchronously
gpu_kernel<<<1, 64, 0, stream>>>(deviceGrid);
// Launch the host "kernel" (synchronously)
cpu_kernel(cpuGrid);
}