Skip to content

Commit fb4516e

Browse files
kmusethkmuseth
authored
Squash NanoVDB and vdb_tool contributions (#2241)
Signed-off-by: Ken Museth <ken.museth@gmail.com>
1 parent a532de5 commit fb4516e

48 files changed

Lines changed: 11017 additions & 1835 deletions

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎nanovdb/nanovdb/NanoVDB.h‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4736,7 +4736,7 @@ using OnIndexGrid = Grid<OnIndexTree>;
47364736
* @endcode
47374737
*/
47384738

4739-
/// @brief Use this function, which depends a pointer to GridData, to call
4739+
/// @brief Use this function, which depends on a pointer to GridData, to call
47404740
/// other functions that depend on a NanoGrid of a known ValueType.
47414741
/// @details This function allows for generic programming by converting GridData
47424742
/// to a NanoGrid of the type encoded in GridData::mGridType.
@@ -5780,6 +5780,26 @@ class ChannelAccessor : public DefaultReadAccessor<IndexT>
57805780

57815781
}; // ChannelAccessor
57825782

5783+
/// @brief Generic Accessor type that maps to either a ReadAccessor or ChannelAccessor
5784+
/// @tparam BuildT Build type, e.g. float or ValueOnIndex
5785+
/// @tparam ValueT Value type, e.g. float or Vec3f
5786+
template <typename BuildT, typename ValueT>
5787+
using AccType = typename util::conditional<BuildTraits<BuildT>::is_index,
5788+
ChannelAccessor<ValueT, BuildT>, DefaultReadAccessor<BuildT>>::type;
5789+
5790+
/// @brief Generic template functions that return an Accessor to either an index grid or a regular grid
5791+
template <typename GridT, typename ValueT>
5792+
inline __hostdev__ auto getAccessor(const GridT &grid, ValueT *sideCar = nullptr)
5793+
{
5794+
using BuildT = typename GridT::BuildType;
5795+
if constexpr(BuildTraits<BuildT>::is_index) {
5796+
return sideCar ? ChannelAccessor<ValueT, BuildT>(grid, sideCar) : ChannelAccessor<ValueT, BuildT>(grid);
5797+
} else {
5798+
static_assert(util::is_same<ValueT, typename GridT::ValueType>::value, "wrong ValueT for regular GridT");
5799+
return DefaultReadAccessor<BuildT>(grid);
5800+
}
5801+
}
5802+
57835803
#if 0
57845804
// This MiniGridHandle class is only included as a stand-alone example. Note that aligned_alloc is a C++17 feature!
57855805
// Normally we recommend using GridHandle defined in util/GridHandle.h but this minimal implementation could be an

‎nanovdb/nanovdb/examples/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -109,6 +109,7 @@ nanovdb_example(NAME "ex_bump_pool_buffer")
109109
nanovdb_example(NAME "ex_collide_level_set")
110110
nanovdb_example(NAME "ex_raytrace_fog_volume")
111111
nanovdb_example(NAME "ex_raytrace_level_set")
112+
nanovdb_example(NAME "ex_raytrace_iso_surface")
112113
nanovdb_example(NAME "ex_dilate_nanovdb_cuda" OPENVDB)
113114
nanovdb_example(NAME "ex_merge_nanovdb_cuda" OPENVDB)
114115
nanovdb_example(NAME "ex_refine_nanovdb_cuda" OPENVDB)

‎nanovdb/nanovdb/examples/ex_raytrace_iso_surface/common.h‎

