openvdb/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu at f269d2a1dda2ba4709e2131e0086a1e50596118d · harrism/openvdb · GitHub

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
// Copyright Contributors to the OpenVDB Project
// SPDX-License-Identifier: Apache-2.0

#ifndef _USE_MATH_DEFINES
#define _USE_MATH_DEFINES
#endif
#include <cmath>
#include <chrono>

#if defined(NANOVDB_USE_CUDA)
#include <nanovdb/cuda/DeviceBuffer.h>
using BufferT = nanovdb::cuda::DeviceBuffer;
#else
using BufferT = nanovdb::HostBuffer;
#endif
#include <nanovdb/GridHandle.h>
#include <nanovdb/io/IO.h>
#include <nanovdb/math/Ray.h>
#include <nanovdb/math/HDDA.h>

#include "common.h"

void runNanoVDB(nanovdb::GridHandle<BufferT>& handle, int numIterations, int width, int height, BufferT& imageBuffer, bool usePersistentThreads)
{
    float *h_outImage = reinterpret_cast<float*>(imageBuffer.data());
    RenderOp renderOp(handle, width, height);

    auto kernel = [&](auto *h_grid){
        float sum = 0;
        for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImage(false/*useCuda*/, h_outImage, h_grid));
        std::cout << "Average of " << numIterations << " renderings (NanoVDB-Host) = " << (sum/numIterations) << " ms" << std::endl;
        renderOp.saveImage("raytrace_iso_surface-nanovdb-host.pfm", (float*)imageBuffer.data());

#if defined(NANOVDB_USE_CUDA)
        handle.deviceUpload();
        using BuildT = typename nanovdb::util::remove_pointer_t<decltype(h_grid)>::BuildType;
        auto* d_grid = handle.deviceGrid<BuildT>();
        if (!d_grid) throw std::runtime_error("GridHandle does not contain a valid device grid");
        imageBuffer.deviceUpload();
        float* d_outImage = reinterpret_cast<float*>(imageBuffer.deviceData());
        sum = 0;
        if (usePersistentThreads) {
            int* d_nextPixel = nullptr;
            NANOVDB_CUDA_CHECK_ERROR(cudaMalloc(&d_nextPixel, sizeof(int)), __FILE__, __LINE__);
            for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImagePersistent(d_outImage, d_grid, d_nextPixel));
            NANOVDB_CUDA_CHECK_ERROR(cudaFree(d_nextPixel), __FILE__, __LINE__);
            std::cout << "Average of " << numIterations << " renderings (NanoVDB-Cuda-Persistent) = " << (sum/numIterations) << " ms " << std::endl;
            imageBuffer.deviceDownload();
            renderOp.saveImage("raytrace_iso_surface-nanovdb-cuda-persistent.pfm", (float*)imageBuffer.data());
        } else {
            for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImage(true/*useCuda*/, d_outImage, d_grid));
            std::cout << "Average of " << numIterations << " renderings (NanoVDB-Cuda) = " << (sum/numIterations) << " ms " << std::endl;
            imageBuffer.deviceDownload();
            renderOp.saveImage("raytrace_iso_surface-nanovdb-cuda.pfm", (float*)imageBuffer.data());
        }
#endif
    };// kernel

    if (auto *h_grid = handle.grid<float>()) {
        kernel(h_grid);
    } else if (auto *h_grid = handle.grid<nanovdb::ValueIndex>()) {
        kernel(h_grid);
    } else if (auto *h_grid = handle.grid<nanovdb::ValueOnIndex>()) {
        kernel(h_grid);
    } else {
        throw std::runtime_error("GridHandle does not contain a valid device grid");
    }
}// runNanoVDB