add vdbio

09/01_texture/01/CMakeLists.txt

@@ -9,3 +9,8 @@ add_executable(main main.cu)
 target_include_directories(main PUBLIC . ../../include)
 target_compile_options(main PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>)
 target_compile_options(main PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>)
+
+target_sources(main PUBLIC vdbio.cpp)
+find_package(OpenVDB REQUIRED)
+message(STATUS "Found OpenVDB ${OpenVDB_VERSION} at ${OpenVDB_LIBRARIES}")
+target_link_libraries(main PUBLIC OpenVDB::openvdb)

09/01_texture/01/main.cu

@@ -1,152 +1,29 @@
 #include <cstdio>
 #include <vector>
-#include <memory>
 #include <cuda_runtime.h>
 #include "helper_cuda.h"
 #include "CudaAllocator.h"
+#include "CudaArray.h"
 #include "ticktock.h"
-
-struct DisableCopy {
-    DisableCopy() = default;
-    DisableCopy(DisableCopy const &) = delete;
-    DisableCopy &operator=(DisableCopy const &) = delete;
-    DisableCopy(DisableCopy &&) = delete;
-    DisableCopy &operator=(DisableCopy &&) = delete;
-};
-
-template <class T>
-class CudaArray {
-    struct BuildArgs {
-        std::array<unsigned int, 3> const dim{};
-        int flags = 0; // or cudaArraySurfaceLoadStore
-    };
-
-    struct Impl : DisableCopy {
-        cudaArray *m_cuArray{};
-        std::array<unsigned int, 3> m_dim{};
-
-        explicit Impl(BuildArgs _args) : m_dim(_args.dim) {
-            cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<T>();  // or cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned)
-            checkCudaErrors(cudaMalloc3DArray(&m_cuArray, &channelDesc, make_cudaExtent(m_dim[0], m_dim[1], m_dim[2]), _args.flags));
-        }
-
-        void assign(T *_data) {
-            cudaMemcpy3DParms copy3DParams{};
-            copy3DParams.srcPtr = make_cudaPitchedPtr((void *)_data, m_dim[0] * sizeof(T), m_dim[1], m_dim[2]);
-            copy3DParams.dstArray = m_cuArray;
-            copy3DParams.extent = make_cudaExtent(m_dim[0], m_dim[1], m_dim[2]);
-            copy3DParams.kind = cudaMemcpyHostToDevice;
-            checkCudaErrors(cudaMemcpy3D(&copy3DParams));
-        }
-
-        ~Impl() {
-            checkCudaErrors(cudaFreeArray(m_cuArray));
-        }
-    };
-
-    std::shared_ptr<Impl> impl;
-
-public:
-    explicit CudaArray(BuildArgs _args) : impl(std::make_shared<Impl>(_args)) {
-    }
-
-    CudaArray &assign(T *_data) const {
-        impl->assign(_data);
-        return *this;
-    }
-
-    operator cudaArray *() const {
-        return impl->m_cuArray;
-    }
-};
-
-template <class T>
-class CudaSurface {
-    struct Impl : DisableCopy {
-        cudaSurfaceObject_t m_cuSuf{};
-        CudaArray<T> m_cuarr;
-
-        explicit Impl(CudaArray<T> _cuarr) : m_cuarr(_cuarr) {
-            cudaResourceDesc resDesc{};
-            resDesc.resType = cudaResourceTypeArray;
-
-            resDesc.res.array.array = m_cuarr;
-            cudaCreateSurfaceObject(&m_cuSuf, &resDesc);
-        }
-
-        ~Impl() {
-            checkCudaErrors(cudaDestroySurfaceObject(m_cuSuf));
-        }
-    };
-
-    std::shared_ptr<Impl> impl;
-
-public:
-    explicit CudaSurface(CudaArray<T> _cuarr) : impl(std::make_shared<Impl>(_cuarr)) {
-    }
-
-    CudaArray<T> &getArray() const {
-        return impl->m_cuarr;
-    }
-
-    operator cudaSurfaceObject_t() const {
-        return impl->m_cuSuf;
-    }
-};
-
-template <class T>
-class CudaTexture {
-    struct Impl : DisableCopy {
-        cudaTextureObject_t m_cuTex{};
-        CudaArray<T> m_cuarr;
-
-        explicit Impl(CudaArray<T> _cuarr) : m_cuarr(_cuarr) {
-            cudaResourceDesc resDesc{};
-            resDesc.resType = cudaResourceTypeArray;
-            resDesc.res.array.array = m_cuarr;
-
-            cudaTextureDesc texDesc{};
-            texDesc.addressMode[0] = cudaAddressModeClamp; // or cudaAddressModeWrap
-            texDesc.addressMode[1] = cudaAddressModeClamp; // or cudaAddressModeWrap
-            texDesc.addressMode[2] = cudaAddressModeClamp; // or cudaAddressModeWrap
-            texDesc.filterMode = cudaFilterModePoint;      // or cudaFilterModeLinear
-            texDesc.readMode = cudaReadModeElementType;    // or cudaReadModeNormalizedFloat
-            texDesc.normalizedCoords = false;              // or true
-
-            checkCudaErrors(cudaCreateTextureObject(&m_cuTex, &resDesc, &texDesc, NULL));
-        }
-
-        ~Impl() {
-            checkCudaErrors(cudaDestroyTextureObject(m_cuTex));
-        }
-    };
-
-    std::shared_ptr<Impl> impl;
-
-public:
-    explicit CudaTexture(CudaArray<T> _cuarr) : impl(std::make_shared<Impl>(_cuarr)) {
-    }
-
-    CudaArray<T> &getArray() const {
-        return impl->m_cuarr;
-    }
-
-    operator cudaTextureObject_t() const {
-        return impl->m_cuTex;
-    }
-};
+#include "ycmcudahelp.h"
 
 __global__ void kernel(cudaSurfaceObject_t out, cudaTextureObject_t in) {
-    int x = 0, y = 0;
+    int x = 0, y = 0, z = 0;
     float fx = 0, fy = 0, fz = 0;
     float value = tex3D<float>(in, fx, fy, fz);
-    value += 1;
-    surf2Dwrite(value, out, x, y);
-    // or cudaBoundaryModeTrap, cudaBoundaryModeClamp
+    value = 3;
+    surf3Dwrite<float>(value, out, x, y, z, cudaBoundaryModeTrap);  // or cudaBoundaryModeZero, cudaBoundaryModeClamp
 }
 
 int main() {
-    CudaSurface<float> out(CudaArray<float>({{1, 1, 1}, cudaArraySurfaceLoadStore}));
-    CudaTexture<float> in(CudaArray<float>({{1, 1, 1}, 0}));
+    unsigned int n = 2;
+    auto out = CudaSurface<float>::make(CudaArray<float>::make({{n, n, n}, cudaArraySurfaceLoadStore}));
+    auto in = CudaTexture<float>::make(CudaArray<float>::make({{n, n, n}, 0}));
+    kernel<<<1, 1>>>(out, in);
+    std::vector<float> arr(n * n * n);
+    out.getArray().copyOut(arr.data());
+    for (int i = 0; i < arr.size(); i++) {
+        printf("%f\n", arr[i]);
+    }
     return 0;
 }