Lines changed: 187 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,187 @@
1+
// Copyright Contributors to the OpenVDB Project
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
#pragma once
5+
6+
#define _USE_MATH_DEFINES
7+
#include <cmath>
8+
#include <chrono>
9+
#include <fstream>
10+
#include <nanovdb/NanoVDB.h>
11+
#include "ComputePrimitives.h"
12+
13+
struct RenderOp;
14+
template<typename GridT>
15+
__global__ void renderIsoSurfacePersistentKernel(RenderOp renderOp, float* image, const GridT* grid, int numPixels, int* nextPixel);
16+
17+
struct RenderOp
18+
{
19+
using Vec3T = nanovdb::math::Vec3<float>;
20+
using RayT = nanovdb::math::Ray<float>;
21+
int mWidth, mHeight;
22+
float mDx, mIso, mWBBoxDimZ;
23+
Vec3T mWBBoxCenter;
24+
25+
template<typename BufferT>
26+
RenderOp(nanovdb::GridHandle<BufferT>& handle, int width, int height)
27+
{
28+
mWidth = width;
29+
mHeight = height;
30+
const auto *metaData = handle.gridMetaData();
31+
mDx = float(metaData->voxelSize()[0]);
32+
mIso = mDx;
33+
mWBBoxDimZ = (float)metaData->worldBBox().dim()[2] * 2;
34+
mWBBoxCenter = Vec3T(metaData->worldBBox().min() + metaData->worldBBox().dim() * 0.5f);
35+
}
36+
37+
template<typename GridT>
38+
inline float renderImage(bool useCuda, float* image, const GridT* grid)
39+
{
40+
using ClockT = std::chrono::high_resolution_clock;
41+
auto t0 = ClockT::now();
42+
43+
computeForEach(
44+
useCuda, mWidth * mHeight, 512, __FILE__, __LINE__, [this, image, grid] __hostdev__(int start, int end) {
45+
(*this)(start, end, image, grid);
46+
});
47+
computeSync(useCuda, __FILE__, __LINE__);
48+
49+
auto t1 = ClockT::now();
50+
auto duration = std::chrono::duration_cast<std::chrono::microseconds>(t1 - t0).count() / 1000.f;
51+
return duration;
52+
}
53+
54+
template <typename GridT>
55+
inline __hostdev__ void operator()(int start, int end, float* image, const GridT* grid) const
56+
{
57+
static_assert(nanovdb::util::is_same<typename GridT::BuildType, float, nanovdb::ValueOnIndex, nanovdb::ValueIndex>::value, "only works for float and OnIndex grids");
58+
auto acc = nanovdb::getAccessor<GridT, float>(*grid);
59+
for (int i = start; i < end; ++i) {
60+
this->renderPixel(i, image, grid, acc);
61+
}
62+
}
63+
64+
template <typename GridT, typename AccT>
65+
inline __hostdev__ void renderPixel(int i, float* image, const GridT* grid, AccT& acc) const
66+
{
67+
float t0, v;
68+
nanovdb::Coord ijk;
69+
RayT iRay = this->getIndexRay(i, grid);
70+
if (nanovdb::math::isoCrossing(iRay, acc, ijk, v, t0, mIso)) {// intersect...
71+
this->composite(image, i, (t0 * mDx) / (mWBBoxDimZ * 2), 1.0f);
72+
} else {
73+
this->composite(image, i, 0.0f, 0.0f);// write background value.
74+
}
75+
}
76+
77+
template<typename GridT>
78+
inline float renderImagePersistent(float* image, const GridT* grid, int* nextPixel) const
79+
{
80+
int device = 0;
81+
NANOVDB_CUDA_CHECK_ERROR(cudaGetDevice(&device), __FILE__, __LINE__);
82+
83+
cudaDeviceProp properties;
84+
NANOVDB_CUDA_CHECK_ERROR(cudaGetDeviceProperties(&properties, device), __FILE__, __LINE__);
85+
86+
constexpr int blockSize = 256;
87+
// Launch a small, fixed pool of blocks that persists on the GPU and
88+
// pulls pixel work from a global counter instead of launching one
89+
// logical thread per pixel up front.
90+
int blockCount = properties.multiProcessorCount * 4;
91+
if (blockCount < 1) blockCount = 1;
92+
93+
// Reset the work queue before each timed render. The kernel advances
94+
// this counter by one warp of pixels at a time.
95+
NANOVDB_CUDA_CHECK_ERROR(cudaMemset(nextPixel, 0, sizeof(int)), __FILE__, __LINE__);
96+
97+
using ClockT = std::chrono::high_resolution_clock;
98+
auto t0 = ClockT::now();
99+
100+
const int numPixels = mWidth * mHeight;
101+
renderIsoSurfacePersistentKernel<GridT><<<blockCount, blockSize>>>(*this, image, grid, numPixels, nextPixel);
102+
NANOVDB_CUDA_CHECK_ERROR(cudaGetLastError(), __FILE__, __LINE__);
103+
NANOVDB_CUDA_CHECK_ERROR(cudaDeviceSynchronize(), __FILE__, __LINE__);
104+
105+
auto t1 = ClockT::now();
106+
auto duration = std::chrono::duration_cast<std::chrono::microseconds>(t1 - t0).count() / 1000.f;
107+
return duration;
108+
}
109+
110+
template <typename GridT>
111+
inline __hostdev__ RayT getIndexRay(int i, const GridT *grid) const
112+
{
113+
// perspective camera along Z-axis...
114+
const uint32_t x = i % mWidth, y = i / mWidth;
115+
const float fov = 45.f;
116+
const float u = (float(x) + 0.5f) / mWidth;
117+
const float v = (float(y) + 0.5f) / mHeight;
118+
const float aspect = mWidth / float(mHeight);
119+
const float Px = (2.f * u - 1.f) * tanf(fov / 2 * 3.14159265358979323846f / 180.f) * aspect;
120+
const float Py = (2.f * v - 1.f) * tanf(fov / 2 * 3.14159265358979323846f / 180.f);
121+
const Vec3T origin = mWBBoxCenter + Vec3T(0, 0, mWBBoxDimZ);
122+
Vec3T dir(Px, Py, -1.f);
123+
dir.normalize();
124+
RayT wRay(origin, dir);
125+
return wRay.worldToIndexF(*grid);// transform the ray to the grid's index-space.
126+
}
127+
128+
inline __hostdev__ void composite(float* outImage, int offset, float value, float alpha) const
129+
{
130+
const uint32_t x = offset % mWidth, y = offset / mWidth;
131+
132+
// checkerboard background...
133+
const int mask = 1 << 7;
134+
const float bg = ((x & mask) ^ (y & mask)) ? 1.0f : 0.5f;
135+
outImage[offset] = alpha * value + (1.0f - alpha) * bg;
136+
}
137+
138+
inline void saveImage(const std::string& filename, const float* image) const
139+
{
140+
const auto isLittleEndian = []() -> bool {
141+
static int x = 1;
142+
static bool result = reinterpret_cast<uint8_t*>(&x)[0] == 1;
143+
return result;
144+
};
145+
146+
float scale = 1.0f;
147+
if (isLittleEndian()) scale = -scale;
148+
149+
std::fstream fs(filename, std::ios::out | std::ios::binary);
150+
if (!fs.is_open()) throw std::runtime_error("Unable to open file: " + filename);
151+
152+
fs << "Pf\n"
153+
<< mWidth << "\n"
154+
<< mHeight << "\n"
155+
<< scale << "\n";
156+
157+
for (int i = 0; i < mWidth * mHeight; ++i) {
158+
float r = image[i];
159+
fs.write((char*)&r, sizeof(float));
160+
}
161+
}
162+
};
163+
164+
template<typename GridT>
165+
__global__ void renderIsoSurfacePersistentKernel(RenderOp renderOp, float* image, const GridT* grid, int numPixels, int* nextPixel)
166+
{
167+
static_assert(nanovdb::util::is_same<typename GridT::BuildType, float, nanovdb::ValueOnIndex, nanovdb::ValueIndex>::value, "only works for float and OnIndex grids");
168+
auto acc = nanovdb::getAccessor<GridT, float>(*grid);
169+
const unsigned int lane = threadIdx.x & 31u;
170+
171+
// Keep the fixed set of launched threads busy until all pixels have been assigned.
172+
while (true) {
173+
int base = 0;
174+
// Each warp asks the shared counter for the next batch of 32 pixels.
175+
// Only lane 0 updates the counter; __shfl_sync copies lane 0's result
176+
// to the other lanes in the warp.
177+
if (lane == 0) base = atomicAdd(nextPixel, 32);
178+
base = __shfl_sync(0xFFFFFFFFu, base, 0);
179+
180+
// Each lane renders one pixel from the batch: lane 0 renders base,
181+
// lane 1 renders base + 1, and so on.
182+
const int i = base + int(lane);
183+
if (i >= numPixels) break;
184+
185+
renderOp.renderPixel(i, image, grid, acc);
186+
}
187+
}

