From 404deb763e6b5255eb6094bafd8dbc7c4912e395 Mon Sep 17 00:00:00 2001 From: WU Qi Date: Sun, 17 Dec 2023 21:14:24 -0800 Subject: [PATCH] Propagate updates from internal branch (1eb506) --- apps/main_app.cpp | 35 +- ovr/common/cuda/cuda_buffer.h | 10 +- ovr/common/cuda/cuda_misc.h | 96 +++- ovr/devices/optix7/device_impl.cpp | 22 +- ovr/devices/optix7/device_impl.h | 2 +- ovr/devices/optix7/render/method_optix.cpp | 89 ++++ ovr/devices/optix7/render/method_optix.cu | 493 ++++++++++++++++++ ovr/devices/optix7/render/method_optix.h | 41 ++ ovr/devices/optix7/render/optix_program.h | 286 ++++++++++ .../optix7/render/raytracing_shaders.h | 106 ++++ ovr/devices/ospray/device_impl.cpp | 219 +++++--- ovr/devices/ospray/device_impl.h | 43 +- ovr/renderer.cpp | 88 ++-- ovr/renderer.h | 57 +- ovr/scene.h | 84 ++- ovr/serializer/serializer_vidi3d.cpp | 16 +- 16 files changed, 1506 insertions(+), 181 deletions(-) create mode 100644 ovr/devices/optix7/render/method_optix.cpp create mode 100644 ovr/devices/optix7/render/method_optix.cu create mode 100644 ovr/devices/optix7/render/method_optix.h create mode 100644 ovr/devices/optix7/render/optix_program.h create mode 100644 ovr/devices/optix7/render/raytracing_shaders.h diff --git a/apps/main_app.cpp b/apps/main_app.cpp index 8807f48..763f3d8 100644 --- a/apps/main_app.cpp +++ b/apps/main_app.cpp @@ -532,13 +532,40 @@ main(int ac, const char** av) ovr::Scene scene; if (ac < 2) { - scene = create_example_scene(); - scene.camera = { /*from*/ vec3f(0.f, 0.f, -1200.f), - /* at */ vec3f(0.f, 0.f, 0.f), - /* up */ vec3f(0.f, 1.f, 0.f) }; + // scene = create_example_scene(); + // scene.camera = { /*from*/ vec3f(0.f, 0.f, -1200.f), + // /* at */ vec3f(0.f, 0.f, 0.f), + // /* up */ vec3f(0.f, 1.f, 0.f) }; + throw std::runtime_error("no scene file specified"); } else { scene = ovr::scene::create_scene(std::string(av[1])); + + // TODO hack for testing isosurface rendering + + // const int32_t volume_raw_id = scene.instances[0].models[0].volume_model.volume_texture; + // scene.instances[0].models[0].volume_model.volume_texture = volume_raw_id; + + // ovr::scene::Texture volume_tfn; + // volume_tfn.type = ovr::scene::Texture::TRANSFER_FUNCTION_TEXTURE; + // volume_tfn.transfer_function.transfer_function = scene.instances[0].models[0].volume_model.transfer_function; + // volume_tfn.transfer_function.volume_texture = volume_raw_id; + // scene.textures.push_back(volume_tfn); + // const int32_t volume_tfn_id = scene.textures.size() - 1; + + // ovr::scene::Material volume_mtl; + // volume_mtl.type = ovr::scene::Material::OBJ_MATERIAL; + // volume_mtl.obj.map_kd = volume_tfn_id; + // scene.materials.push_back(volume_mtl); + // const int32_t volume_mtl_id = scene.materials.size() - 1; + + // ovr::scene::Model model; + // model.type = ovr::scene::Model::GEOMETRIC_MODEL; + // model.geometry_model.geometry.type = ovr::scene::Geometry::ISOSURFACE_GEOMETRY; + // model.geometry_model.geometry.isosurfaces.volume_texture = volume_raw_id; + // model.geometry_model.geometry.isosurfaces.isovalues = { 0.5f }; + // model.geometry_model.mtl = volume_mtl_id; + // scene.instances[0].models[0] = model; } MainWindow::FrameLayer layer; diff --git a/ovr/common/cuda/cuda_buffer.h b/ovr/common/cuda/cuda_buffer.h index 37139e3..4215c38 100644 --- a/ovr/common/cuda/cuda_buffer.h +++ b/ovr/common/cuda/cuda_buffer.h @@ -94,7 +94,7 @@ struct CUDABuffer { CUDA_CHECK(cudaMallocAsync((void**)&d_ptr, sizeInBytes, stream)); - util::total_n_bytes_allocated() += sizeInBytes; + util::tot_nbytes_allocated() += sizeInBytes; #ifdef CUDA_BUFFER_VERBOSE_MEMORY_ALLOCS printf("[mem] CUDABuffer alloc %s\n", util::prettyBytes(sizeInBytes).c_str()); @@ -108,7 +108,7 @@ struct CUDABuffer { { if (owned_data && d_ptr) { CUDA_CHECK(cudaFreeAsync(d_ptr, stream)); - util::total_n_bytes_allocated() -= sizeInBytes; + util::tot_nbytes_allocated() -= sizeInBytes; #ifdef CUDA_BUFFER_VERBOSE_MEMORY_ALLOCS printf("[mem] CUDABuffer free %s\n", util::prettyBytes(sizeInBytes).c_str()); #endif @@ -312,7 +312,7 @@ createCudaArray3D(void* dataPtr, const int3& dims) // allocate 3D CUDA array cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc(); CUDA_CHECK(cudaMalloc3DArray(&dataArr, &channel_desc, make_cudaExtent(dims.x, dims.y, dims.z))); - util::total_n_bytes_allocated() += (size_t)dims.x * dims.y * dims.z * sizeof(Type); + util::tot_nbytes_allocated() += (size_t)dims.x * dims.y * dims.z * sizeof(Type); #ifdef CUDA_BUFFER_VERBOSE_MEMORY_ALLOCS printf("[mem] 3DTex %s\n", util::prettyBytes((size_t)dims.x * dims.y * dims.z * sizeof(Type)).c_str()); #endif @@ -339,7 +339,7 @@ createCudaArray1D(const void* dataPtr, const size_t& size) // Allocate actually a 2D CUDA array of shape N x 1 cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc(); CUDA_CHECK(cudaMallocArray(&dataArr, &channel_desc, size, 1)); - util::total_n_bytes_allocated() += size * sizeof(Type); + util::tot_nbytes_allocated() += size * sizeof(Type); #ifdef CUDA_BUFFER_VERBOSE_MEMORY_ALLOCS printf("[mem] 3DTex %s\n", util::prettyBytes(size * sizeof(Type)).c_str()); #endif @@ -370,7 +370,7 @@ allocateCudaArray1D(const void* dataPtr, const size_t& size) // Allocate a 1D CUDA array of size N cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc(); CUDA_CHECK(cudaMallocArray(&dataArr, &channel_desc, size)); - util::total_n_bytes_allocated() += size * sizeof(Type); + util::tot_nbytes_allocated() += size * sizeof(Type); #ifdef CUDA_BUFFER_VERBOSE_MEMORY_ALLOCS printf("[mem] Array1D %s\n", util::prettyBytes(size * sizeof(Type)).c_str()); #endif diff --git a/ovr/common/cuda/cuda_misc.h b/ovr/common/cuda/cuda_misc.h index 474fbb3..e0303d6 100644 --- a/ovr/common/cuda/cuda_misc.h +++ b/ovr/common/cuda/cuda_misc.h @@ -93,6 +93,9 @@ #define OPTIX_CHECK_NOEXCEPT(call) \ do { \ OptixResult res = call; \ + if (res != OPTIX_SUCCESS) { \ + fprintf(stderr, "OptiX call (%s) failed with %d (line %d)\n", #call, res, __LINE__); \ + } \ } while (0) @@ -400,14 +403,101 @@ static void getUsedGPUMemory(unsigned long long* out) *out = total-free; } -inline std::atomic& total_n_bytes_allocated() +inline std::atomic& tot_nbytes_allocated() noexcept { static std::atomic s{0}; return s; } +inline std::atomic& max_nbytes_allocated() noexcept { static std::atomic s{0}; return s; } +inline void update_max_nbytes_allocated(const size_t& value) noexcept { - static std::atomic s_total_n_bytes_allocated{0}; - return s_total_n_bytes_allocated; + auto& maximum_value = max_nbytes_allocated(); + size_t prev_value = maximum_value; + while(prev_value < value && !maximum_value.compare_exchange_weak(prev_value, value)) {} +} + +} + +inline cudaError_t cudaTrackedFree(void *devPtr, size_t size) { + size_t tot = ::util::tot_nbytes_allocated().fetch_sub(size); +#ifdef VNR_VERBOSE_MEMORY_ALLOCS + printf("[mem] Linear -%s\n", util::prettyBytes(size).c_str()); +#endif + ::util::update_max_nbytes_allocated(tot); + return cudaFree(devPtr); } +template +inline cudaError_t cudaTrackedMalloc(T **devPtr, size_t size) { + size_t tot = ::util::tot_nbytes_allocated().fetch_add(size); +#ifdef VNR_VERBOSE_MEMORY_ALLOCS + printf("[mem] Linear +%s\n", util::prettyBytes(size).c_str()); +#endif + ::util::update_max_nbytes_allocated(tot); + return cudaMalloc(devPtr, size); } +inline cudaError_t cudaTrackedFreeAsync(void *devPtr, size_t size, cudaStream_t stream) { + size_t tot = ::util::tot_nbytes_allocated().fetch_sub(size); +#ifdef VNR_VERBOSE_MEMORY_ALLOCS + printf("[mem] Linear -%s\n", util::prettyBytes(size).c_str()); +#endif + ::util::update_max_nbytes_allocated(tot); + return cudaFreeAsync(devPtr, stream); +} + +template +inline cudaError_t cudaTrackedMallocAsync(T **devPtr, size_t size, cudaStream_t stream) { + size_t tot = ::util::tot_nbytes_allocated().fetch_add(size); +#ifdef VNR_VERBOSE_MEMORY_ALLOCS + printf("[mem] Linear +%s\n", util::prettyBytes(size).c_str()); +#endif + ::util::update_max_nbytes_allocated(tot); + return cudaMallocAsync(devPtr, size, stream); +} + +inline cudaError_t CUDARTAPI cudaTrackedFreeArray(cudaArray_t array) { + // Gets info about the specified cudaArray. + cudaChannelFormatDesc desc; + cudaExtent extent; + unsigned int flags; + CUDA_CHECK(cudaArrayGetInfo(&desc, &extent, &flags, array)); + // Computes the size of the array + size_t size = (desc.x+desc.y+desc.z+desc.w); + if (extent.width == 0 && extent.height == 0 && extent.depth == 0) { + size = 0; + } + else { + if (extent.width) size *= (size_t)extent.width; + if (extent.height) size *= (size_t)extent.height; + if (extent.depth) size *= (size_t)extent.depth; + } + size_t tot = ::util::tot_nbytes_allocated().fetch_sub(size); +#ifdef VNR_VERBOSE_MEMORY_ALLOCS + printf("[mem] Array -%s\n", util::prettyBytes(size).c_str()); +#endif + ::util::update_max_nbytes_allocated(tot); + return cudaFreeArray(array); +} + +inline cudaError_t cudaTrackedMallocArray(cudaArray_t *array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height = 0, unsigned int flags = 0) { + size_t size = (size_t)width * (desc->x+desc->y+desc->z+desc->w); + if (height) size *= (size_t)height; + size_t tot = ::util::tot_nbytes_allocated().fetch_add(size); +#ifdef VNR_VERBOSE_MEMORY_ALLOCS + printf("[mem] Array +%s\n", util::prettyBytes(size).c_str()); +#endif + ::util::update_max_nbytes_allocated(tot); + return cudaMallocArray(array, desc, width, height, flags); +} + +inline cudaError_t cudaTrackedMalloc3DArray(cudaArray_t *array, const struct cudaChannelFormatDesc* desc, struct cudaExtent extent, unsigned int flags = 0) { + size_t size = (size_t)extent.width * (size_t)extent.height * (size_t)extent.depth * (desc->x+desc->y+desc->z+desc->w); + size_t tot = ::util::tot_nbytes_allocated().fetch_add(size); +#ifdef VNR_VERBOSE_MEMORY_ALLOCS + printf("[mem] 3D Array +%s\n", util::prettyBytes(size).c_str()); +#endif + ::util::update_max_nbytes_allocated(tot); + return cudaMalloc3DArray(array, desc, extent, flags); +} + + namespace ovr { namespace misc = ::util; } diff --git a/ovr/devices/optix7/device_impl.cpp b/ovr/devices/optix7/device_impl.cpp index 89a54ec..2eaa7c0 100644 --- a/ovr/devices/optix7/device_impl.cpp +++ b/ovr/devices/optix7/device_impl.cpp @@ -281,34 +281,24 @@ DeviceOptix7::Impl::mapframe(FrameBufferData* fb) } void -DeviceOptix7::Impl::buildScene(Scene& scene) +DeviceOptix7::Impl::buildScene(const Scene& scene) { - assert(scene.instances.size() == 1); - assert(scene.instances[0].models.size() == 1); - assert(scene.instances[0].models[0].type == scene::Model::VOLUMETRIC_MODEL); - assert(scene.instances[0].models[0].volume_model.volume.type == scene::Volume::STRUCTURED_REGULAR_VOLUME); - - // if (!scene.lights.empty()) - // params.light_directional_pos = scene.lights[0].position; - - auto& scene_tfn = scene.instances[0].models[0].volume_model.transfer_function; - auto& scene_volume = scene.instances[0].models[0].volume_model.volume.structured_regular; + auto& scene_volume = parse_single_volume_scene(scene, scene::Volume::STRUCTURED_REGULAR_VOLUME).structured_regular; vec3f scale = scene_volume.grid_spacing * vec3f(scene_volume.data->dims); vec3f translate = scene_volume.grid_origin; - std::cout << "scale " << scale.x << " " << scale.y << " " << scale.z << std::endl; - std::cout << "translate " << translate.x << " " << translate.y << " " << translate.z << std::endl; + // std::cout << "scale " << scale.x << " " << scale.y << " " << scale.z << std::endl; + // std::cout << "translate " << translate.x << " " << translate.y << " " << translate.z << std::endl; // TODO support other parameters // auto v = StructuredRegularVolume(); v.matrix = affine3f::translate(translate) * affine3f::scale(scale); - v.load_from_array3d_scalar(scene_volume.data); - v.set_transfer_function(scene_tfn.color, scene_tfn.opacity, scene_tfn.value_range); v.set_sampling_rate(scene.volume_sampling_rate); - volumes.emplace_back(std::move(v)); + + commit(); } vec3i diff --git a/ovr/devices/optix7/device_impl.h b/ovr/devices/optix7/device_impl.h index fa0d6c9..5ea4cd7 100644 --- a/ovr/devices/optix7/device_impl.h +++ b/ovr/devices/optix7/device_impl.h @@ -50,7 +50,7 @@ struct DeviceOptix7::Impl { void buildSBT(); /*! constructs the scene data */ - void buildScene(Scene& scene); + void buildScene(const Scene& scene); /*! build the instance acceleration structure */ OptixTraversableHandle buildTLAS(); diff --git a/ovr/devices/optix7/render/method_optix.cpp b/ovr/devices/optix7/render/method_optix.cpp new file mode 100644 index 0000000..0cace48 --- /dev/null +++ b/ovr/devices/optix7/render/method_optix.cpp @@ -0,0 +1,89 @@ +//. ======================================================================== // +//. // +//. Copyright 2019-2022 Qi Wu // +//. // +//. Licensed under the MIT License // +//. // +//. ======================================================================== // +//. ======================================================================== // +//. Copyright 2018-2019 Ingo Wald // +//. // +//. Licensed under the Apache License, Version 2.0 (the "License"); // +//. you may not use this file except in compliance with the License. // +//. You may obtain a copy of the License at // +//. // +//. http://www.apache.org/licenses/LICENSE-2.0 // +//. // +//. Unless required by applicable law or agreed to in writing, software // +//. distributed under the License is distributed on an "AS IS" BASIS, // +//. WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +//. See the License for the specific language governing permissions and // +//. limitations under the License. // +//. ======================================================================== // + +#include "method_optix.h" + +namespace vnr { + +extern "C" char embedded_ptx_code_optix[]; + +// ------------------------------------------------------------------ +// +// ------------------------------------------------------------------ + +MethodOptiX::MethodOptiX() : OptixProgram(embedded_ptx_code_optix, RAY_TYPE_COUNT) +{ + OptixProgram::HitGroupShaders s; + + shader_raygen = "__raygen__default"; + + shader_misses.push_back("__miss__radiance"); + shader_misses.push_back("__miss__shadow"); + + /* objects */ + { + OptixProgram::ObjectGroup group; + group.type = VOLUME_STRUCTURED_REGULAR; + + s.shader_CH = "__closesthit__volume_radiance"; + s.shader_AH = "__anyhit__volume_radiance"; + s.shader_IS = "__intersection__volume"; + group.hitgroup.push_back(s); + + s.shader_CH = "__closesthit__volume_shadow"; + s.shader_AH = "__anyhit__volume_shadow"; + s.shader_IS = "__intersection__volume"; + group.hitgroup.push_back(s); + + shader_objects.push_back(group); + } + + params_buffer.resize(sizeof(DefaultUserData), /*stream=*/0); +} + +void +MethodOptiX::render(cudaStream_t stream, const LaunchParams& _params, ShadingMode mode) +{ + DefaultUserData params = _params; + { + params.mode = mode; + params.traversable = ias[0].traversable; + params.geometry_traversable = ias[1].traversable; + } + + params_buffer.upload_async(¶ms, 1, stream); + + OPTIX_CHECK(optixLaunch(/*! pipeline we're launching launch: */ + pipeline.handle, + stream, + /*! parameters and SBT */ + params_buffer.d_pointer(), + params_buffer.sizeInBytes, + &sbt, + /*! dimensions of the launch: */ + params.frame.size.x, + params.frame.size.y, + 1)); +} + +} // namespace ovr diff --git a/ovr/devices/optix7/render/method_optix.cu b/ovr/devices/optix7/render/method_optix.cu new file mode 100644 index 0000000..fa90075 --- /dev/null +++ b/ovr/devices/optix7/render/method_optix.cu @@ -0,0 +1,493 @@ +//. ======================================================================== // +//. // +//. Copyright 2019-2022 Qi Wu // +//. // +//. Licensed under the MIT License // +//. // +//. ======================================================================== // +//. ======================================================================== // +//. Copyright 2018-2019 Ingo Wald // +//. // +//. Licensed under the Apache License, Version 2.0 (the "License"); // +//. you may not use this file except in compliance with the License. // +//. You may obtain a copy of the License at // +//. // +//. http://www.apache.org/licenses/LICENSE-2.0 // +//. // +//. Unless required by applicable law or agreed to in writing, software // +//. distributed under the License is distributed on an "AS IS" BASIS, // +//. WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +//. See the License for the specific language governing permissions and // +//. limitations under the License. // +//. ======================================================================== // + +#include "method_optix.h" +#include "raytracing_shaders.h" + +#include +#include + +#include "dda.h" + +#ifndef ADAPTIVE_SAMPLING +#error "ADAPTIVE_SAMPLING is not defined" +#endif + +namespace vnr { + +/*! launch parameters in constant memory, filled in by optix upon optixLaunch + (this gets filled in from the buffer we pass to optixLaunch) */ +extern "C" __constant__ DefaultUserData optixLaunchParams; + +struct RadiancePayload +{ + float alpha = 0.f; + vec3f color = 0.f; + + void* rng = nullptr; + // float t_max = 0.f; +}; + +struct ShadowPayload +{ + float alpha = 0.f; + + void* rng = nullptr; + // float t_max = 0.f; +}; + +//------------------------------------------------------------------------------ +// helpers +// ------------------------------------------------------------------------------ + +inline __device__ float +sample_size_scaler(const float ss, const float t0, const float t1) { + const int32_t N = (t1-t0) / ss + 1; + return (t1-t0) / N; + // return ss; +} + +template +inline __device__ void +raymarching_iterator(const DeviceVolume& self, + const vec3f& org, const vec3f& dir, + const float tMin, const float tMax, + const float step, const F& body, + bool debug = false) +{ +#if ADAPTIVE_SAMPLING + + const auto& dims = self.macrocell_dims; + const vec3f m_org = org * self.macrocell_spacings_rcp; + const vec3f m_dir = dir * self.macrocell_spacings_rcp; + dda::dda3(m_org, m_dir, tMin, tMax, dims, debug, [&](const vec3i& cell, float t0, float t1) { + // calculate max opacity + float r = opacityUpperBound(self, cell); + if (fabsf(r) <= float_epsilon) return true; // the cell is empty + // estimate a step size + const auto ss = sample_size_scaler(adaptiveSamplingRate(step, r), t0, t1); + // iterate within the interval + vec2f t = vec2f(t0, min(t1, t0 + ss)); + while (t.y > t.x) { + if (!body(t)) return false; + t.x = t.y; + t.y = min(t.x + ss, t1); + } + return true; + }); + +#else + + vec2f t = vec2f(tMin, min(tMax, tMin + step)); + while ((t.y > t.x) && body(t)) { + t.x = t.y; + t.y = min(t.x + step, tMax); + } + +#endif +} + +inline __device__ float +transmittance(void* rng, const vec3f& org, const vec3f& dir) +{ + ShadowPayload shadow; + shadow.rng = rng; + // shadow.t_max = 0.f; + + uint32_t u0, u1; + packPointer(&shadow, u0, u1); + + optixTrace(optixLaunchParams.traversable, + /*org=*/org, /*dir=*/dir, + /*tmin=*/0.f, + /*tmax=*/float_large, + /*time=*/0.0f, + OptixVisibilityMask(255), /* not just volume */ + OPTIX_RAY_FLAG_DISABLE_ANYHIT, + MethodOptiX::SHADOW_RAY_TYPE, // SBT offset + MethodOptiX::RAY_TYPE_COUNT, // SBT stride + MethodOptiX::SHADOW_RAY_TYPE, // miss SBT index + u0, u1); + + return 1.f - shadow.alpha; +} + +//------------------------------------------------------------------------------ +// helpers +// ------------------------------------------------------------------------------ + +inline __device__ void +shadeVolume_radiance(const DeviceVolume& self, + // world space ray + const vec3f _org, + const vec3f _dir, + const float tMin, + const float tMax, + // output + RadiancePayload& payload) +{ + const auto& params = optixLaunchParams; + + if (tMin >= tMax) { + return; + } + + const auto otw = getXfmOTW(); + const auto wto = getXfmWTO(); + const auto wtc = getXfmWorldToCamera(params); + + // to object space + const vec3f org = xfmPoint(wto, _org); + const vec3f dir = xfmVector(wto, _dir); + + // vec3f gradient = 0.f; + vec3f highestContributionOrg = 0.f; + vec3f highestContributionColor = 0.f; + float highestContributionAlpha = 0.f; + + const auto& shadingScale = params.scivis_shading_scale; + const auto& gradientStep = self.grad_step; + + auto rng = (RandomTEA*)payload.rng; + const float jitter = rng->get_floats().x; + + // start marching + raymarching_iterator(self, org, dir, tMin, tMax, self.step, [&](const vec2f& t) + { + // sample data value + const auto p = org + lerp(jitter, t.x, t.y) * dir; // object space position + const auto sampleValue = sampleVolume(self.volume, p); + + // classification + vec3f sampleColor; + float sampleAlpha; + sampleTransferFunction(self.tfn, sampleValue, sampleColor, sampleAlpha); + opacityCorrection(self, t.y - t.x, sampleAlpha); + + // sample gradient + // (from OSPRay) assume that opacity directly correlates to volume scalar + // field, i.e. that "outside" has lower values; because the gradient point + // towards increasing values we need to flip it. + // object space + const vec3f No = -sampleGradient(self.volume, p, sampleValue, gradientStep); + const vec3f Nw = xfmNormal(otw, No); // world space + // const vec3f Nc = xfmNormal(wtc, Nw); // camera space + + const float tr = 1.f - payload.alpha; + + // shade volume + if (params.mode == MethodOptiX::GRADIENT_SHADING) { + const vec3f shadingColor = shade_scivis_light(_dir, Nw, sampleColor, + params.mat_gradient_shading, + params.light_ambient, + params.light_directional_rgb, + params.light_directional_dir); + sampleColor = lerp(shadingScale, sampleColor, shadingColor); + } + + else if (params.mode == MethodOptiX::FULL_SHADOW) { + const float shadow = transmittance(payload.rng, xfmPoint(otw, p), normalize(params.light_directional_dir)); + // const vec3f shadingColor = shade_scivis_light(_dir, Nw, sampleColor, + // params.mat_full_shadow, + // params.light_ambient, + // params.light_directional_rgb, + // params.light_directional_dir); + sampleColor = lerp(shadingScale, sampleColor, sampleColor * shadow); + } + + else if (params.mode == MethodOptiX::SINGLE_SHADE_HEURISTIC) { + // remember point of highest density for deferred shading + if (highestContributionAlpha < (1.f - payload.alpha) * sampleAlpha) { + highestContributionOrg = p; + highestContributionColor = sampleColor; + highestContributionAlpha = (1.f - payload.alpha) * sampleAlpha; + } + // gradient += tr * Nw; // accumulate gradient for SSH + } + + // blending + payload.alpha += tr * sampleAlpha; + payload.color += tr * sampleColor * sampleAlpha; + + return payload.alpha < nearly_one; + }); + + // At a single, point cast rays for contribution. + if (highestContributionAlpha > 0.f) { + const float shadow = transmittance(payload.rng, xfmPoint(otw, highestContributionOrg), normalize(params.light_directional_dir)); + // const vec3f shadingColor = shade_scivis_light(_dir, gradient, highestContributionColor, + // params.mat_single_shade_heuristic, + // params.light_ambient, + // params.light_directional_rgb, + // params.light_directional_dir); + payload.color = lerp(shadingScale, payload.color, highestContributionColor * payload.alpha * shadow); + } +} + +inline __device__ void +shadeVolume_shadow(const DeviceVolume& self, + // world space ray + const vec3f _org, + const vec3f _dir, + const float tMin, + const float tMax, + // performance tuning + const float samplingScale, + // output + ShadowPayload& payload) +{ + if (tMin >= tMax) { + return; + } + + const Array3DScalar& volume = self.volume; + const auto wto = getXfmWTO(); + + // to object space + const vec3f org = xfmPoint(wto, _org); + const vec3f dir = xfmVector(wto, _dir); + + auto rng = (RandomTEA*)payload.rng; + const float jitter = rng->get_floats().x; + + // start marching + raymarching_iterator(self, org, dir, tMin, tMax, samplingScale * self.step, [&](const vec2f& t) + { + // sample data value + const auto p = org + lerp(jitter, t.x, t.y) * dir; // object space position + const auto sampleValue = sampleVolume(volume, p); + + // classification + vec3f sampleColor; + float sampleAlpha; + sampleTransferFunction(self.tfn, sampleValue, sampleColor, sampleAlpha); + opacityCorrection(self, t.y - t.x, sampleAlpha); + + // blending + payload.alpha += (1.f - payload.alpha) * sampleAlpha; + + return payload.alpha < nearly_one; + }); +} + +//------------------------------------------------------------------------------ +// closesthit +// ------------------------------------------------------------------------------ + +extern "C" __global__ void +__closesthit__volume_radiance() +{ + const auto& self = getProgramData(); + auto data = *getPRD(); + + const float t0 = optixGetRayTmax(); + const float t1 = __int_as_float(optixGetAttribute_1()); + + const vec3f rayOrg = optixGetWorldRayOrigin(); + const vec3f rayDir = optixGetWorldRayDirection(); + + // return pre-multiplied color + shadeVolume_radiance(self, rayOrg, rayDir, t0, t1, data); + + // finalize + // data.t_max = t1 + float_epsilon; + *getPRD() = data; +} + +extern "C" __global__ void +__closesthit__volume_shadow() +{ + const auto& self = getProgramData(); + auto data = *getPRD(); + + const float t0 = optixGetRayTmax(); + const float t1 = __int_as_float(optixGetAttribute_1()); + + const vec3f rayOrg = optixGetWorldRayOrigin(); + const vec3f rayDir = optixGetWorldRayDirection(); + + // return alpha + shadeVolume_shadow(self, rayOrg, rayDir, t0, t1, optixLaunchParams.raymarching_shadow_sampling_scale, data); + + // finalize + // data.t_max = t1 + float_epsilon; + *getPRD() = data; +} + +//------------------------------------------------------------------------------ +// anyhit +// ------------------------------------------------------------------------------ + +extern "C" __global__ void +__anyhit__volume_radiance() +{ +} + +extern "C" __global__ void +__anyhit__volume_shadow() +{ +} + +//------------------------------------------------------------------------------ +// miss program that gets called for any ray that did not have a +// valid intersection +// +// as with the anyhit/closest hit programs, in this example we only +// need to have _some_ dummy function to set up a valid SBT +// ------------------------------------------------------------------------------ + +extern "C" __global__ void +__miss__radiance() +{ + // RadiancePayload& payload = *getPRD(); + // payload.t_max = optixGetRayTmax(); +} + +extern "C" __global__ void +__miss__shadow() +{ + // ShadowPayload& payload = *getPRD(); + // payload.t_max = optixGetRayTmax(); +} + +//------------------------------------------------------------------------------ +// ray gen program - the actual rendering happens in here +//------------------------------------------------------------------------------ + +inline __device__ void +render_volume(vec3f org, vec3f dir, void* rng, float& _alpha, vec3f& _color) +{ + RadiancePayload payload; + payload.rng = rng; + + uint32_t u0, u1; + packPointer(&payload, u0, u1); + + /* trace non-volumes */ + // struct + // { + // float alpha = 0.f; + // vec3f color = 0.f; + // } background; + + // optixTrace(optixLaunchParams.traversable, + // /*org=*/org, + // /*dir=*/dir, + // /*tmin=*/0.f, + // /*tmax=*/float_large, + // /*time=*/0.0f, + // OptixVisibilityMask(~VISIBILITY_MASK_VOLUME), /* non-volume */ + // OPTIX_RAY_FLAG_DISABLE_ANYHIT, + // MethodOptiX::RADIANCE_RAY_TYPE, // SBT offset + // MethodOptiX::RAY_TYPE_COUNT, // SBT stride + // MethodOptiX::RADIANCE_RAY_TYPE, // miss SBT index + // u0, + // u1); + + // background.alpha = payload.alpha; + // background.color = payload.color; + + // payload.alpha = 0; + // payload.color = 0; + // payload.t_max = 0.f; + + /* trace volumes */ + optixTrace(optixLaunchParams.traversable, + /*org=*/org, + /*dir=*/dir, + /*tmin=*/0.f, + /*tmax=*/float_large, + /*time=*/0.0f, + OptixVisibilityMask(VISIBILITY_MASK_VOLUME), /* just volume */ + OPTIX_RAY_FLAG_DISABLE_ANYHIT, + MethodOptiX::RADIANCE_RAY_TYPE, // SBT offset + MethodOptiX::RAY_TYPE_COUNT, // SBT stride + MethodOptiX::RADIANCE_RAY_TYPE, // miss SBT index + u0, + u1); + + _alpha = payload.alpha; + _color = payload.color; +} + +extern "C" __global__ void +__raygen__default() +{ + // compute a test pattern based on pixel ID + const int ix = optixGetLaunchIndex().x; + const int iy = optixGetLaunchIndex().y; + + // pixel index + const uint32_t fbIndex = ix + iy * optixLaunchParams.frame.size.x; + + // random number generator + RandomTEA rng_state(optixLaunchParams.frame_index, fbIndex); + + // normalized screen plane position, in [0,1]^2 + const auto& camera = optixLaunchParams.camera; + const vec2f screen(vec2f((float)ix + .5f, (float)iy + .5f) / vec2f(optixLaunchParams.frame.size)); + + // generate ray + vec3f rayOrg = camera.position; + vec3f rayDir = normalize(camera.direction + /* -z axis */ + (screen.x - 0.5f) * camera.horizontal + /* x shift */ + (screen.y - 0.5f) * camera.vertical); /* y shift */ + + // render + float alpha = 0; + vec3f color = 0; + + if (optixLaunchParams.mode == MethodOptiX::DEBUG) { + RadiancePayload payload; + payload.rng = (void*)&rng_state; + + uint32_t u0, u1; + packPointer(&payload, u0, u1); + + optixTrace(optixLaunchParams.geometry_traversable, + /**/ rayOrg, + /**/ rayDir, + /* tmin */ 0.f, + /* tmax */ float_large, + /* rayTime */ 0.0f, + OptixVisibilityMask(255), // only have one volume + OPTIX_RAY_FLAG_DISABLE_ANYHIT, + MethodOptiX::RADIANCE_RAY_TYPE, // SBT offset + MethodOptiX::RAY_TYPE_COUNT, // SBT stride + MethodOptiX::RADIANCE_RAY_TYPE, // miss SBT index + u0, + u1); + + alpha = payload.alpha; + color = payload.color; + } + + else { + render_volume(rayOrg, rayDir, (void*)&rng_state, alpha, color); + } + + // and write to frame buffer ... + writePixelColor(optixLaunchParams, vec4f(color, alpha), fbIndex); +} + +} // namespace ovr diff --git a/ovr/devices/optix7/render/method_optix.h b/ovr/devices/optix7/render/method_optix.h new file mode 100644 index 0000000..d49c0b5 --- /dev/null +++ b/ovr/devices/optix7/render/method_optix.h @@ -0,0 +1,41 @@ +//. ======================================================================== // +//. // +//. Copyright 2019-2022 Qi Wu // +//. // +//. Licensed under the MIT License // +//. // +//. ======================================================================== // + +#pragma once + +#include "optix_program.h" + +#include + +namespace vnr { + +class MethodOptiX : public OptixProgram +{ +public: + enum { RADIANCE_RAY_TYPE = 0, SHADOW_RAY_TYPE, RAY_TYPE_COUNT }; + + enum ShadingMode { NO_SHADING = 0, GRADIENT_SHADING, FULL_SHADOW, SINGLE_SHADE_HEURISTIC, DEBUG }; + + MethodOptiX(); + ~MethodOptiX() { params_buffer.free(0); } + void render(cudaStream_t stream, const LaunchParams& params, ShadingMode mode); + +private: + CUDABuffer params_buffer; +}; + +struct DefaultUserData : LaunchParams +{ + DefaultUserData(const LaunchParams& p) : LaunchParams(p) {} + + MethodOptiX::ShadingMode mode; + OptixTraversableHandle traversable{}; + OptixTraversableHandle geometry_traversable{}; +}; + +}; // namespace ovr diff --git a/ovr/devices/optix7/render/optix_program.h b/ovr/devices/optix7/render/optix_program.h new file mode 100644 index 0000000..f7d697a --- /dev/null +++ b/ovr/devices/optix7/render/optix_program.h @@ -0,0 +1,286 @@ +#pragma once + +#include "../instantvnr_types.h" + +#include +#include +#include + +namespace vnr { + +// ------------------------------------------------------------------ +// Object Definitions +// ------------------------------------------------------------------ + +enum ObjectType { + VOLUME_STRUCTURED_REGULAR, +}; + +#ifndef __NVCC__ +static std::string +object_type_string(ObjectType t) +{ + switch (t) { + case VOLUME_STRUCTURED_REGULAR: return "VOLUME_STRUCTURED_REGULAR"; + default: throw std::runtime_error("unknown object type"); + } +} +#endif + +enum { + VISIBILITY_MASK_GEOMETRY = 0x1, + VISIBILITY_MASK_VOLUME = 0x2, +}; + + +struct OptixProgram +{ +public: + struct HitGroupShaders + { + std::string shader_CH; + std::string shader_AH; + std::string shader_IS; + }; + + struct ObjectGroup + { + ObjectType type; + std::vector hitgroup; + }; + + struct InstanceHandler + { + ObjectType type; // object type + uint32_t idx = 0; // object index + OptixInstance handler; + }; + + OptixProgram(std::string ptx_code, uint32_t num_ray_types) : ptx_code(ptx_code), num_ray_types(num_ray_types) {} + + virtual ~OptixProgram() + { + program.raygen_buffer.free(0); + program.miss_buffer.free(0); + program.hitgroup_buffer.free(0); + } + + void init(OptixDeviceContext, std::map> records, std::vector> blas); + +protected: + /*! creates the module that contains all the programs we are going + to use. in this simple example, we use a single module from a + single .cu file, using a single embedded ptx string */ + void createModule(OptixDeviceContext); + + /*! does all setup for the raygen program(s) we are going to use */ + void createRaygenPrograms(OptixDeviceContext); + + /*! does all setup for the miss program(s) we are going to use */ + void createMissPrograms(OptixDeviceContext); + + /*! does all setup for the hitgroup program(s) we are going to use */ + void createHitgroupPrograms(OptixDeviceContext); + + /*! assembles the full pipeline of all programs */ + void createPipeline(OptixDeviceContext); + + /*! constructs the shader binding table */ + void createSBT(OptixDeviceContext, const std::map>&); + + /*! build the top level acceleration structure */ + void createTLAS(OptixDeviceContext, const std::vector>&); + +protected: + /*! @{ the pipeline we're building */ + struct + { + OptixPipeline handle{}; + OptixPipelineCompileOptions compile_opts{}; + OptixPipelineLinkOptions link_opts{}; + } pipeline; + /*! @} */ + + /*! @{ the module that contains out device programs */ + struct + { + OptixModule handle{}; + OptixModuleCompileOptions compile_opts{}; + } module; + /*! @} */ + + /*! @{ vector of all our program(group)s, and the SBT built around them */ + struct + { + std::vector raygens; + CUDABuffer raygen_buffer; + std::vector misses; + CUDABuffer miss_buffer; + std::vector hitgroups; + CUDABuffer hitgroup_buffer; + } program; + + OptixShaderBindingTable sbt = {}; + /*! @} */ + + std::string shader_raygen; + std::vector shader_misses; + std::vector shader_objects; + + std::map> sbt_offset_table; + + struct IasData + { + std::vector instances; /*! the ISA handlers */ + CUDABuffer instances_buffer; /*! one buffer for all ISAs on GPU */ + CUDABuffer as_buffer; /*! buffer that keeps the (final, compacted) accel structure */ + OptixTraversableHandle traversable; // <- the output + + ~IasData() + { + instances_buffer.free(0); + as_buffer.free(0); + } + }; + std::shared_ptr ias; + + const std::string ptx_code; + const uint32_t num_ray_types; +}; + +struct OptixProgramDenoiser +{ +private: + OptixDenoiser denoiser = nullptr; + CUDABuffer denoiserScratch; + CUDABuffer denoiserState; + + CUDABuffer inputBuffer; + +public: + vec4f* d_ptr() const + { + return (vec4f*)inputBuffer.d_pointer(); + } + + void process(OptixDeviceContext optixContext, cudaStream_t stream, bool accumulate, const uint32_t& frameID, const vec2i& frameSize, vec4f* output) + { + OptixDenoiserParams denoiserParams; + denoiserParams.denoiseAlpha = 1; + denoiserParams.hdrIntensity = (CUdeviceptr)0; + if (accumulate) + denoiserParams.blendFactor = 1.f / (frameID); + else + denoiserParams.blendFactor = 0.0f; + + // ------------------------------------------------------- + OptixImage2D inputLayer; + inputLayer.data = inputBuffer.d_pointer(); + /// Width of the image (in pixels) + inputLayer.width = frameSize.x; + /// Height of the image (in pixels) + inputLayer.height = frameSize.y; + /// Stride between subsequent rows of the image (in bytes). + inputLayer.rowStrideInBytes = frameSize.x * sizeof(float4); + /// Stride between subsequent pixels of the image (in bytes). + /// For now, only 0 or the value that corresponds to a dense packing of pixels (no gaps) is supported. + inputLayer.pixelStrideInBytes = sizeof(float4); + /// Pixel format. + inputLayer.format = OPTIX_PIXEL_FORMAT_FLOAT4; + + // ------------------------------------------------------- + OptixImage2D outputLayer; + outputLayer.data = (CUdeviceptr)output; + /// Width of the image (in pixels) + outputLayer.width = frameSize.x; + /// Height of the image (in pixels) + outputLayer.height = frameSize.y; + /// Stride between subsequent rows of the image (in bytes). + outputLayer.rowStrideInBytes = frameSize.x * sizeof(float4); + /// Stride between subsequent pixels of the image (in bytes). + /// For now, only 0 or the value that corresponds to a dense packing of pixels (no gaps) is supported. + outputLayer.pixelStrideInBytes = sizeof(float4); + /// Pixel format. + outputLayer.format = OPTIX_PIXEL_FORMAT_FLOAT4; + +#if OPTIX_VERSION >= 70300 + OptixDenoiserGuideLayer denoiserGuideLayer = {}; + + OptixDenoiserLayer denoiserLayer = {}; + denoiserLayer.input = inputLayer; + denoiserLayer.output = outputLayer; + + OPTIX_CHECK(optixDenoiserInvoke(denoiser, + /*stream*/ stream, + &denoiserParams, + denoiserState.d_pointer(), + denoiserState.sizeInBytes, + &denoiserGuideLayer, + &denoiserLayer, + 1, + /*inputOffsetX*/ 0, + /*inputOffsetY*/ 0, + denoiserScratch.d_pointer(), + denoiserScratch.sizeInBytes)); +#else + OPTIX_CHECK(optixDenoiserInvoke(denoiser, + /*stream*/ stream, + &denoiserParams, + denoiserState.d_pointer(), + denoiserState.sizeInBytes, + &inputLayer, + 1, + /*inputOffsetX*/ 0, + /*inputOffsetY*/ 0, + &outputLayer, + denoiserScratch.d_pointer(), + denoiserScratch.sizeInBytes)); +#endif + } + + void resize(OptixDeviceContext optixContext, cudaStream_t stream, const vec2i& newSize) + { + if (denoiser) { + OPTIX_CHECK(optixDenoiserDestroy(denoiser)); + }; + + inputBuffer.resize(newSize.x * newSize.y * sizeof(vec4f), stream); + + // ------------------------------------------------------------------ + // create the denoiser: + OptixDenoiserOptions denoiserOptions = {}; + +#if OPTIX_VERSION >= 70300 + OPTIX_CHECK(optixDenoiserCreate(optixContext, OPTIX_DENOISER_MODEL_KIND_LDR, &denoiserOptions, &denoiser)); +#else + denoiserOptions.inputKind = OPTIX_DENOISER_INPUT_RGB; + +#if OPTIX_VERSION < 70100 + // these only exist in 7.0, not 7.1 + denoiserOptions.pixelFormat = OPTIX_PIXEL_FORMAT_FLOAT4; +#endif + OPTIX_CHECK(optixDenoiserCreate(optixContext, &denoiserOptions, &denoiser)); + OPTIX_CHECK(optixDenoiserSetModel(denoiser, OPTIX_DENOISER_MODEL_KIND_LDR, NULL, 0)); +#endif + + // .. then compute and allocate memory resources for the denoiser + OptixDenoiserSizes denoiserReturnSizes; + OPTIX_CHECK(optixDenoiserComputeMemoryResources(denoiser, newSize.x, newSize.y, &denoiserReturnSizes)); + +#if OPTIX_VERSION < 70100 + denoiserScratch.resize(denoiserReturnSizes.recommendedScratchSizeInBytes, stream); +#else + denoiserScratch.resize(std::max(denoiserReturnSizes.withOverlapScratchSizeInBytes, denoiserReturnSizes.withoutOverlapScratchSizeInBytes), stream); +#endif + denoiserState.resize(denoiserReturnSizes.stateSizeInBytes, stream); + + // ------------------------------------------------------------------ + OPTIX_CHECK(optixDenoiserSetup(denoiser, stream, newSize.x, newSize.y, + denoiserState.d_pointer(), + denoiserState.sizeInBytes, + denoiserScratch.d_pointer(), + denoiserScratch.sizeInBytes)); + } +}; + +} // namespace vnr diff --git a/ovr/devices/optix7/render/raytracing_shaders.h b/ovr/devices/optix7/render/raytracing_shaders.h new file mode 100644 index 0000000..5571ee0 --- /dev/null +++ b/ovr/devices/optix7/render/raytracing_shaders.h @@ -0,0 +1,106 @@ +//. ======================================================================== // +//. // +//. Copyright 2019-2022 Qi Wu // +//. // +//. Licensed under the MIT License // +//. // +//. ======================================================================== // +#pragma once + +#include "raytracing.h" +#include + +namespace vnr { + +static __forceinline__ __device__ void* +unpackPointer(uint32_t i0, uint32_t i1) +{ + const auto uptr = static_cast(i0) << 32U | i1; + void* ptr = reinterpret_cast(uptr); + return ptr; +} + +static __forceinline__ __device__ void +packPointer(void* ptr, uint32_t& i0, uint32_t& i1) +{ + const auto uptr = reinterpret_cast(ptr); + i0 = uptr >> 32U; + i1 = uptr & 0x00000000ffffffff; +} + +template +static __forceinline__ __device__ T* +getPRD() +{ + const uint32_t u0 = optixGetPayload_0(); + const uint32_t u1 = optixGetPayload_1(); + return reinterpret_cast(unpackPointer(u0, u1)); +} + +template +__device__ inline const T& +getProgramData() +{ + return **((const T**)optixGetSbtDataPointer()); +} + +inline __device__ affine3f +getXfmWTO() +{ + float mat[12]; /* 3x4 row major */ + optixGetWorldToObjectTransformMatrix(mat); + affine3f xfm; + xfm.l.vx = vec3f(mat[0], mat[4], mat[8]); + xfm.l.vy = vec3f(mat[1], mat[5], mat[9]); + xfm.l.vz = vec3f(mat[2], mat[6], mat[10]); + xfm.p = vec3f(mat[3], mat[7], mat[11]); + return xfm; +} + +inline __device__ affine3f +getXfmOTW() +{ + float mat[12]; /* 3x4 row major */ + optixGetObjectToWorldTransformMatrix(mat); + affine3f xfm; + xfm.l.vx = vec3f(mat[0], mat[4], mat[8]); + xfm.l.vy = vec3f(mat[1], mat[5], mat[9]); + xfm.l.vz = vec3f(mat[2], mat[6], mat[10]); + xfm.p = vec3f(mat[3], mat[7], mat[11]); + return xfm; +} + +inline vec2f __device__ +projectToScreen(const vec3f p, const DeviceCamera& camera) +{ + vec3f wsvec = p - camera.position; + vec2f screen; + const float r = length(camera.horizontal); + const float t = length(camera.vertical); + screen.x = dot(wsvec, normalize(camera.horizontal)) / r; + screen.y = dot(wsvec, normalize(camera.vertical)) / t; + return screen + 0.5f; +} + +//------------------------------------------------------------------------------ +// intersection program that computes customized intersections for an AABB +// ------------------------------------------------------------------------------ + +extern "C" __global__ void +__intersection__volume() +{ + const vec3f org = optixGetObjectRayOrigin(); + const vec3f dir = optixGetObjectRayDirection(); + + float t0 = optixGetRayTmin(); + float t1 = optixGetRayTmax(); + + const auto& self = getProgramData(); + + if (intersectVolume(t0, t1, org, dir, self)) { + optixReportIntersection(t0, 0, /* user defined attributes, for now set to 0 */ + __float_as_int(t0), __float_as_int(t1)); + } +} + +} diff --git a/ovr/devices/ospray/device_impl.cpp b/ovr/devices/ospray/device_impl.cpp index e7efcb6..1691aec 100644 --- a/ovr/devices/ospray/device_impl.cpp +++ b/ovr/devices/ospray/device_impl.cpp @@ -70,8 +70,7 @@ namespace { template void -ospSetVectorAsData(OSPObject o, const char* n, OSPDataType type, const std::vector& array) -{ +ospSetVectorAsData(OSPObject o, const char* n, OSPDataType type, const std::vector& array) { if (!array.empty()) { OSPData shared = ospNewSharedData(array.data(), type, array.size()); OSPData data = ospNewData(type, array.size()); @@ -84,8 +83,7 @@ ospSetVectorAsData(OSPObject o, const char* n, OSPDataType type, const std::vect template OSPData -ospNewData2D(OSPDataType type, const std::vector& array) -{ +ospNewData2D(OSPDataType type, const std::vector& array) { OSPData shared = ospNewSharedData(array.data(), type, array.size()); OSPData data = ospNewData(type, array.size()); ospCopyData(shared, data); @@ -93,8 +91,7 @@ ospNewData2D(OSPDataType type, const std::vector& array) } inline OSPDataType -to_ospray_data_type(ValueType type) -{ +to_ospray_data_type(ValueType type) { switch (type) { case VALUE_TYPE_UINT8: return OSP_UCHAR; case VALUE_TYPE_INT8: return OSP_CHAR; @@ -104,33 +101,26 @@ to_ospray_data_type(ValueType type) case VALUE_TYPE_INT32: return OSP_INT; case VALUE_TYPE_FLOAT: return OSP_FLOAT; case VALUE_TYPE_DOUBLE: return OSP_DOUBLE; - case VALUE_TYPE_FLOAT2: return OSP_VEC2F; case VALUE_TYPE_FLOAT3: return OSP_VEC3F; case VALUE_TYPE_FLOAT4: return OSP_VEC4F; - default: throw std::runtime_error("unknown type encountered"); } } inline OSPData -create_ospray_array1d_scalar(array_1d_scalar_t input) -{ +create_ospray_array1d_scalar(array_1d_scalar_t input) { return ospNewSharedData1D(input->data(), to_ospray_data_type(input->type), input->dims.v); } inline OSPData -create_ospray_array1d_float3(array_1d_float4_t input) -{ - if (input->type != value_type()) - throw std::runtime_error("type mismatch"); - +create_ospray_array1d_float3(array_1d_float4_t input) { + if (input->type != value_type()) throw std::runtime_error("type mismatch"); return ospNewSharedData1DStride(input->data(), OSP_VEC3F, input->dims.v, sizeof(vec4f)); } inline OSPData -create_ospray_array3d_scalar(array_3d_scalar_t input) -{ +create_ospray_array3d_scalar(array_3d_scalar_t input) { return ospNewSharedData3D(input->data(), to_ospray_data_type(input->type), /**/ input->dims.x, input->dims.y, input->dims.z); } @@ -141,8 +131,7 @@ create_ospray_array3d_scalar(array_3d_scalar_t input) // ------------------------------------------------------------------ OSPVolume -DeviceOSPRay::Impl::create_ospray_volume(scene::Volume::VolumeStructuredRegular handler) -{ +DeviceOSPRay::Impl::create_ospray_volume(scene::Volume::VolumeStructuredRegular handler) { auto data = create_ospray_array3d_scalar(handler.data); OSPVolume volume = ospNewVolume("structuredRegular"); ospSetParam(volume, "gridOrigin", OSP_VEC3F, &handler.grid_origin); @@ -155,8 +144,7 @@ DeviceOSPRay::Impl::create_ospray_volume(scene::Volume::VolumeStructuredRegular } OSPVolume -DeviceOSPRay::Impl::create_ospray_volume(scene::Volume handler) -{ +DeviceOSPRay::Impl::create_ospray_volume(scene::Volume handler) { switch (handler.type) { case scene::Volume::STRUCTURED_REGULAR_VOLUME: return create_ospray_volume(handler.structured_regular); default: throw std::runtime_error("unknown volume type"); @@ -164,22 +152,21 @@ DeviceOSPRay::Impl::create_ospray_volume(scene::Volume handler) } OSPTransferFunction -DeviceOSPRay::Impl::create_ospray_transfer_function(scene::TransferFunction handler) -{ +DeviceOSPRay::Impl::create_ospray_transfer_function(scene::TransferFunction handler) { OSPTransferFunction tfn = ospNewTransferFunction("piecewiseLinear"); ospSetParam(tfn, "valueRange", OSP_VEC2F, &handler.value_range); ospSetObject(tfn, "color", create_ospray_array1d_float3(handler.color)); ospSetObject(tfn, "opacity", create_ospray_array1d_scalar(handler.opacity)); ospCommit(tfn); + ospray.tfns.push_back(tfn); return tfn; } OSPGeometry -DeviceOSPRay::Impl::create_ospray_geometry(scene::Geometry::GeometryTriangles handler) -{ - OSPGeometry mesh = ospNewGeometry("mesh"); +DeviceOSPRay::Impl::create_ospray_geometry(scene::Geometry::GeometryTriangles handler) { // TODO finish the implementation // + OSPGeometry mesh = ospNewGeometry("mesh"); assert(handler.position->type == ovr::VALUE_TYPE_FLOAT3); auto position = ospNewSharedData1D(handler.position->data(), OSP_VEC3F, handler.position->dims.v); @@ -250,62 +237,128 @@ DeviceOSPRay::Impl::create_ospray_geometry(scene::Geometry::GeometryTriangles ha } ospCommit(mesh); - ospRelease(position); ospRelease(index); - if (faces_texcoord) ospRelease(faces_texcoord); - if (faces_normal) ospRelease(faces_normal); - if (faces_color) ospRelease(faces_color); - + if (faces_normal) ospRelease(faces_normal); + if (faces_color) ospRelease(faces_color); if (verts_texcoord) ospRelease(verts_texcoord); - if (verts_normal) ospRelease(verts_normal); - if (verts_color) ospRelease(verts_color); - + if (verts_normal) ospRelease(verts_normal); + if (verts_color) ospRelease(verts_color); return mesh; } OSPGeometry -DeviceOSPRay::Impl::create_ospray_geometry(scene::Geometry handler) -{ +DeviceOSPRay::Impl::create_ospray_geometry(scene::Geometry::GeometryIsosurfaces handler) { + OSPGeometry geom = ospNewGeometry("isosurface"); + OSPVolume volume = ospray.get_volume(handler.volume_texture); + ospSetVectorAsData(geom, "isovalue", OSP_FLOAT, handler.isovalues); + ospSetObject(geom, "volume", volume); + ospCommit(geom); + return geom; +} + +OSPGeometry +DeviceOSPRay::Impl::create_ospray_geometry(scene::Geometry handler) { switch (handler.type) { - case scene::Geometry::TRIANGLES_GEOMETRY: return create_ospray_geometry(handler.triangles); + case scene::Geometry::TRIANGLES_GEOMETRY: return create_ospray_geometry(handler.triangles); + case scene::Geometry::ISOSURFACE_GEOMETRY: return create_ospray_geometry(handler.isosurfaces); default: throw std::runtime_error("unknown geometry type"); } } -OSPVolumetricModel -DeviceOSPRay::Impl::create_ospray_volumetric_model(scene::Model::VolumetricModel handler) -{ - auto volume = create_ospray_volume(handler.volume); - auto tfn = create_ospray_transfer_function(handler.transfer_function); +// ------------------------------------------------------------------ +// +// ------------------------------------------------------------------ +OSPTexture +DeviceOSPRay::Impl::create_ospray_texture(scene::Texture::TransferFunctionTexture handler) { + OSPTexture texture = ospNewTexture("volume"); + OSPTransferFunction tfn = create_ospray_transfer_function(handler.transfer_function); + OSPVolume volume = ospray.get_volume(handler.volume_texture); + ospSetObject(texture, "transferFunction", tfn); + ospSetObject(texture, "volume", volume); + ospCommit(texture); + return texture; +} + +OSPTexOrVol +DeviceOSPRay::Impl::create_ospray_texture(scene::Texture handler) { + using namespace scene; + OSPTexOrVol ret; + switch (handler.type) { + case Texture::VOLUME_TEXTURE: ret.vol = create_ospray_volume(handler.volume.volume); break; + case Texture::TRANSFER_FUNCTION_TEXTURE: ret.tex = create_ospray_texture(handler.transfer_function); break; + default: throw std::runtime_error("unknown texture type: " + std::to_string(handler.type)); + } + return ret; +} + +// ------------------------------------------------------------------ +// +// ------------------------------------------------------------------ + +OSPMaterial +DeviceOSPRay::Impl::create_ospray_material(scene::Material::ObjMaterial handler) { + OSPMaterial mtl = ospNewMaterial(NULL, "obj"); + ospSetVec3f(mtl, "kd", handler.kd.x, handler.kd.y, handler.kd.z); + ospSetVec3f(mtl, "ks", handler.ks.x, handler.ks.y, handler.ks.z); + ospSetFloat(mtl, "ns", handler.ns); + ospSetFloat(mtl, "d", handler.d); + ospSetVec3f(mtl, "tf", handler.tf.x, handler.tf.y, handler.tf.z); + if (handler.map_kd != -1) { + ospSetObject(mtl, "map_kd", ospray.get_texture(handler.map_kd)); + } + if (handler.map_bump != -1) { + ospSetObject(mtl, "map_bump", ospray.get_texture(handler.map_bump)); + } + ospCommit(mtl); + return mtl; +} + +OSPMaterial +DeviceOSPRay::Impl::create_ospray_material(scene::Material handler) { + using namespace scene; + switch (handler.type) { + case Material::OBJ_MATERIAL: return create_ospray_material(handler.obj); + default: throw std::runtime_error("unknown material type"); + } +} + +// ------------------------------------------------------------------ +// +// ------------------------------------------------------------------ + +OSPVolumetricModel +DeviceOSPRay::Impl::create_ospray_volumetric_model(scene::Model::VolumetricModel handler) { + OSPTransferFunction tfn = create_ospray_transfer_function(handler.transfer_function); + OSPVolume volume = ospray.get_volume(handler.volume_texture); OSPVolumetricModel model = ospNewVolumetricModel(volume); ospSetObject(model, "transferFunction", tfn); ospSetFloat(model, "gradientShadingScale", 1.f); ospCommit(model); ospRelease(volume); - - ospray.tfns.push_back(tfn); - return model; } OSPGeometricModel -DeviceOSPRay::Impl::create_ospray_geometric_model(scene::Model::GeometricModel handler) -{ +DeviceOSPRay::Impl::create_ospray_geometric_model(scene::Model::GeometricModel handler) { auto geometry = create_ospray_geometry(handler.geometry); OSPGeometricModel model = ospNewGeometricModel(geometry); - OSPMaterial mtl = ospNewMaterial(NULL, "obj"); - ospSetObject(model, "material", mtl); + if (handler.mtl == -1) { + OSPMaterial mtl = ospNewMaterial(NULL, "obj"); + ospSetObject(model, "material", mtl); + } + else { + ospSetObject(model, "material", ospray.materials[handler.mtl]); + } ospCommit(model); ospRelease(geometry); return model; } OSPInstance -DeviceOSPRay::Impl::create_ospray_instance(scene::Instance handler) -{ +DeviceOSPRay::Impl::create_ospray_instance(scene::Instance handler) { std::vector volume_group; std::vector geometry_group; for (auto m : handler.models) { @@ -325,12 +378,8 @@ DeviceOSPRay::Impl::create_ospray_instance(scene::Instance handler) ospSetVectorAsData(group, "volume", OSP_VOLUMETRIC_MODEL, volume_group); ospCommit(group); - for (auto m : geometry_group) { - ospRelease(m); - } - for (auto m : volume_group) { - ospRelease(m); - } + for (auto m : geometry_group) { ospRelease(m); } + for (auto m : volume_group) { ospRelease(m); } // put the group into an instance (give the group a world transform) OSPInstance instance = ospNewInstance(group); @@ -342,8 +391,7 @@ DeviceOSPRay::Impl::create_ospray_instance(scene::Instance handler) return instance; } -DeviceOSPRay::Impl::~Impl() -{ +DeviceOSPRay::Impl::~Impl() { for (auto tfn : ospray.tfns) { ospRelease(tfn); } @@ -361,8 +409,7 @@ DeviceOSPRay::Impl::~Impl() } void -DeviceOSPRay::Impl::init(int argc, const char** argv, DeviceOSPRay* p) -{ +DeviceOSPRay::Impl::init(int argc, const char** argv, DeviceOSPRay* p) { if (parent) { if (parent != p) throw std::runtime_error("[ospray] a different parent is provided"); } @@ -408,8 +455,7 @@ DeviceOSPRay::Impl::init(int argc, const char** argv, DeviceOSPRay* p) } void -DeviceOSPRay::Impl::commit_renderer() -{ +DeviceOSPRay::Impl::commit_renderer() { auto& scene = parent->current_scene; if (parent->params.path_tracing.update() || !ospray.renderer) { @@ -455,8 +501,7 @@ DeviceOSPRay::Impl::commit_renderer() } void -DeviceOSPRay::Impl::commit_framebuffer() -{ +DeviceOSPRay::Impl::commit_framebuffer() { bool recreate = !ospray.framebuffer; if (parent->params.fbsize.update()) { @@ -494,8 +539,7 @@ DeviceOSPRay::Impl::commit_framebuffer() } void -DeviceOSPRay::Impl::commit_camera() -{ +DeviceOSPRay::Impl::commit_camera() { if (parent->params.camera.update() || !ospray.camera) { const Camera& camera = parent->params.camera.ref(); @@ -536,8 +580,7 @@ DeviceOSPRay::Impl::commit_camera() } void -DeviceOSPRay::Impl::commit_transfer_function() -{ +DeviceOSPRay::Impl::commit_transfer_function() { if (parent->params.tfn.update()) { framebuffer_should_reset_accum = true; @@ -562,6 +605,7 @@ DeviceOSPRay::Impl::commit_transfer_function() OSPData color_data = ospNewData2D(OSP_VEC3F, tfn_color); OSPData opacity_data = ospNewData2D(OSP_FLOAT, tfn_opacity); + // TODO: can only support one transfer function for now, limited by the framework for (auto& tfn : ospray.tfns) { ospSetParam(tfn, "valueRange", OSP_VEC2F, &tfn_range); ospSetObject(tfn, "color", color_data); @@ -572,11 +616,22 @@ DeviceOSPRay::Impl::commit_transfer_function() } void -DeviceOSPRay::Impl::build_scene() -{ +DeviceOSPRay::Impl::build_scene() { // put the instance in the world const auto& scene = parent->current_scene; + // create all standalone volumes & textures + ospray.texorvols.reserve(scene.textures.size()); + for (auto t : scene.textures) { + ospray.texorvols.push_back(create_ospray_texture(t)); + } + + // create all materials + ospray.materials.reserve(scene.materials.size()); + for (auto m : scene.materials) { + ospray.materials.push_back(create_ospray_material(m)); + } + // create all ospray instances std::vector instances; for (auto i : scene.instances) { @@ -594,7 +649,11 @@ DeviceOSPRay::Impl::build_scene() } else if (l.type == scene::Light::DIRECTIONAL) { light = ospNewLight("distant"); - ospSetVec3f(light, "direction", l.directional.direction.x, l.directional.direction.y, l.directional.direction.z); + ospSetVec3f(light, "direction", + l.directional.direction.x, + l.directional.direction.y, + l.directional.direction.z + ); } else { throw std::runtime_error("OSPray: unknown light type"); @@ -644,14 +703,12 @@ DeviceOSPRay::Impl::build_scene() } void -DeviceOSPRay::Impl::swap() -{ +DeviceOSPRay::Impl::swap() { framebuffer_index = (framebuffer_index + 1) % 2; } void -DeviceOSPRay::Impl::commit() -{ +DeviceOSPRay::Impl::commit() { if (parent->params.focus_center.update()) { framebuffer_should_reset_accum = true; } @@ -667,6 +724,10 @@ DeviceOSPRay::Impl::commit() const auto& scene = parent->current_scene; commit_transfer_function(); + for (auto& m : ospray.materials) { + ospCommit(m); + } + commit_framebuffer(); commit_renderer(); @@ -684,8 +745,7 @@ DeviceOSPRay::Impl::commit() } void -DeviceOSPRay::Impl::render() -{ +DeviceOSPRay::Impl::render() { frame_index++; if (parent->params.sparse_sampling.ref()) { @@ -721,8 +781,7 @@ DeviceOSPRay::Impl::render() } void -DeviceOSPRay::Impl::mapframe(FrameBufferData* fb) -{ +DeviceOSPRay::Impl::mapframe(FrameBufferData* fb) { const size_t num_bytes = framebuffer_size_latest.long_product(); fb->rgba->set_data((void*)framebuffer_rgba_ptr, num_bytes * sizeof(vec4f), CrossDeviceBuffer::DEVICE_CPU); } diff --git a/ovr/devices/ospray/device_impl.h b/ovr/devices/ospray/device_impl.h index 9731bab..c5d02f9 100644 --- a/ovr/devices/ospray/device_impl.h +++ b/ovr/devices/ospray/device_impl.h @@ -27,6 +27,12 @@ namespace ovr::ospray { +// In OVR, we consider volumes as special textures +struct OSPTexOrVol { + OSPVolume vol = nullptr; + OSPTexture tex = nullptr; +}; + struct DeviceOSPRay::Impl { DeviceOSPRay* parent{ nullptr }; @@ -53,13 +59,24 @@ struct DeviceOSPRay::Impl { // scene building functions OSPVolume create_ospray_volume(scene::Volume::VolumeStructuredRegular handler); OSPVolume create_ospray_volume(scene::Volume handler); - OSPTransferFunction create_ospray_transfer_function(scene::TransferFunction handler); + OSPGeometry create_ospray_geometry(scene::Geometry::GeometryTriangles handler); + OSPGeometry create_ospray_geometry(scene::Geometry::GeometryIsosurfaces handler); OSPGeometry create_ospray_geometry(scene::Geometry handler); + + OSPTexture create_ospray_texture(scene::Texture::TransferFunctionTexture handler); + + OSPMaterial create_ospray_material(scene::Material::ObjMaterial handler); + OSPMaterial create_ospray_material(scene::Material handler); + + OSPTransferFunction create_ospray_transfer_function(scene::TransferFunction handler); OSPVolumetricModel create_ospray_volumetric_model(scene::Model::VolumetricModel handler); OSPGeometricModel create_ospray_geometric_model(scene::Model::GeometricModel handler); OSPInstance create_ospray_instance(scene::Instance handler); + // our extension + OSPTexOrVol create_ospray_texture(scene::Texture handler); + protected: struct { OSPCamera camera{ 0 }; @@ -67,7 +84,31 @@ struct DeviceOSPRay::Impl { OSPWorld world{ 0 }; OSPRenderer renderer{ 0 }; std::vector tfns; + std::vector texorvols; + std::vector materials; OSPData sparse_samples; + + OSPVolume get_volume(int32_t idx) { + if (idx < 0 || idx >= texorvols.size()) { + throw std::runtime_error("invalid texture index: " + std::to_string(idx) + " / " + std::to_string(texorvols.size())); + } + auto ret = texorvols[idx]; + if (ret.vol == nullptr) { + throw std::runtime_error("texture " + std::to_string(idx) + " is not a volume texture"); + } + return ret.vol; + } + + OSPTexture get_texture(int32_t idx) { + if (idx < 0 || idx >= texorvols.size()) { + throw std::runtime_error("invalid texture index: " + std::to_string(idx) + " / " + std::to_string(texorvols.size())); + } + auto ret = texorvols[idx]; + if (ret.tex == nullptr) { + throw std::runtime_error("texture " + std::to_string(idx) + " is not a texture"); + } + return ret.tex; + } } ospray; uint32_t framebuffer_channels = OSP_FB_COLOR | OSP_FB_VARIANCE /*| OSP_FB_NORMAL*/ /*| OSP_FB_ACCUM*/; diff --git a/ovr/renderer.cpp b/ovr/renderer.cpp index 42e4080..3b150f9 100644 --- a/ovr/renderer.cpp +++ b/ovr/renderer.cpp @@ -60,47 +60,47 @@ create_renderer(std::string name) throw std::runtime_error("unknown device name: " + name); } -Scene -create_example_scene() -{ - static std::random_device rng; // Will be used to obtain a seed for the random number engine - static std::mt19937 generator(rng()); // Standard mersenne_twister_engine seeded with rng() - static std::uniform_int_distribution<> index(0, (int)(colormap::name.size() - 1ULL)); - - scene::TransferFunction tfn; - - // /* S3D data, range = (-7.6862547926737625e-19, 0.00099257915280759335) */ - // scene::Volume volume; - // volume.type = scene::Volume::STRUCTURED_REGULAR_VOLUME; - // volume.structured_regular.data = - // CreateArray3DScalarFromFile("data/3.900E-04_H2O2.raw", vec3i(600, 750, 500), VALUE_TYPE_FLOAT, 0, false); - // volume.structured_regular.grid_origin = -vec3i(600, 750, 500) / 2.f; - // tfn.value_range = vec2f(0.f, 0.0002f); - // tfn.color = CreateColorMap(colormap::name[index(generator)]); - // tfn.opacity = CreateArray1DScalar(std::vector{ 0.f, 1.f }); - - /* 0.0070086f, 12.1394f */ - scene::Volume volume; - volume.type = scene::Volume::STRUCTURED_REGULAR_VOLUME; - volume.structured_regular.data = - CreateArray3DScalarFromFile(std::string("data/vorts1.data"), vec3i(128, 128, 128), VALUE_TYPE_FLOAT, 0, false); - volume.structured_regular.grid_origin = -vec3i(128, 128, 128) * 2.f; - volume.structured_regular.grid_spacing = 4.f; - tfn.value_range = vec2f(0.0, 12.0); - tfn.color = CreateColorMap("sequential2/summer"); - tfn.opacity = CreateArray1DScalar(std::vector{ 0.f, 0.f, 0.01f, 0.2f, 0.01f, 0.f, 0.f }); - - scene::Model model; - model.type = scene::Model::VOLUMETRIC_MODEL; - model.volume_model.volume = volume; - model.volume_model.transfer_function = tfn; - - scene::Instance instance; - instance.models.push_back(model); - instance.transform = affine3f::translate(vec3f(0)); - - scene::Scene scene; - scene.instances.push_back(instance); - - return scene; -} +// Scene +// create_example_scene() +// { +// static std::random_device rng; // Will be used to obtain a seed for the random number engine +// static std::mt19937 generator(rng()); // Standard mersenne_twister_engine seeded with rng() +// static std::uniform_int_distribution<> index(0, (int)(colormap::name.size() - 1ULL)); +// +// scene::TransferFunction tfn; +// +// // /* S3D data, range = (-7.6862547926737625e-19, 0.00099257915280759335) */ +// // scene::Volume volume; +// // volume.type = scene::Volume::STRUCTURED_REGULAR_VOLUME; +// // volume.structured_regular.data = +// // CreateArray3DScalarFromFile("data/3.900E-04_H2O2.raw", vec3i(600, 750, 500), VALUE_TYPE_FLOAT, 0, false); +// // volume.structured_regular.grid_origin = -vec3i(600, 750, 500) / 2.f; +// // tfn.value_range = vec2f(0.f, 0.0002f); +// // tfn.color = CreateColorMap(colormap::name[index(generator)]); +// // tfn.opacity = CreateArray1DScalar(std::vector{ 0.f, 1.f }); +// +// /* 0.0070086f, 12.1394f */ +// scene::Volume volume; +// volume.type = scene::Volume::STRUCTURED_REGULAR_VOLUME; +// volume.structured_regular.data = +// CreateArray3DScalarFromFile(std::string("data/vorts1.data"), vec3i(128, 128, 128), VALUE_TYPE_FLOAT, 0, false); +// volume.structured_regular.grid_origin = -vec3i(128, 128, 128) * 2.f; +// volume.structured_regular.grid_spacing = 4.f; +// tfn.value_range = vec2f(0.0, 12.0); +// tfn.color = CreateColorMap("sequential2/summer"); +// tfn.opacity = CreateArray1DScalar(std::vector{ 0.f, 0.f, 0.01f, 0.2f, 0.01f, 0.f, 0.f }); +// +// scene::Model model; +// model.type = scene::Model::VOLUMETRIC_MODEL; +// model.volume_model.volume = volume; +// model.volume_model.transfer_function = tfn; +// +// scene::Instance instance; +// instance.models.push_back(model); +// instance.transform = affine3f::translate(vec3f(0)); +// +// scene::Scene scene; +// scene.instances.push_back(instance); +// +// return scene; +// } diff --git a/ovr/renderer.h b/ovr/renderer.h index 5294b95..a4d436e 100644 --- a/ovr/renderer.h +++ b/ovr/renderer.h @@ -49,6 +49,31 @@ namespace ovr { using vidi::TransactionalValue; +using scene::Camera; + +inline int +count_tfn(const scene::Scene& scene, scene::TransferFunction& scene_tfn) +{ + int count = 0; + // There are two places transfer function can be stored: + // 1) in a volume model + for (const auto& instance : scene.instances) { + for (const auto& model : instance.models) { + if (model.type == scene::Model::VOLUMETRIC_MODEL) { + scene_tfn = model.volume_model.transfer_function; + count++; + } + } + } + // 2) in a texture + for (const auto& texture : scene.textures) { + if (texture.type == scene::Texture::TRANSFER_FUNCTION_TEXTURE) { + scene_tfn = texture.transfer_function.transfer_function; + count++; + } + } + return count; +} /*! a sample OptiX-7 renderer that demonstrates how to set up context, module, programs, pipeline, SBT, etc, and perform a @@ -62,13 +87,10 @@ struct MainRenderer { double render_time; struct FrameBufferData { - // vec4f* rgba{}; - // vec3f* grad{}; std::shared_ptr rgba; std::shared_ptr grad; - FrameBufferData() - { + FrameBufferData() { rgba = std::make_shared(); grad = std::make_shared(); } @@ -180,6 +202,11 @@ struct MainRenderer { params.volume_sampling_rate = volume_sampling_rate; } + void set_volume_density_scale(float volume_density_scale) + { + params.volume_density_scale = volume_density_scale; + } + void set_mat_ambient(float ambient) { params.ambient = ambient; @@ -221,7 +248,7 @@ struct MainRenderer { } protected: - void set_scene(Scene scene); + void set_scene(const Scene& scene); virtual void init(int argc, const char** argv) = 0; protected: @@ -231,6 +258,7 @@ struct MainRenderer { TransactionalValue sample_per_pixel; TransactionalValue volume_sampling_rate; + TransactionalValue volume_density_scale; TransactionalValue ambient; TransactionalValue diffuse; @@ -246,9 +274,11 @@ struct MainRenderer { TransactionalValue focus_scale; TransactionalValue base_noise; TransactionalValue add_lights; + TransactionalValue sparse_sampling; TransactionalValue path_tracing; TransactionalValue photonmapping; + TransactionalValue frame_accumulation; TransactionalValue camera; @@ -267,15 +297,17 @@ MainRenderer::init(int argc, const char** argv, Scene scene, Camera camera) } inline void -MainRenderer::set_scene(Scene scene) +MainRenderer::set_scene(const Scene& scene) { // TODO generalize to support multiple transfer functions // - assert(scene.instances.size() == 1); - if (scene.instances[0].models.size() == 1 && scene.instances[0].models[0].type == scene::Model::VOLUMETRIC_MODEL) - { - assert(scene.instances[0].models[0].volume_model.volume.type == scene::Volume::STRUCTURED_REGULAR_VOLUME); - scene::TransferFunction scene_tfn = scene.instances[0].models[0].volume_model.transfer_function; + scene::TransferFunction scene_tfn; + int count = count_tfn(scene, scene_tfn); + if (count > 1) { + std::cerr << "ERROR: found multiple transfer functions, they will be treated as one" << std::endl; + } + // TODO: find a better way to set transfer function // + if (count > 0) { const float* data_o = scene_tfn.opacity->data_typed(); const size_t size_o = scene_tfn.opacity->dims.v; @@ -312,6 +344,3 @@ MainRenderer::set_scene(Scene scene) std::shared_ptr create_renderer(std::string name); - -ovr::Scene -create_example_scene(); diff --git a/ovr/scene.h b/ovr/scene.h index e0fd863..6614f45 100644 --- a/ovr/scene.h +++ b/ovr/scene.h @@ -186,11 +186,18 @@ using array_1d_float3_t = std::shared_ptr; using array_1d_float4_t = std::shared_ptr; using array_3d_scalar_t = std::shared_ptr; +using array_1d_t = std::shared_ptr>; +using array_2d_t = std::shared_ptr>; +using array_3d_t = std::shared_ptr>; + #if defined(__cplusplus) // ------------------------------------------------------------------ // Scene Definitions // ------------------------------------------------------------------ + +namespace scene { + struct Camera { // camera position - *from* where we are looking vec3f from; @@ -202,7 +209,7 @@ struct Camera { ORTHOGRAPHIC, } type = PERSPECTIVE; - // TODO: + // TODO: implement missing features // affine3f transform: additional world-space transform, overridden by motion.* arrays // float nearClip: 10-6 near clipping distance // vec2f imageStart: (0,0) start of image region (lower left corner) @@ -210,7 +217,7 @@ struct Camera { struct PerspectiveCamera { float fovy = 60.f; - // TODO: + // TODO: implement missing features // float aspect; // float apertureRadius; // float focusDistance; @@ -223,8 +230,6 @@ struct Camera { } orthographic; }; -namespace scene { - struct TransferFunction { array_1d_float4_t color; array_1d_scalar_t opacity; @@ -232,7 +237,7 @@ struct TransferFunction { }; struct Volume { - enum { + enum VolumeType { STRUCTURED_REGULAR_VOLUME, } type; @@ -243,9 +248,43 @@ struct Volume { } structured_regular; }; +struct Texture { + enum { + VOLUME_TEXTURE, + TRANSFER_FUNCTION_TEXTURE, + } type; + + struct VolumeTexture { + Volume volume; + } volume; + + struct TransferFunctionTexture { + TransferFunction transfer_function; + int32_t volume_texture = -1; + } transfer_function; +}; + +struct Material { + enum { + OBJ_MATERIAL, + } type; + + struct ObjMaterial { + vec3f kd = vec3f(0.8f); // diffuse reflectivity + vec3f ks = vec3f(0.0f); // specular reflectivity + float ns = 10.f; // specular exponent + float d = 1.f; // opacity + vec3f tf = vec3f(1.f); // transparency filter + // texture maps + int32_t map_kd = -1; + int32_t map_bump = -1; + } obj; +}; + struct Geometry { enum { TRIANGLES_GEOMETRY, + ISOSURFACE_GEOMETRY, } type; struct GeometryTriangles { /* TODO */ @@ -257,6 +296,11 @@ struct Geometry { array_1d_float4_t color; } verts, faces; } triangles; + + struct GeometryIsosurfaces { + int32_t volume_texture; + std::vector isovalues; + } isosurfaces; }; struct Model { @@ -267,11 +311,13 @@ struct Model { struct VolumetricModel { TransferFunction transfer_function; - Volume volume; + int32_t volume_texture; + // Volume volume; } volume_model; struct GeometricModel { Geometry geometry; + int32_t mtl = -1; } geometry_model; }; @@ -304,9 +350,12 @@ struct Light { }; struct Scene { + std::vector textures; + std::vector materials; + std::vector instances; std::vector lights; - Camera camera; + scene::Camera camera; int ao_samples = 0; int spp = 1; @@ -323,7 +372,7 @@ struct Scene { } // namespace scene -using Scene = scene::Scene; +using scene::Scene; // ------------------------------------------------------------------ // Factory Functions @@ -357,6 +406,25 @@ CreateColorMap(const std::string& name); array_3d_scalar_t CreateArray3DScalarFromFile(const std::string& filename, vec3i dims, ValueType type, size_t offset, bool is_big_endian); +// ------------------------------------------------------------------ +// +// ------------------------------------------------------------------ + +inline const scene::Volume& +parse_single_volume_scene(const scene::Scene& scene, scene::Volume::VolumeType vtype = scene::Volume::STRUCTURED_REGULAR_VOLUME) +{ + if (scene.instances.size() != 1) throw std::runtime_error("expect only one instance"); + if (scene.instances[0].models.size() != 1) throw std::runtime_error("expect only one model"); + if (scene.instances[0].models[0].type != scene::Model::VOLUMETRIC_MODEL) throw std::runtime_error("expect a volume model"); + // check texture type + int32_t tex = scene.instances[0].models[0].volume_model.volume_texture; + if (tex < 0 && tex >= scene.textures.size()) throw std::runtime_error("invalid texture index: " + std::to_string(tex)); + if (scene.textures[tex].type != scene::Texture::VOLUME_TEXTURE) throw std::runtime_error("expect a volume texture"); + // check volume type + if (scene.textures[tex].volume.volume.type != vtype) throw std::runtime_error("expect a volume of type: " + std::to_string((int)vtype)); + return scene.textures[tex].volume.volume; +} + #endif // defined(__cplusplus) } // namespace ovr diff --git a/ovr/serializer/serializer_vidi3d.cpp b/ovr/serializer/serializer_vidi3d.cpp index 2276545..e7f549c 100644 --- a/ovr/serializer/serializer_vidi3d.cpp +++ b/ovr/serializer/serializer_vidi3d.cpp @@ -279,7 +279,7 @@ create_scene_tfn(const json& jsview, ValueType type) ovr::scene::Volume create_scene_volume(const json& jsdata, std::string workdir) { - scene::Volume volume{}; + ovr::scene::Volume volume{}; const auto format = jsdata[FORMAT].get(); @@ -307,12 +307,12 @@ create_scene_volume(const json& jsdata, std::string workdir) return volume; } -Camera +ovr::scene::Camera create_scene_camera(const json& jsview) { const auto& jscamera = jsview[CAMERA]; - Camera camera; + ovr::scene::Camera camera; camera.from = scalar_from_json(jscamera[EYE]); camera.at = scalar_from_json(jscamera[CENTER]); @@ -334,6 +334,8 @@ using namespace ovr::vidi3d; Scene create_json_scene_vidi3d(json root, std::string workdir) { + Scene scene; + Instance instance; instance.transform = affine3f::translate(vec3f(0)); @@ -351,15 +353,19 @@ create_json_scene_vidi3d(json root, std::string workdir) } } + Texture texture; + texture.type = Texture::VOLUME_TEXTURE; + texture.volume.volume = volume; + scene.textures.push_back(texture); + Model model; model.type = Model::VOLUMETRIC_MODEL; - model.volume_model.volume = volume; + model.volume_model.volume_texture = scene.textures.size() - 1; model.volume_model.transfer_function = tfn; instance.models.push_back(model); } - Scene scene; scene.instances.push_back(instance); if (root[VIEW].contains(LIGHT_SOURCE)) {