09/01_texture/01/vdbio.cpp

@@ -0,0 +1,2 @@
+#define VDBIO_IMPLEMENTATION
+#include "vdbio.h"

09/01_texture/01/vdbio.h

@@ -0,0 +1,81 @@
+#pragma once
+
+#include <string>
+#include <functional>
+#include <array>
+
+template <class ValT>
+struct _impl_writevdb {
+    std::string const &path;
+    uint32_t sizex, sizey, sizez;
+    std::function<void(ValT *, uint32_t, uint32_t)> sampler;
+
+    void operator()() const;
+};
+
+template <class ValT, class FuncT>
+void writevdb(std::string const &path, uint32_t sizex, uint32_t sizey, uint32_t sizez, FuncT const &func) {
+    _impl_writevdb<ValT>{path, sizex, sizey, sizez, [sizex, &func] (ValT *tmp, uint32_t y, uint32_t z) {
+            for (uint32_t x = 0; x < sizex; x++) {
+                tmp[x] = func(x, y, z);
+            }
+    }}();
+}
+
+#ifdef VDBIO_IMPLEMENTATION
+#include <openvdb/openvdb.h>
+#include <openvdb/tools/Dense.h>
+#include "vdbio.h"
+
+namespace {
+
+template <class T>
+struct vdbtraits {
+    static T convert(T const &val) {
+        return val;
+    }
+};
+
+struct vdbtraits<float> {
+    using grid_type = openvdb::FloatGrid;
+};
+
+struct vdbtraits<std::array<float, 3>> {
+    using grid_type = openvdb::Vec3fGrid;
+
+    template <class T>
+    static typename grid_type::ValueType convert(T const &val) {
+        return {val[0], val[1], val[2]};
+    }
+};
+
+}
+
+template <class ValT>
+void _impl_writevdb<ValT>::operator()() const {
+    auto dummy = [] {
+        if constexpr (std::is_same_v<ValT, std::array<float, 3>>)
+            return std::decay<openvdb::Vec3fGrid>{};
+        else
+            return std::decay<openvdb::FloatGrid>{};
+    }();
+    using GridT = typename vdbtraits<T>::grid_type;
+    openvdb::tools::Dense<typename GridT::ValueType> dens(openvdb::Coord(sizex, sizey, sizez));
+    std::vector<ValT> tmp(sizex);
+    for (uint32_t z = 0; z < sizez; z++) {
+        for (uint32_t y = 0; y < sizey; y++) {
+            sampler(tmp.data(), y, z);
+            for (uint32_t x = 0; x < sizex; x++) {
+                dens.setValue(x, y, z, vdbtraits<T>::convert(tmp[x]));
+            }
+        }
+    }
+    auto grid = GridT::create();
+    typename GridT::ValueType tolerance{0};
+    openvdb::tools::copyFromDense(dens, grid->tree(), tolerance);
+    openvdb::io::File(path).write({grid});
+}
+
+template struct _impl_writevdb<float>;
+template struct _impl_writevdb<std::array<float, 3>>;
+#endif