‎nanovdb/nanovdb/examples/ex_raytrace_iso_surface/main.cc‎

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,52 @@
1+
// Copyright Contributors to the OpenVDB Project
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
#include <algorithm>
5+
#include <cstring>
6+
#include <iostream>
7+
#include <nanovdb/io/IO.h>
8+
#include <nanovdb/tools/CreatePrimitives.h>
9+
#include <nanovdb/cuda/DeviceBuffer.h>
10+
11+
#if defined(NANOVDB_USE_CUDA)
12+
using BufferT = nanovdb::cuda::DeviceBuffer;
13+
#else
14+
using BufferT = nanovdb::HostBuffer;
15+
#endif
16+
17+
extern void runNanoVDB(nanovdb::GridHandle<BufferT>& handle, int numIterations, int width, int height, BufferT& imageBuffer, bool usePersistentThreads);
18+
19+
int main(int ac, char** av)
20+
{
21+
try {
22+
bool usePersistentThreads = false;
23+
const char* gridName = nullptr;
24+
for (int i = 1; i < ac; ++i) {
25+
if (std::strcmp(av[i], "--persistent") == 0) {
26+
usePersistentThreads = true;
27+
} else if (!gridName) {
28+
gridName = av[i];
29+
} else {
30+
throw std::runtime_error("Usage: ex_raytrace_iso_surface [--persistent] [grid.nvdb]");
31+
}
32+
}
33+
nanovdb::GridHandle<BufferT> handle;
34+
if (gridName) {
35+
handle = nanovdb::io::readGrid<BufferT>(gridName);
36+
std::cout << "Loaded NanoVDB grid[" << handle.gridMetaData()->shortGridName() << "]...\n";
37+
} else {
38+
handle = nanovdb::tools::createLevelSetSphere<float, BufferT>(100.0f, nanovdb::Vec3d(-20, 0, 0), 1.0, 3.0, nanovdb::Vec3d(0), "sphere");
39+
}
40+
41+
const int numIterations = 50;
42+
const int width = 4096;
43+
const int height = 4096;
44+
BufferT imageBuffer(width * height * sizeof(float));
45+
46+
runNanoVDB(handle, numIterations, width, height, imageBuffer, usePersistentThreads);
47+
}
48+
catch (const std::exception& e) {
49+
std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl;
50+
}
51+
return 0;
52+
}

‎nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu‎

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
// Copyright Contributors to the OpenVDB Project
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
#define _USE_MATH_DEFINES
5+
#include <cmath>
6+
#include <chrono>
7+
8+
#if defined(NANOVDB_USE_CUDA)
9+
#include <nanovdb/cuda/DeviceBuffer.h>
10+
using BufferT = nanovdb::cuda::DeviceBuffer;
11+
#else
12+
using BufferT = nanovdb::HostBuffer;
13+
#endif
14+
#include <nanovdb/GridHandle.h>
15+
#include <nanovdb/io/IO.h>
16+
#include <nanovdb/math/Ray.h>
17+
#include <nanovdb/math/HDDA.h>
18+
19+
#include "common.h"
20+
21+
void runNanoVDB(nanovdb::GridHandle<BufferT>& handle, int numIterations, int width, int height, BufferT& imageBuffer, bool usePersistentThreads)
22+
{
23+
float *h_outImage = reinterpret_cast<float*>(imageBuffer.data());
24+
RenderOp renderOp(handle, width, height);
25+
26+
auto kernel = [&](auto *h_grid){
27+
float sum = 0;
28+
for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImage(false/*useCuda*/, h_outImage, h_grid));
29+
std::cout << "Average of " << numIterations << " renderings (NanoVDB-Host) = " << (sum/numIterations) << " ms" << std::endl;
30+
renderOp.saveImage("raytrace_iso_surface-nanovdb-host.pfm", (float*)imageBuffer.data());
31+
32+
#if defined(NANOVDB_USE_CUDA)
33+
handle.deviceUpload();
34+
using BuildT = typename nanovdb::util::remove_pointer_t<decltype(h_grid)>::BuildType;
35+
auto* d_grid = handle.deviceGrid<BuildT>();
36+
if (!d_grid) throw std::runtime_error("GridHandle does not contain a valid device grid");
37+
imageBuffer.deviceUpload();
38+
float* d_outImage = reinterpret_cast<float*>(imageBuffer.deviceData());
39+
sum = 0;
40+
if (usePersistentThreads) {
41+
int* d_nextPixel = nullptr;
42+
NANOVDB_CUDA_CHECK_ERROR(cudaMalloc(&d_nextPixel, sizeof(int)), __FILE__, __LINE__);
43+
for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImagePersistent(d_outImage, d_grid, d_nextPixel));
44+
NANOVDB_CUDA_CHECK_ERROR(cudaFree(d_nextPixel), __FILE__, __LINE__);
45+
std::cout << "Average of " << numIterations << " renderings (NanoVDB-Cuda-Persistent) = " << (sum/numIterations) << " ms " << std::endl;
46+
imageBuffer.deviceDownload();
47+
renderOp.saveImage("raytrace_iso_surface-nanovdb-cuda-persistent.pfm", (float*)imageBuffer.data());
48+
} else {
49+
for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImage(true/*useCuda*/, d_outImage, d_grid));
50+
std::cout << "Average of " << numIterations << " renderings (NanoVDB-Cuda) = " << (sum/numIterations) << " ms " << std::endl;
51+
imageBuffer.deviceDownload();
52+
renderOp.saveImage("raytrace_iso_surface-nanovdb-cuda.pfm", (float*)imageBuffer.data());
53+
}
54+
#endif
55+
};// kernel
56+
57+
if (auto *h_grid = handle.grid<float>()) {
58+
kernel(h_grid);
59+
} else if (auto *h_grid = handle.grid<nanovdb::ValueIndex>()) {
60+
kernel(h_grid);
61+
} else if (auto *h_grid = handle.grid<nanovdb::ValueOnIndex>()) {
62+
kernel(h_grid);
63+
} else {
64+
throw std::runtime_error("GridHandle does not contain a valid device grid");
65+
}
66+
}// runNanoVDB

‎nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu‎

Lines changed: 6 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -20,15 +20,14 @@ using BufferT = nanovdb::HostBuffer;
2020

2121
void runNanoVDB(nanovdb::GridHandle<BufferT>& handle, int numIterations, int width, int height, BufferT& imageBuffer)
2222
{
23-
using GridT = nanovdb::FloatGrid;
23+
using GridT = nanovdb::FloatGrid;
2424
using CoordT = nanovdb::Coord;
25-
using RealT = float;
26-
using Vec3T = nanovdb::math::Vec3<RealT>;
27-
using RayT = nanovdb::math::Ray<RealT>;
25+
using RealT = float;
26+
using Vec3T = nanovdb::math::Vec3<RealT>;
27+
using RayT = nanovdb::math::Ray<RealT>;
2828

2929
auto *h_grid = handle.grid<float>();
30-
if (!h_grid)
31-
throw std::runtime_error("GridHandle does not contain a valid host grid");
30+
if (!h_grid) throw std::runtime_error("GridHandle does not contain a valid host grid");
3231

3332
float* h_outImage = reinterpret_cast<float*>(imageBuffer.data());
3433

@@ -58,7 +57,7 @@ void runNanoVDB(nanovdb::GridHandle<BufferT>& handle, int numIterations, int wid
5857
float t0;
5958
CoordT ijk;
6059
float v;
61-
if (nanovdb::math::ZeroCrossing(iRay, acc, ijk, v, t0)) {
60+
if (nanovdb::math::zeroCrossing(iRay, acc, ijk, v, t0)) {
6261
// write distance to surface. (we assume it is a uniform voxel)
6362
float wT0 = t0 * float(grid->voxelSize()[0]);
6463
compositeOp(image, i, width, height, wT0 / (wBBoxDimZ * 2), 1.0f);

0 commit comments

Comments
 (0)