diff --git a/README.md b/README.md index 091466a9..fefeb94b 100644 --- a/README.md +++ b/README.md @@ -140,7 +140,6 @@ The renderer implementation has the following limitations at this time: * Volume absorption and scattering assumes homogeneous volume coefficients. There's simply no volume data primitive in the scene to sample heterogeneous coefficients from. * Volume scattering emission intensity not implemented. Goes along with heterogeneous volume scattering support. * Geometry displacement not implemented. To be revisited once Displacement Micro Meshes (DMM) are supported by OptiX. -* Hair BSDF not implemented. The renderer currently supports one texture coordinate, hair BSDF requires two to be fully featured and that'll increase the cost for every shader, so it's deferred into a separate example. * UV-tile and animation textures not implemented. * rounded_corner_normal() not implemented. That cannot be done by the automatic code generation for the geometry.normal expression. * Spot and light profile `global_distribution: true` not supported by the MDL SDK code generation. Lights will be black then. The renderer supports own point lights with spot and IES distribution though. @@ -151,6 +150,19 @@ Everything else inside the MDL specifications should just work! ![MDL_renderer with MDL materials](./apps/MDL_renderer/MDL_renderer_demo.png) ![MDL_renderer with vMaterials](./apps/MDL_renderer/MDL_renderer_vMaterials.png) +The MDL_renderer has now been updated to also support cubic B-spline curve primitives and the MDL Hair BSDF. + +Because that requires two texture coordinates to be fully featured, the `NUM_TEXTURE_SPACES` define has been added to the `config.h` to allow switching between one and two texture coordinates. If you do not need the hair BSDF, you can set `NUM_TEXTURE_SPACES` to 1 for a little more performance. + +The MDL hair BSDF supports a fully parameterized fiber surface accessible via the `state::texture_coordinate(0)` providing (uFiber, vFiber, thickness) values, which allows implementing parameter changes along the whole fiber and even aound it. The provided `mdl/bsdf_hair_uv.mdl` material shows this by placing tiny arrows on the fibers pointing from root to tip. + +Additionally the second texture coordinate `state::texture_coordinate(1)` defines a fixed texture coordinate per fiber, which allows coloring of individual fibers depending on some texture value. The image below used a Perlin noise function to produce highlights in the hair, resp. a 2D texture to color the fibers of the `fur.hair` model (included). + +The renderer currently loads only `*.hair` models which do not have texture coordinates. The example auto-generates a 2D coordinate with a cubemap projection from the root points' center coordinate. There are better ways to do this when actually growing hair from surfaces, not done in this example. Transparency and color values of *.hair files are ignored. The assigned MDL hair material defines these properties. + +![MDL_renderer with hair rendering](./apps/MDL_renderer/MDL_renderer_highlights.png) +![MDL_renderer with fur rendering](./apps/MDL_renderer/MDL_renderer_fur.png) + **User Interaction inside the examples**: * Left Mouse Button + Drag = Orbit (around center of interest) * Middle Mouse Button + Drag = Pan (The mouse ratio field in the GUI defines how many pixels is one unit.) @@ -331,6 +343,11 @@ For a lot more complex materials (this scene requires about 5.4 GB of VRAM), the * `MDL_renderer.exe -s system_mdl_vMaterials.txt -d scene_mdl_vMaterials.txt` +For the curves rendering with MDL hair BSDF materials, issue the command line. That will display a sphere with cubic B-spline curves using a red hair material lit by an area light from above. Please read the `scene_mdl_hair.txt? for other possible material and model configurations. + +* `MDL_renderer.exe -s system_mdl_hair.txt -d scene_mdl_hair.txt` + + # Pull Requests NVIDIA is happy to review and consider pull requests for merging into the main tree of the optix_apps for bug fixes and features. Before providing a pull request to NVIDIA, please note the following: diff --git a/apps/MDL_renderer/CMakeLists.txt b/apps/MDL_renderer/CMakeLists.txt index c5238343..74eda838 100644 --- a/apps/MDL_renderer/CMakeLists.txt +++ b/apps/MDL_renderer/CMakeLists.txt @@ -118,6 +118,7 @@ set( HEADERS inc/CheckMacros.h inc/CompileResult.h inc/Device.h + inc/Hair.h inc/LightGUI.h inc/LoaderIES.h inc/MaterialMDL.h @@ -141,7 +142,9 @@ set( SOURCES src/Assimp.cpp src/Box.cpp src/Camera.cpp + src/Curves.cpp src/Device.cpp + src/Hair.cpp src/LoaderIES.cpp src/main.cpp src/MaterialMDL.cpp @@ -181,6 +184,8 @@ set( SHADERS_HEADERS ${CMAKE_CURRENT_SOURCE_DIR}/shaders/camera_definition.h ${CMAKE_CURRENT_SOURCE_DIR}/shaders/compositor_data.h ${CMAKE_CURRENT_SOURCE_DIR}/shaders/config.h + ${CMAKE_CURRENT_SOURCE_DIR}/shaders/curve.h + ${CMAKE_CURRENT_SOURCE_DIR}/shaders/curve_attributes.h ${CMAKE_CURRENT_SOURCE_DIR}/shaders/function_indices.h ${CMAKE_CURRENT_SOURCE_DIR}/shaders/light_definition.h ${CMAKE_CURRENT_SOURCE_DIR}/shaders/material_definition_mdl.h diff --git a/apps/MDL_renderer/MDL_renderer_fur.png b/apps/MDL_renderer/MDL_renderer_fur.png new file mode 100644 index 00000000..2b0173fc Binary files /dev/null and b/apps/MDL_renderer/MDL_renderer_fur.png differ diff --git a/apps/MDL_renderer/MDL_renderer_highlights.png b/apps/MDL_renderer/MDL_renderer_highlights.png new file mode 100644 index 00000000..7e1db689 Binary files /dev/null and b/apps/MDL_renderer/MDL_renderer_highlights.png differ diff --git a/apps/MDL_renderer/inc/Device.h b/apps/MDL_renderer/inc/Device.h index 8f210d13..f63102dc 100644 --- a/apps/MDL_renderer/inc/Device.h +++ b/apps/MDL_renderer/inc/Device.h @@ -161,7 +161,6 @@ struct DeviceAttribute int concurrentManagedAccess; int computePreemptionSupported; int canUseHostPointerForRegisteredMem; - int canUseStreamMemOps; int canUse64BitStreamMemOps; int canUseStreamWaitValueNor; int cooperativeLaunch; @@ -229,6 +228,9 @@ enum ProgramGroupId // 3 = emission, cutout PGID_HIT_RADIANCE_3, PGID_HIT_SHADOW_3, + // 4 = cubic B-spline curves. + PGID_HIT_CURVES, + PGID_HIT_CURVES_SHADOW, // Direct Callables // Lens shader @@ -252,7 +254,8 @@ enum ProgramGroupId enum PrimitiveType { PT_UNKNOWN, // It's an error when this is still set. - PT_TRIANGLES + PT_TRIANGLES, + PT_CURVES }; @@ -271,7 +274,7 @@ struct GeometryData info = {}; } - PrimitiveType primitiveType; // Triangles ony in this renderer. + PrimitiveType primitiveType; // Triangles or cubic B-spline curves. Used to pick the correct hit record inside the SBT. int owner; // The device index which originally allocated all device side memory below. Needed when sharing GeometryData, resp. when freeing it. OptixTraversableHandle traversable; // The traversable handle for this GAS. Assigned to the Instance above it. CUdeviceptr d_attributes; // Array of VertexAttribute structs. @@ -341,6 +344,7 @@ class Device void initLights(const std::vector& lightsGUI, const std::vector& geometryData, const unsigned int stride, const unsigned int index); // Must be called after initScene()! GeometryData createGeometry(std::shared_ptr geometry); + GeometryData createGeometry(std::shared_ptr geometry); void destroyGeometry(GeometryData& data); void createInstance(const GeometryData& geometryData, const InstanceData& data, const float matrix[12]); void createTLAS(); diff --git a/apps/MDL_renderer/inc/Hair.h b/apps/MDL_renderer/inc/Hair.h new file mode 100644 index 00000000..329488b7 --- /dev/null +++ b/apps/MDL_renderer/inc/Hair.h @@ -0,0 +1,129 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +#pragma once + +#ifndef HAIR_H +#define HAIR_H + +// Hair file format and models courtesy of Cem Yuksel: http://www.cemyuksel.com/research/hairmodels/ + +// For the float3 +#include + +#include +#include + +#define HAS_SEGMENTS_ARRAY (1 << 0) +#define HAS_POINTS_ARRAY (1 << 1) +#define HAS_THICKNESS_ARRAY (1 << 2) +#define HAS_TRANSPARENCY_ARRAY (1 << 3) +#define HAS_COLOR_ARRAY (1 << 4) + +// Note that the Hair class setter interfaces are unused in this example but allow generating *.hair files easily. +// +// At minimum that needs an array of float3 points and the matching header values. +// This renderer doesn't care about the transparency and color inside the *.hair model file. +// The material is defined by the assigned MDL hair BSDF material. +// Not setting the color might result in black in other applications. +// +// Example for setting individual cubic B-spline curves with four control points each. +// ... // Build your std::vector pointsArray here. +//Hair hair; +//hair.setNumStrands(numStrands); // Number of hair strands inside the file. +//hair.setNumSegments(3); // Cubic B-Splines with 4 control points each are 3 segments. +//hair.setPointsArray(pointsArray); // This defines the m_header.numPoints as well. +//hair.setThickness(thickness); // Constant thickness. Use setThicknessArray() if values differ along strands. +//hair.save(filename); + +class Hair +{ + // HAIR File Header (128 Bytes) + struct Header + { + char signature[4]; // Bytes 0-3 Must be "HAIR" in ascii code (48 41 49 52) + unsigned int numStrands; // Bytes 4-7 Number of hair strands as unsigned int + unsigned int numPoints; // Bytes 8-11 Total number of points of all strands as unsigned int + unsigned int bits; // Bytes 12-15 Bit array of data in the file + // Bit-0 is 1 if the file has segments array. + // Bit-1 is 1 if the file has points array (this bit must be 1). + // Bit-2 is 1 if the file has thickness array. + // Bit-3 is 1 if the file has transparency array. + // Bit-4 is 1 if the file has color array. + // Bit-5 to Bit-31 are reserved for future extension (must be 0). + unsigned int numSegments; // Bytes 16-19 Default number of segments of hair strands as unsigned int + // If the file does not have a segments array, this default value is used. + float thickness; // Bytes 20-23 Default thickness hair strands as float + // If the file does not have a thickness array, this default value is used. + float transparency; // Bytes 24-27 Default transparency hair strands as float + // If the file does not have a transparency array, this default value is used. + float3 color; // Bytes 28-39 Default color hair strands as float array of size 3 + // If the file does not have a thickness array, this default value is used. + char information[88]; // Bytes 40-127 File information as char array of size 88 in ascii + }; + +public: + Hair(); + //~Hair(); + + bool load(const std::string& filename); + bool save(const std::string& filename); + + void setNumStrands(const unsigned int num); + unsigned int getNumStrands() const; + + void setNumSegments(const unsigned int num); + void setSegmentsArray(const std::vector& segments); + unsigned int getNumSegments(const unsigned int idxStrand) const; + + void setPointsArray(const std::vector& points); // This also sets m_header.numPoints! + float3 getPoint(const unsigned int idx) const; + + void setThickness(const float thickness); + void setThicknessArray(const std::vector& thickness); + float getThickness(const unsigned int idx) const; + + void setTransparency(const float transparency); + void setTransparencyArray(const std::vector& transparency); + float getTransparency(const unsigned int idx) const; + + void setColor(const float3 color); + void setColorArray(const std::vector& color); + float3 getColor(const unsigned int idx) const; + +private: + Header m_header; + + std::vector m_segmentsArray; // Empty or size numStrands. + std::vector m_pointsArray; // Size numPoints. + std::vector m_thicknessArray; // Empty or size numPoints. + std::vector m_transparencyArray; // Empty or size numPoints. + std::vector m_colorArray; // Empty or size numPoints. +}; + +#endif // HAIR_H diff --git a/apps/MDL_renderer/inc/SceneGraph.h b/apps/MDL_renderer/inc/SceneGraph.h index 99f2e5d2..271e82e2 100644 --- a/apps/MDL_renderer/inc/SceneGraph.h +++ b/apps/MDL_renderer/inc/SceneGraph.h @@ -34,6 +34,7 @@ // For the vector types. #include +#include "shaders/curve_attributes.h" #include "shaders/vertex_attributes.h" #include "shaders/vector_math.h" @@ -49,7 +50,8 @@ namespace sg { NT_GROUP, NT_INSTANCE, - NT_TRIANGLES + NT_TRIANGLES, + NT_CURVES }; class Node @@ -93,9 +95,31 @@ namespace sg private: std::vector m_attributes; - std::vector m_indices; // If m_indices.size() == 0, m_attributes are independent primitives. // Not actually supported in this renderer implementation. + std::vector m_indices; // If m_indices.size() == 0, m_attributes are independent primitives (not actually supported in this renderer implementation!) }; + class Curves : public Node + { + public: + Curves(const unsigned int id); + //~Curves(); + + sg::NodeType getType() const; + + bool createHair(std::string const& filename, const float scale); + + void setAttributes(std::vector const& attributes); + std::vector const& getAttributes() const; + + void setIndices(std::vector const&); + std::vector const& getIndices() const; + + private: + std::vector m_attributes; + std::vector m_indices; + }; + + class Instance : public Node { public: diff --git a/apps/MDL_renderer/inc/ShaderConfiguration.h b/apps/MDL_renderer/inc/ShaderConfiguration.h index 286a8de1..3a368382 100644 --- a/apps/MDL_renderer/inc/ShaderConfiguration.h +++ b/apps/MDL_renderer/inc/ShaderConfiguration.h @@ -63,6 +63,7 @@ struct ShaderConfiguration bool use_volume_scattering; bool is_cutout_opacity_constant; bool use_cutout_opacity; + bool is_hair_bsdf_valid; // The constant expression values: bool thin_walled; diff --git a/apps/MDL_renderer/shaders/config.h b/apps/MDL_renderer/shaders/config.h index 2aae1b49..4f46c63e 100644 --- a/apps/MDL_renderer/shaders/config.h +++ b/apps/MDL_renderer/shaders/config.h @@ -61,4 +61,23 @@ #define INTEROP_MODE_TEX 1 #define INTEROP_MODE_PBO 2 +// The number of supported texture coordinate slots inside MDL shaders (state::texture_*(i)). +// This makes the Mdl_state bigger which will cost performance! +// It's also less convenient to use the TBN ortho-normal basis. +// hair_bsdf() requires two texture coordinates to communicate the intersection results +// and a per fiber texture coordinate which can be used to color each hair individually. +// That's the whole reason for this define. +// HACK The renderer's vertex attributes are not acctually managing two distinct texture coordinates. +// The Mdl_state will simply duplicate the data of texture coordinate slot 0 to slot 1 everywhere except for the hair_bsdf(). +#define NUM_TEXTURE_SPACES 2 + +// The number of float4 elements inside the texture_results cache. Default is 16. +// Used to configure the MDL backend's code generation with set_option("num_texture_results", ...) +// This value influences how many things can be precalculated inside the init() function. +// If the number of result elements in this array is lower than what is required, +// the expressions for the remaining results will be compiled into the sample() and evaluate() functions +// which will make the compilation and runtime performance slower. +// For very resource-heavy materials, experiment with bigger values. +#define NUM_TEXTURE_RESULTS 16 + #endif // CONFIG_H diff --git a/apps/MDL_renderer/shaders/curve.h b/apps/MDL_renderer/shaders/curve.h new file mode 100644 index 00000000..461119cc --- /dev/null +++ b/apps/MDL_renderer/shaders/curve.h @@ -0,0 +1,438 @@ +// +// Copyright (c) 2021-2023, NVIDIA CORPORATION. All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions +// are met: +// * Redistributions of source code must retain the above copyright +// notice, this list of conditions and the following disclaimer. +// * Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimer in the +// documentation and/or other materials provided with the distribution. +// * Neither the name of NVIDIA CORPORATION nor the names of its +// contributors may be used to endorse or promote products derived +// from this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY +// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR +// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR +// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, +// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, +// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR +// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY +// OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +// (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#pragma once + +#ifndef CURVE_H +#define CURVE_H + +#include "config.h" + +#include "vector_math.h" + +// First order polynomial interpolator +// +struct LinearInterpolator +{ + __device__ __forceinline__ LinearInterpolator() + { + } + + __device__ __forceinline__ void initialize(const float4* q) + { + p[0] = q[0]; + p[1] = q[1] - q[0]; + } + + __device__ __forceinline__ float4 position4(float u) const + { + return p[0] + u * p[1]; // Horner scheme + } + + __device__ __forceinline__ float3 position3(float u) const + { + return make_float3(position4(u)); + } + + __device__ __forceinline__ float radius(const float& u) const + { + return position4(u).w; + } + + __device__ __forceinline__ float4 velocity4(float u) const + { + return p[1]; + } + + __device__ __forceinline__ float3 velocity3(float u) const + { + return make_float3(velocity4(u)); + } + + __device__ __forceinline__ float derivative_of_radius(float u) const + { + return velocity4(u).w; + } + + __device__ __forceinline__ float3 acceleration3(float u) const + { + return make_float3(0.f); + } + __device__ __forceinline__ float4 acceleration4(float u) const + { + return make_float4(0.f); + } + + + float4 p[2]; +}; + + +// +// Second order polynomial interpolator +// +struct QuadraticInterpolator +{ + __device__ __forceinline__ QuadraticInterpolator() + { + } + + __device__ __forceinline__ void initializeFromBSpline(const float4* q) + { + // Bspline-to-Poly = Matrix([[1/2, -1, 1/2], + // [-1, 1, 0], + // [1/2, 1/2, 0]]) + p[0] = (q[0] - 2.0f * q[1] + q[2]) / 2.0f; + p[1] = (-2.0f * q[0] + 2.0f * q[1]) / 2.0f; + p[2] = (q[0] + q[1]) / 2.0f; + } + + __device__ __forceinline__ void export2BSpline(float4 bs[3]) const + { + // inverse of initializeFromBSpline + // Bspline-to-Poly = Matrix([[1/2, -1, 1/2], + // [-1, 1, 0], + // [1/2, 1/2, 0]]) + // invert to get: + // Poly-to-Bspline = Matrix([[0, -1/2, 1], + // [0, 1/2, 1], + // [2, 3/2, 1]]) + bs[0] = p[0] - p[1] / 2; + bs[1] = p[0] + p[1] / 2; + bs[2] = p[0] + 1.5f * p[1] + 2 * p[2]; + } + + __device__ __forceinline__ float4 position4(float u) const + { + return (p[0] * u + p[1]) * u + p[2]; // Horner scheme + } + + __device__ __forceinline__ float3 position3(float u) const + { + return make_float3(position4(u)); + } + + __device__ __forceinline__ float radius(float u) const + { + return position4(u).w; + } + + __device__ __forceinline__ float4 velocity4(float u) const + { + return 2.0f * p[0] * u + p[1]; + } + + __device__ __forceinline__ float3 velocity3(float u) const + { + return make_float3(velocity4(u)); + } + + __device__ __forceinline__ float derivative_of_radius(float u) const + { + return velocity4(u).w; + } + + __device__ __forceinline__ float4 acceleration4(float u) const + { + return 2.0f * p[0]; + } + + __device__ __forceinline__ float3 acceleration3(float u) const + { + return make_float3(acceleration4(u)); + } + + + float4 p[3]; +}; + +// +// Third order polynomial interpolator +// +// Storing {p0, p1, p2, p3} for evaluation: +// P(u) = p0 * u^3 + p1 * u^2 + p2 * u + p3 +// +struct CubicInterpolator +{ + __device__ __forceinline__ CubicInterpolator() + { + } + // TODO: Initialize from polynomial weights. Check that sample doesn't rely on + // legacy behavior. + // __device__ __forceinline__ CubicBSplineSegment( const float4* q ) { initializeFromBSpline( q ); } + + + __device__ __forceinline__ void initializeFromBSpline(const float4* q) + { + // Bspline-to-Poly = Matrix([[-1/6, 1/2, -1/2, 1/6], + // [ 1/2, -1, 1/2, 0], + // [-1/2, 0, 1/2, 0], + // [ 1/6, 2/3, 1/6, 0]]) + // Original: + //p[0] = ( q[0] * ( -1.0f ) + q[1] * ( 3.0f ) + q[2] * ( -3.0f ) + q[3] ) / 6.0f; + //p[1] = ( q[0] * ( 3.0f ) + q[1] * ( -6.0f ) + q[2] * ( 3.0f ) ) / 6.0f; + //p[2] = ( q[0] * ( -3.0f ) + q[2] * ( 3.0f ) ) / 6.0f; + //p[3] = ( q[0] * ( 1.0f ) + q[1] * ( 4.0f ) + q[2] * ( 1.0f ) ) / 6.0f; + // Optimized: + const float o_s = 1.0f / 6.0f; + p[0] = (q[3] - q[0]) * o_s + (q[1] - q[2]) * 0.5f; + p[1] = (q[0] + q[2]) * 0.5f - q[1]; + p[2] = (q[2] - q[0]) * 0.5f; + p[3] = (q[0] + q[1] * 4.0f + q[2]) * o_s; + } + + __device__ __forceinline__ void export2BSpline(float4 bs[4]) const + { + // inverse of initializeFromBSpline + // Bspline-to-Poly = Matrix([[-1/6, 1/2, -1/2, 1/6], + // [ 1/2, -1, 1/2, 0], + // [-1/2, 0, 1/2, 0], + // [ 1/6, 2/3, 1/6, 0]]) + // invert to get: + // Poly-to-Bspline = Matrix([[0, 2/3, -1, 1], + // [0, -1/3, 0, 1], + // [0, 2/3, 1, 1], + // [6, 11/3, 2, 1]]) + bs[0] = (p[1] * (2.0f) + p[2] * (-1.0f) + p[3]) / 3.0f; + bs[1] = (p[1] * (-1.0f) + p[3]) / 3.0f; + bs[2] = (p[1] * (2.0f) + p[2] * (1.0f) + p[3]) / 3.0f; + bs[3] = (p[0] + p[1] * (11.0f) + p[2] * (2.0f) + p[3]) / 3.0f; + } + + + __device__ __forceinline__ void initializeFromCatrom(const float4* q) + { + // Catrom-to-Poly = Matrix([[-1/2, 3/2, -3/2, 1/2], + // [1, -5/2, 2, -1/2], + // [-1/2, 0, 1/2, 0], + // [0, 1, 0, 0]]) + p[0] = (-1.0f * q[0] + (3.0f) * q[1] + (-3.0f) * q[2] + (1.0f) * q[3]) / 2.0f; + p[1] = (2.0f * q[0] + (-5.0f) * q[1] + (4.0f) * q[2] + (-1.0f) * q[3]) / 2.0f; + p[2] = (-1.0f * q[0] + (1.0f) * q[2]) / 2.0f; + p[3] = ((2.0f) * q[1]) / 2.0f; + } + + __device__ __forceinline__ void export2Catrom(float4 cr[4]) const + { + // Catrom-to-Poly = Matrix([[-1/2, 3/2, -3/2, 1/2], + // [1, -5/2, 2, -1/2], + // [-1/2, 0, 1/2, 0], + // [0, 1, 0, 0]]) + // invert to get: + // Poly-to-Catrom = Matrix([[1, 1, -1, 1], + // [0, 0, 0, 1], + // [1, 1, 1, 1], + // [6, 4, 2, 1]]) + cr[0] = (p[0] * 6.f / 6.f) - (p[1] * 5.f / 6.f) + (p[2] * 2.f / 6.f) + (p[3] * 1.f / 6.f); + cr[1] = (p[0] * 6.f / 6.f); + cr[2] = (p[0] * 6.f / 6.f) + (p[1] * 1.f / 6.f) + (p[2] * 2.f / 6.f) + (p[3] * 1.f / 6.f); + cr[3] = (p[0] * 6.f / 6.f) + (p[3] * 6.f / 6.f); + } + + __device__ __forceinline__ float4 position4(float u) const + { + return (((p[0] * u) + p[1]) * u + p[2]) * u + p[3]; // Horner scheme + } + + __device__ __forceinline__ float3 position3(float u) const + { + // rely on compiler and inlining for dead code removal + return make_float3(position4(u)); + } + __device__ __forceinline__ float radius(float u) const + { + return position4(u).w; + } + + __device__ __forceinline__ float4 velocity4(float u) const + { + // adjust u to avoid problems with triple knots. + if (u == 0) + u = 0.000001f; + if (u == 1) + u = 0.999999f; + return ((3.0f * p[0] * u) + 2.0f * p[1]) * u + p[2]; + } + + __device__ __forceinline__ float3 velocity3(float u) const + { + return make_float3(velocity4(u)); + } + + __device__ __forceinline__ float derivative_of_radius(float u) const + { + return velocity4(u).w; + } + + __device__ __forceinline__ float4 acceleration4(float u) const + { + return 6.0f * p[0] * u + 2.0f * p[1]; // Horner scheme + } + + __device__ __forceinline__ float3 acceleration3(float u) const + { + return make_float3(acceleration4(u)); + } + + float4 p[4]; +}; + + +// Compute curve primitive surface normal in object space. +// +// Template parameters: +// CurveType - A B-Spline evaluator class. +// type - 0 ~ cylindrical approximation (correct if radius' == 0) +// 1 ~ conic approximation (correct if curve'' == 0) +// other ~ the bona fide surface normal +// +// Parameters: +// bc - A B-Spline evaluator object. +// u - segment parameter of hit-point. +// ps - hit-point on curve's surface in object space; usually +// computed like this. +// float3 ps = ray_orig + t_hit * ray_dir; +// the resulting point is slightly offset away from the +// surface. For this reason (Warning!) ps gets modified by this +// method, projecting it onto the surface +// in case it is not already on it. (See also inline +// comments.) +// +template +__device__ __forceinline__ float3 surfaceNormal(const CurveType& bc, float u, float3& ps) +{ + float3 normal; + if (u == 0.0f) + { + normal = -bc.velocity3(0); // special handling for flat endcaps + } + else if (u == 1.0f) + { + normal = bc.velocity3(1); // special handling for flat endcaps + } + else + { + // ps is a point that is near the curve's offset surface, + // usually ray.origin + ray.direction * rayt. + // We will push it exactly to the surface by projecting it to the plane(p,d). + // The function derivation: + // we (implicitly) transform the curve into coordinate system + // {p, o1 = normalize(ps - p), o2 = normalize(curve'(t)), o3 = o1 x o2} in which + // curve'(t) = (0, length(d), 0); ps = (r, 0, 0); + float4 p4 = bc.position4(u); + float3 p = make_float3(p4); + float r = p4.w; // == length(ps - p) if ps is already on the surface + float4 d4 = bc.velocity4(u); + float3 d = make_float3(d4); + float dr = d4.w; + float dd = dot(d, d); + + float3 o1 = ps - p; // dot(modified_o1, d) == 0 by design: + o1 -= (dot(o1, d) / dd) * d; // first, project ps to the plane(p,d) + o1 *= r / length(o1); // and then drop it to the surface + ps = p + o1; // fine-tuning the hit point + if (type == 0) + { + normal = o1; // cylindrical approximation + } + else + { + if (type != 1) + { + dd -= dot(bc.acceleration3(u), o1); + } + normal = dd * o1 - (dr * r) * d; + } + } + return normalize(normal); +} + +template +__device__ __forceinline__ float3 surfaceNormal(const LinearInterpolator& bc, float u, float3& ps) +{ + float3 normal; + if (u == 0.0f) + { + normal = ps - (float3&) (bc.p[0]); // special handling for round endcaps + } + else if (u >= 1.0f) + { + // reconstruct second control point (Note: the interpolator pre-transforms + // the control-points to speed up repeated evaluation. + const float3 p1 = (float3&) (bc.p[1]) + (float3&) (bc.p[0]); + normal = ps - p1; // special handling for round endcaps + } + else + { + // ps is a point that is near the curve's offset surface, + // usually ray.origin + ray.direction * rayt. + // We will push it exactly to the surface by projecting it to the plane(p,d). + // The function derivation: + // we (implicitly) transform the curve into coordinate system + // {p, o1 = normalize(ps - p), o2 = normalize(curve'(t)), o3 = o1 x o2} in which + // curve'(t) = (0, length(d), 0); ps = (r, 0, 0); + float4 p4 = bc.position4(u); + float3 p = make_float3(p4); + float r = p4.w; // == length(ps - p) if ps is already on the surface + float4 d4 = bc.velocity4(u); + float3 d = make_float3(d4); + float dr = d4.w; + float dd = dot(d, d); + + float3 o1 = ps - p; // dot(modified_o1, d) == 0 by design: + o1 -= (dot(o1, d) / dd) * d; // first, project ps to the plane(p,d) + o1 *= r / length(o1); // and then drop it to the surface + ps = p + o1; // fine-tuning the hit point + if (type == 0) + { + normal = o1; // cylindrical approximation + } + else + { + normal = dd * o1 - (dr * r) * d; + } + } + return normalize(normal); +} + +// Compute curve primitive tangent in object space. +// +// Template parameters: +// CurveType - A B-Spline evaluator class. +// +// Parameters: +// bc - A B-Spline evaluator object. +// u - segment parameter of tangent location on curve. +// +template +__device__ __forceinline__ float3 curveTangent(const CurveType& bc, float u) +{ + float3 tangent = bc.velocity3(u); + return normalize(tangent); +} + +#endif // CURVE_H diff --git a/apps/MDL_renderer/shaders/curve_attributes.h b/apps/MDL_renderer/shaders/curve_attributes.h new file mode 100644 index 00000000..ff8726bb --- /dev/null +++ b/apps/MDL_renderer/shaders/curve_attributes.h @@ -0,0 +1,41 @@ +/* + * Copyright (c) 2019-2023, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +#pragma once + +#ifndef CURVE_ATTRIBUTES_H +#define CURVE_ATTRIBUTES_H + +struct CurveAttributes +{ + float4 vertex; // .xyz = 3D position in object space, .w = radius in world space. + float4 reference; // .xyz = Constant reference vector along the fiber. Used to calculate circular intersection value vFiber [0.0, 1.0] around the hair. .w = unused. + float4 texcoord; // .xyz = 3D texture coordinates. .w = interpolant along the whole fiber. Used to calculate uFiber in range [0.0, 1.0] from root to tip. +}; + +#endif // CURVE_ATTRIBUTES_H diff --git a/apps/MDL_renderer/shaders/hit.cu b/apps/MDL_renderer/shaders/hit.cu index e6f3e429..a8249c61 100644 --- a/apps/MDL_renderer/shaders/hit.cu +++ b/apps/MDL_renderer/shaders/hit.cu @@ -39,6 +39,8 @@ #include "shader_common.h" #include "transform.h" #include "random_number_generators.h" +#include "curve.h" +#include "curve_attributes.h" // Contained in per_ray_data.h: //#include @@ -73,6 +75,28 @@ extern "C" __constant__ SystemData sysData; // This shader handles every supported feature of the the renderer. extern "C" __global__ void __closesthit__radiance() { + // Get the current rtPayload pointer from the unsigned int payload registers p0 and p1. + PerRayData* thePrd = mergePointer(optixGetPayload_0(), optixGetPayload_1()); + + thePrd->flags |= FLAG_HIT; // Required to distinguish surface hits from random walk miss. + + thePrd->distance = optixGetRayTmax(); // Return the current path segment distance, needed for absorption calculations in the integrator. + + // PRECISION Calculate this from the object space vertex positions and transform to world for better accuracy when needed. + // Same as: thePrd->pos = optixGetWorldRayOrigin() + optixGetWorldRayDirection() * optixGetRayTmax(); + thePrd->pos += thePrd->wi * thePrd->distance; + + // If we're inside a volume and hit something, the path throughput needs to be modulated + // with the transmittance along this segment before adding surface or light radiance! + if (0 < thePrd->idxStack) // This assumes the first stack entry is vaccuum. + { + thePrd->throughput *= expf(thePrd->sigma_t * -thePrd->distance); + + // Increment the volume scattering random walk counter. + // Unused when FLAG_VOLUME_SCATTERING is not set. + ++thePrd->walk; + } + GeometryInstanceData theData = sysData.geometryInstanceData[optixGetInstanceId()]; // theData.ids: .x = idMaterial, .y = idLight, .z = idObject @@ -88,7 +112,7 @@ extern "C" __global__ void __closesthit__radiance() const TriangleAttributes& attr1 = attributes[tri.y]; const TriangleAttributes& attr2 = attributes[tri.z]; - const float2 theBarycentrics = optixGetTriangleBarycentrics(); // beta and gamma + const float2 theBarycentrics = optixGetTriangleBarycentrics(); // .x = beta, .y = gamma const float alpha = 1.0f - theBarycentrics.x - theBarycentrics.y; float4 objectToWorld[3]; @@ -96,47 +120,48 @@ extern "C" __global__ void __closesthit__radiance() getTransforms(optixGetTransformListHandle(0), objectToWorld, worldToObject); // Single instance level transformation list only. + // Object space vertex attributes at the hit point. + float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; float3 ng = cross(attr1.vertex - attr0.vertex, attr2.vertex - attr0.vertex); - float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; - - const float3 tc = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - // Transform into internal space == world space. + // Transform attributes into internal space == world space. + ns = normalize(transformNormal(worldToObject, ns)); ng = normalize(transformNormal(worldToObject, ng)); + // This is actually the geometry tangent which for the runtime generated geometry objects + // (plane, box, sphere, torus) match exactly with the texture space tangent. + // FIXME Generate these from the triangle's texture derivatives instead, but that's more expensive. + // Mind that tangents and bitangents are transformed as vectors, not normals, because they lie inside the surface's plane. tg = normalize(transformVector(objectToWorld, tg)); - ns = normalize(transformNormal(worldToObject, ns)); - - TBN tbn(tg, ns); // Calculate an otho-normal system respective to the shading normal. - - // Get the current rtPayload pointer from the unsigned int payload registers p0 and p1. - PerRayData* thePrd = mergePointer(optixGetPayload_0(), optixGetPayload_1()); - - thePrd->flags |= FLAG_HIT; // Required to distinguish surface hits from random walk miss. - - thePrd->distance = optixGetRayTmax(); // Return the current path segment distance, needed for absorption calculations in the integrator. + // Calculate an ortho-normal system respective to the shading normal. + // Expanding the TBN tbn(tg, ns) constructor because TBN members can't be used as pointers for the Mdl_state with NUM_TEXTURE_SPACES > 1. + float3 bt = normalize(cross(ns, tg)); + tg = cross(bt, ns); // Now the tangent is orthogonal to the shading normal. + + // The Mdl_state holds the texture attributes per texture space in separate arrays. + float3 texture_coordinates[NUM_TEXTURE_SPACES]; + float3 texture_tangents[NUM_TEXTURE_SPACES]; + float3 texture_bitangents[NUM_TEXTURE_SPACES]; + + // NUM_TEXTURE_SPACES is always at least 1. + texture_coordinates[0] = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + texture_bitangents[0] = bt; + texture_tangents[0] = tg; - // PRECISION Calculate this from the object space vertex positions and transform to world for better accuracy when needed. - // Same as: thePrd->pos = optixGetWorldRayOrigin() + optixGetWorldRayDirection() * optixGetRayTmax(); - thePrd->pos += thePrd->wi * thePrd->distance; - - // If we're inside a volume and hit something, the path throughput needs to be modulated - // with the transmittance along this segment before adding surface or light radiance! - if (0 < thePrd->idxStack) // This assumes the first stack entry is vaccuum. - { - thePrd->throughput *= expf(thePrd->sigma_t * -thePrd->distance); - - // Increment the volume scattering random walk counter. - // Unused when FLAG_VOLUME_SCATTERING is not set. - ++thePrd->walk; - } - +#if NUM_TEXTURE_SPACES == 2 + // HACK Simply copy the vertex attributes of texture space 0, simply because there is no second texcood inside TriangleAttributes. + texture_coordinates[1] = texture_coordinates[0]; + texture_bitangents[1] = bt; + texture_tangents[1] = tg; +#endif + + // Setup the Mdl_state. Mdl_state state; // The result of state::normal(). It represents the shading normal as determined by the renderer. // This field will be updated to the result of "geometry.normal" by the material or BSDF init functions, // if requested during code generation with set_option("include_geometry_normal", true) which is the default. - state.normal = tbn.normal; + state.normal = ns; // The result of state::geometry_normal(). // It represents the geometry normal as determined by the renderer. @@ -153,22 +178,20 @@ extern "C" __global__ void __closesthit__radiance() // An array containing the results of state::texture_coordinate(i). // The i-th entry represents the texture coordinates of the i-th texture space at the current position. // Only one element here because "num_texture_spaces" option has been set to 1. - state.text_coords = &tc; + state.text_coords = texture_coordinates; // An array containing the results of state::texture_tangent_u(i). // The i-th entry represents the texture tangent vector of the i-th texture space at the // current position, which points in the direction of the projection of the tangent to the // positive u axis of this texture space onto the plane defined by the original surface normal. - // Only one element because "num_texture_spaces" option has been set to 1. - state.tangent_u = &tbn.tangent; + state.tangent_u = texture_tangents; // An array containing the results of state::texture_tangent_v(i). // The i-th entry represents the texture bitangent vector of the i-th texture space at the // current position, which points in the general direction of the positive v axis of this // texture space, but is orthogonal to both the original surface normal and the tangent // of this texture space. - // Only one element because "num_texture_spaces" option has been set to 1. - state.tangent_v = &tbn.bitangent; + state.tangent_v = texture_bitangents; // The texture results lookup table. // The size must match the backend set_option("num_texture_results") value. @@ -179,8 +202,7 @@ extern "C" __global__ void __closesthit__radiance() // If the number of result elements in this array is lower than what is required, // the expressions for the remaining results will be compiled into the sample() and eval() functions // which will make the compilation and runtime performance slower. - // For very resource-heavy materials, experiment with bigger arrays. - float4 texture_results[16]; + float4 texture_results[NUM_TEXTURE_RESULTS]; state.text_results = texture_results; @@ -500,6 +522,28 @@ extern "C" __global__ void __closesthit__radiance() // PERF Identical to radiance shader above, but used for materials without emission, which is the majority of materials. extern "C" __global__ void __closesthit__radiance_no_emission() { + // Get the current rtPayload pointer from the unsigned int payload registers p0 and p1. + PerRayData* thePrd = mergePointer(optixGetPayload_0(), optixGetPayload_1()); + + thePrd->flags |= FLAG_HIT; // Required to distinguish surface hits from random walk miss. + + thePrd->distance = optixGetRayTmax(); // Return the current path segment distance, needed for absorption calculations in the integrator. + + // PRECISION Calculate this from the object space vertex positions and transform to world for better accuracy when needed. + // Same as: thePrd->pos = optixGetWorldRayOrigin() + optixGetWorldRayDirection() * optixGetRayTmax(); + thePrd->pos += thePrd->wi * thePrd->distance; + + // If we're inside a volume and hit something, the path throughput needs to be modulated + // with the transmittance along this segment before adding surface or light radiance! + if (0 < thePrd->idxStack) // This assumes the first stack entry is vaccuum. + { + thePrd->throughput *= expf(thePrd->sigma_t * -thePrd->distance); + + // Increment the volume scattering random walk counter. + // Unused when FLAG_VOLUME_SCATTERING is not set. + ++thePrd->walk; + } + GeometryInstanceData theData = sysData.geometryInstanceData[optixGetInstanceId()]; // theData.ids: .x = idMaterial, .y = idLight, .z = idObject @@ -523,53 +567,49 @@ extern "C" __global__ void __closesthit__radiance_no_emission() getTransforms(optixGetTransformListHandle(0), objectToWorld, worldToObject); // Single instance level transformation list only. + // Object space vertex attributes at the hit point. + float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; float3 ng = cross(attr1.vertex - attr0.vertex, attr2.vertex - attr0.vertex); - float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; - - const float3 tc = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - // Transform into internal space == world space. + // Transform attributes into internal space == world space. + ns = normalize(transformNormal(worldToObject, ns)); ng = normalize(transformNormal(worldToObject, ng)); tg = normalize(transformVector(objectToWorld, tg)); - ns = normalize(transformNormal(worldToObject, ns)); - - TBN tbn(tg, ns); // Calculate an otho-normal system respective to the shading normal. - - // Get the current rtPayload pointer from the unsigned int payload registers p0 and p1. - PerRayData* thePrd = mergePointer(optixGetPayload_0(), optixGetPayload_1()); - - thePrd->flags |= FLAG_HIT; // Required to distinguish surface hits from random walk miss. - - thePrd->distance = optixGetRayTmax(); // Return the current path segment distance, needed for absorption calculations in the integrator. + // Calculate an ortho-normal system respective to the shading normal. + float3 bt = normalize(cross(ns, tg)); + tg = cross(bt, ns); // Now the tangent is orthogonal to the shading normal. + + // The Mdl_state holds the texture attributes per texture space in separate arrays. + float3 texture_coordinates[NUM_TEXTURE_SPACES]; + float3 texture_tangents[NUM_TEXTURE_SPACES]; + float3 texture_bitangents[NUM_TEXTURE_SPACES]; + + // NUM_TEXTURE_SPACES is always at least 1. + texture_coordinates[0] = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + texture_bitangents[0] = bt; + texture_tangents[0] = tg; - // PRECISION Calculate this from the object space vertex positions and transform to world for better accuracy when needed. - // Same as: thePrd->pos = optixGetWorldRayOrigin() + optixGetWorldRayDirection() * optixGetRayTmax(); - thePrd->pos += thePrd->wi * thePrd->distance; - - // If we're inside a volume and hit something, the path throughput needs to be modulated - // with the transmittance along this segment before adding surface or light radiance! - if (0 < thePrd->idxStack) // This assumes the first stack entry is vaccuum. - { - thePrd->throughput *= expf(thePrd->sigma_t * -thePrd->distance); - - // Increment the volume scattering random walk counter. - // Unused when FLAG_VOLUME_SCATTERING is not set. - ++thePrd->walk; - } - +#if NUM_TEXTURE_SPACES == 2 + // HACK Simply copy the vertex attributes of texture space 0, simply because there is no second texcood inside TriangleAttributes. + texture_coordinates[1] = texture_coordinates[0]; + texture_bitangents[1] = bt; + texture_tangents[1] = tg; +#endif + + // Setup the Mdl_state. Mdl_state state; - float4 texture_results[16]; + float4 texture_results[NUM_TEXTURE_RESULTS]; // For explanations of these fields see comments inside __closesthit__radiance above. - state.normal = tbn.normal; + state.normal = ns; state.geom_normal = ng; state.position = thePrd->pos; state.animation_time = 0.0f; - state.text_coords = &tc; - state.tangent_u = &tbn.tangent; - state.tangent_v = &tbn.bitangent; + state.text_coords = texture_coordinates; + state.tangent_u = texture_tangents; + state.tangent_v = texture_bitangents; state.text_results = texture_results; state.ro_data_segment = nullptr; state.world_to_object = worldToObject; @@ -817,38 +857,62 @@ extern "C" __global__ void __anyhit__radiance_cutout() getTransforms(optixGetTransformListHandle(0), objectToWorld, worldToObject); // Single instance level transformation list only. + // Object space vertex attributes at the hit point. float3 po = attr0.vertex * alpha + attr1.vertex * theBarycentrics.x + attr2.vertex * theBarycentrics.y; + + float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; float3 ng = cross(attr1.vertex - attr0.vertex, attr2.vertex - attr0.vertex); - float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; - - const float3 tc = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - // Transform into internal space == world space. + // Transform attributes into internal space == world space. po = transformPoint(objectToWorld, po); + ns = normalize(transformNormal(worldToObject, ns)); ng = normalize(transformNormal(worldToObject, ng)); + // This is actually the geometry tangent which for the runtime generated geometry objects + // (plane, box, sphere, torus) match exactly with the texture space tangent. + // FIXME Generate these from the triangle's texture derivatives instead, but that's more expensive. + // Mind that tangents and bitangents are transformed as vectors, not normals, because they lie inside the surface's plane. tg = normalize(transformVector(objectToWorld, tg)); - ns = normalize(transformNormal(worldToObject, ns)); - - TBN tbn(tg, ns); // Calculate an otho-normal system respective to the shading normal. - + // Calculate an ortho-normal system respective to the shading normal. + // Expanding the TBN tbn(tg, ns) constructor because TBN members can't be used as pointers for the Mdl_state with NUM_TEXTURE_SPACES > 1. + float3 bt = normalize(cross(ns, tg)); + tg = cross(bt, ns); // Now the tangent is orthogonal to the shading normal. + + // The Mdl_state holds the texture attributes per texture space in separate arrays. + float3 texture_coordinates[NUM_TEXTURE_SPACES]; + float3 texture_tangents[NUM_TEXTURE_SPACES]; + float3 texture_bitangents[NUM_TEXTURE_SPACES]; + + // NUM_TEXTURE_SPACES is always at least 1. + texture_coordinates[0] = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + texture_bitangents[0] = bt; + texture_tangents[0] = tg; + +#if NUM_TEXTURE_SPACES == 2 + // HACK Simply copy the vertex attributes of texture space 0, simply because there is no second texcood inside TriangleAttributes. + texture_coordinates[1] = texture_coordinates[0]; + texture_bitangents[1] = bt; + texture_tangents[1] = tg; +#endif + + // Setup the Mdl_state. Mdl_state state; - float4 texture_results[16]; + float4 texture_results[NUM_TEXTURE_RESULTS]; // For explanations of these fields see comments inside __closesthit__radiance above. - state.normal = tbn.normal; + state.normal = ns; state.geom_normal = ng; state.position = po; state.animation_time = 0.0f; - state.text_coords = &tc; - state.tangent_u = &tbn.tangent; - state.tangent_v = &tbn.bitangent; + state.text_coords = texture_coordinates; + state.tangent_u = texture_tangents; + state.tangent_v = texture_bitangents; state.text_results = texture_results; state.ro_data_segment = nullptr; state.world_to_object = worldToObject; state.object_to_world = objectToWorld; - state.object_id = theData.ids.z; // idObject + state.object_id = theData.ids.z; state.meters_per_scene_unit = 1.0f; const MaterialDefinitionMDL& material = sysData.materialDefinitionsMDL[theData.ids.x]; @@ -917,38 +981,61 @@ extern "C" __global__ void __anyhit__shadow_cutout() // For the radiance ray typ getTransforms(optixGetTransformListHandle(0), objectToWorld, worldToObject); // Single instance level transformation list only. - float3 po = attr0.vertex * alpha + attr1.vertex * theBarycentrics.x + attr2.vertex * theBarycentrics.y; + // Object space vertex attributes at the hit point. + float3 po = attr0.vertex * alpha + attr1.vertex * theBarycentrics.x + attr2.vertex * theBarycentrics.y; + float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; float3 ng = cross(attr1.vertex - attr0.vertex, attr2.vertex - attr0.vertex); - float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - float3 ns = attr0.normal * alpha + attr1.normal * theBarycentrics.x + attr2.normal * theBarycentrics.y; - - const float3 tc = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + float3 tg = attr0.tangent * alpha + attr1.tangent * theBarycentrics.x + attr2.tangent * theBarycentrics.y; - // Transform into internal space == world space. + // Transform attributes into internal space == world space. po = transformPoint(objectToWorld, po); + ns = normalize(transformNormal(worldToObject, ns)); ng = normalize(transformNormal(worldToObject, ng)); + // This is actually the geometry tangent which for the runtime generated geometry objects + // (plane, box, sphere, torus) match exactly with the texture space tangent. + // FIXME Generate these from the triangle's texture derivatives instead, but that's more expensive. + // Mind that tangents and bitangents are transformed as vectors, not normals, because they lie inside the surface's plane. tg = normalize(transformVector(objectToWorld, tg)); - ns = normalize(transformNormal(worldToObject, ns)); - - TBN tbn(tg, ns); // Calculate an otho-normal system respective to the shading normal. - + // Calculate an ortho-normal system respective to the shading normal. + // Expanding the TBN tbn(tg, ns) constructor because TBN members can't be used as pointers for the Mdl_state with NUM_TEXTURE_SPACES > 1. + float3 bt = normalize(cross(ns, tg)); + tg = cross(bt, ns); // Now the tangent is orthogonal to the shading normal. + + // The Mdl_state holds the texture attributes per texture space in separate arrays. + float3 texture_coordinates[NUM_TEXTURE_SPACES]; + float3 texture_tangents[NUM_TEXTURE_SPACES]; + float3 texture_bitangents[NUM_TEXTURE_SPACES]; + + // NUM_TEXTURE_SPACES is always at least 1. + texture_coordinates[0] = attr0.texcoord * alpha + attr1.texcoord * theBarycentrics.x + attr2.texcoord * theBarycentrics.y; + texture_bitangents[0] = bt; + texture_tangents[0] = tg; + +#if NUM_TEXTURE_SPACES == 2 + // HACK Simply copy the vertex attributes of texture space 0, simply because there is no second texcood inside TriangleAttributes. + texture_coordinates[1] = texture_coordinates[0]; + texture_bitangents[1] = bt; + texture_tangents[1] = tg; +#endif + + // Setup the Mdl_state. Mdl_state state; - float4 texture_results[16]; + float4 texture_results[NUM_TEXTURE_RESULTS]; // For explanations of these fields see comments inside __closesthit__radiance above. - state.normal = tbn.normal; + state.normal = ns; state.geom_normal = ng; state.position = po; state.animation_time = 0.0f; - state.text_coords = &tc; - state.tangent_u = &tbn.tangent; - state.tangent_v = &tbn.bitangent; - state.text_results = texture_results; + state.text_coords = texture_coordinates; + state.tangent_u = texture_tangents; + state.tangent_v = texture_bitangents; + state.text_results = texture_results; state.ro_data_segment = nullptr; state.world_to_object = worldToObject; state.object_to_world = objectToWorld; - state.object_id = theData.ids.z; // idObject + state.object_id = theData.ids.z; state.meters_per_scene_unit = 1.0f; const MaterialDefinitionMDL& material = sysData.materialDefinitionsMDL[theData.ids.x]; @@ -1018,18 +1105,11 @@ extern "C" __device__ LightSample __direct_callable__light_mesh(const LightDefin const TriangleAttributes& attr1 = attributes[tri.y]; const TriangleAttributes& attr2 = attributes[tri.z]; - float3 po = attr0.vertex * alpha + attr1.vertex * beta + attr2.vertex * gamma; - float3 ng = cross(attr1.vertex - attr0.vertex, attr2.vertex - attr0.vertex); - float3 tg = attr0.tangent * alpha + attr1.tangent * beta + attr2.tangent * gamma; - float3 ns = attr0.normal * alpha + attr1.normal * beta + attr2.normal * gamma; - - const float3 tc = attr0.texcoord * alpha + attr1.texcoord * beta + attr2.texcoord * gamma; + // Object space vertex attributes at the hit point. + float3 po = attr0.vertex * alpha + attr1.vertex * beta + attr2.vertex * gamma; - // Transform into internal space == world space. + // Transform attributes into internal space == world space. po = transformPoint(light.matrix, po); - ng = normalize(transformNormal(light.matrixInv, ng)); - tg = normalize(transformVector(light.matrix, tg)); - ns = normalize(transformNormal(light.matrixInv, ns)); // Calculate the outgoing direction from light sample position to surface point. lightSample.direction = po - prd->pos; // Sample direction from surface point to light sample position. @@ -1041,21 +1121,48 @@ extern "C" __device__ LightSample __direct_callable__light_mesh(const LightDefin } lightSample.direction *= 1.0f / lightSample.distance; // Normalized vector from light sample position to surface point. - - TBN tbn(tg, ns); // Calculate an otho-normal system respective to the shading normal. + float3 ns = attr0.normal * alpha + attr1.normal * beta + attr2.normal * gamma; + float3 ng = cross(attr1.vertex - attr0.vertex, attr2.vertex - attr0.vertex); + float3 tg = attr0.tangent * alpha + attr1.tangent * beta + attr2.tangent * gamma; + + // Transform attributes into internal space == world space. + ns = normalize(transformNormal(light.matrixInv, ns)); + ng = normalize(transformNormal(light.matrixInv, ng)); + tg = normalize(transformVector(light.matrix, tg)); + + float3 bt = normalize(cross(ns, tg)); + tg = cross(bt, ns); // Now the tangent is orthogonal to the shading normal. + + // The Mdl_state holds the texture attributes per texture space in separate arrays. + float3 texture_coordinates[NUM_TEXTURE_SPACES]; + float3 texture_tangents[NUM_TEXTURE_SPACES]; + float3 texture_bitangents[NUM_TEXTURE_SPACES]; + + // NUM_TEXTURE_SPACES is always at least 1. + texture_coordinates[0] = attr0.texcoord * alpha + attr1.texcoord * beta + attr2.texcoord * gamma; + texture_bitangents[0] = bt; + texture_tangents[0] = tg; + +#if NUM_TEXTURE_SPACES == 2 + // HACK Simply copy the vertex attributes of texture space 0, simply because there is no second texcood inside TriangleAttributes. + texture_coordinates[1] = texture_coordinates[0]; + texture_bitangents[1] = bt; + texture_tangents[1] = tg; +#endif + + // Setup the Mdl_state. Mdl_state state; - float4 texture_results[16]; + float4 texture_results[NUM_TEXTURE_RESULTS]; - // For explanations of these fields see comments inside __closesthit__radiance above. - state.normal = tbn.normal; + state.normal = ns; state.geom_normal = ng; state.position = po; state.animation_time = 0.0f; - state.text_coords = &tc; - state.tangent_u = &tbn.tangent; - state.tangent_v = &tbn.bitangent; + state.text_coords = texture_coordinates; + state.tangent_u = texture_tangents; + state.tangent_v = texture_bitangents; state.text_results = texture_results; state.ro_data_segment = nullptr; state.world_to_object = light.matrixInv; @@ -1164,3 +1271,254 @@ extern "C" __device__ LightSample __direct_callable__light_mesh(const LightDefin return lightSample; } + + +extern "C" __global__ void __closesthit__curves() +{ + // Get the current rtPayload pointer from the unsigned int payload registers p0 and p1. + PerRayData* thePrd = mergePointer(optixGetPayload_0(), optixGetPayload_1()); + + thePrd->flags |= FLAG_HIT; // Required to distinguish surface hits from random walk miss. + + thePrd->distance = optixGetRayTmax(); // Return the current path segment distance, needed for absorption calculations in the integrator. + + // Note that no adjustment to the hit position is done here! + // The OptiX curve primitives are not intersecting with backfaces, which means the sceneEpsilon + // offset on the continuation ray t_min is enough to prevent self-intersections independently + // of the continuation ray direction being a reflection or transmisssion. + //thePrd->pos = optixGetWorldRayOrigin() + optixGetWorldRayDirection() * optixGetRayTmax(); + thePrd->pos += thePrd->wi * thePrd->distance; + + // If we're inside a volume and hit something, the path throughput needs to be modulated + // with the transmittance along this segment before adding surface or light radiance! + if (0 < thePrd->idxStack) // This assumes the first stack entry is vaccuum. + { + thePrd->throughput *= expf(thePrd->sigma_t * -thePrd->distance); + + // Increment the volume scattering random walk counter. + // Unused when FLAG_VOLUME_SCATTERING is not set. + ++thePrd->walk; + } + + float4 objectToWorld[3]; + float4 worldToObject[3]; + + getTransforms(optixGetTransformListHandle(0), objectToWorld, worldToObject); // Single instance level transformation list only. + + GeometryInstanceData theData = sysData.geometryInstanceData[optixGetInstanceId()]; + + const unsigned int thePrimitiveIndex = optixGetPrimitiveIndex(); + + const unsigned int* indices = reinterpret_cast(theData.indices); + const unsigned int index = indices[thePrimitiveIndex]; + + const CurveAttributes* attributes = reinterpret_cast(theData.attributes); + + float4 spline[4]; + + spline[0] = attributes[index ].vertex; + spline[1] = attributes[index + 1].vertex; + spline[2] = attributes[index + 2].vertex; + spline[3] = attributes[index + 3].vertex; + + // Fixed bitangent-like reference vector for the vFiber calculation. + float3 rf = make_float3(attributes[index].reference); + + const float4 t = make_float4(attributes[index ].texcoord.w, + attributes[index + 1].texcoord.w, + attributes[index + 2].texcoord.w, + attributes[index + 3].texcoord.w); + + CubicInterpolator interpolator; + + interpolator.initializeFromBSpline(spline); + + // Convenience function for __uint_as_float( optixGetAttribute_0() ); + const float u = optixGetCurveParameter(); + + // Optimized (see curve.h): + const float o_s = 1.0f / 6.0f; + const float ts0 = (t.w - t.x) * o_s + (t.y - t.z) * 0.5f; + const float ts1 = (t.x + t.z) * 0.5f - t.y; + const float ts2 = (t.z - t.x) * 0.5f; + const float ts3 = (t.x + t.y * 4.0f + t.z) * o_s; + + // Coordinate in range [0.0f, 1.0f] from root to tip along the whole fiber. + const float uFiber = ( ( ( ts0 * u ) + ts1 ) * u + ts2 ) * u + ts3; + + const float4 fiberPosition = interpolator.position4(u); // .xyz = object space position, .w = radius of the curve's center line at the interpolant u. + + float3 tg = interpolator.velocity3(u); // Unnormalized object space tangent along the fiber from root to tip at the interpolant u. + + float3 position = transformPoint(worldToObject, thePrd->pos); // Transform to object space. + float3 ns = surfaceNormal(interpolator, u, position); // This moves position onto the surface of the curve in object space. + + // Transform into renderer internal space == world space. + rf = normalize(transformVector(objectToWorld, rf)); + tg = normalize(transformVector(objectToWorld, tg)); + ns = normalize(transformNormal(worldToObject, ns)); + + // Calculate an ortho-normal system for the fiber. The tangent is the fixed vector here! + // Expand the FBN fbn(tg, ns) constructor because that cannot be used with NUM_TEXTURE_SPACES > 1. + // Tangent is kept, bitangent and normal are calculated. Tangent must be normalized. + float3 bt = normalize(cross(ns, tg)); + ns = cross(tg, bt); // Now the normal is orthogonal to the fiber tangent. + + // Transform the constant reference fiber bitangent vector into the local fiber coordinate system. + // This makes the coordinate [0.0f, 1.0f] around the cross section of the hair relative to the hit point. + // Expanded proj = fbn.transformToLocal(rf); + // Only need the projected y- and z-components.(This works without normalization. + const float2 proj = make_float2(dot(rf, bt), dot(rf, ns)); + + // The (bitangent, normal) plane contains the cross section of the intersection. + // I want vFiber go from 0.0f to 1.0f counter-clockwise around the fiber when looking from the fiber root along the fiber. + // As texture coordinate this means lower-left origin texture coordinates like in OpenGL. + const float vFiber = (atan2f(-proj.x, proj.y) + M_PIf) * 0.5f * M_1_PIf; + + // The Mdl_state holds the texture attributes per texture space in separate arrays. + // NUM_TEXTURE_SPACES is always at least 1. + float3 texture_coordinates[NUM_TEXTURE_SPACES]; + float3 texture_tangents[NUM_TEXTURE_SPACES]; + float3 texture_bitangents[NUM_TEXTURE_SPACES]; + + // In hair shading, texture spaces contain the following values: + // texture_coordinate(0).x: The normalized position of the intersection point along the hair fiber in the range from zero for the root of the fiber to one for the tip of the fiber. + // texture_coordinate(0).y: The normalized position of the intersection point around the hair fiber in the range from zero to one. + // texture_coordinate(0).z: The thickness of the hair fiber at the intersection point in internal space. + texture_coordinates[0] = make_float3(uFiber, vFiber, fiberPosition.w * 2.0f); // .z = thickness = radius * 2.0f + texture_bitangents[0] = bt; + texture_tangents[0] = tg; + +#if NUM_TEXTURE_SPACES == 2 + // PERF Only ever set NUM_TEXTURE_SPACES to 2 if you need the hair BSDF to fully work, + // texture_coordinate(1): A position of the root of the hair fiber, for example, from a texture space of a surface supporting the hair fibers. This position is constant for a fiber. + // Fixed texture coordinate for the fiber. Only loaded when NUM_TEXTURE_SPACES == 2. + texture_coordinates[1] = make_float3(attributes[index].texcoord); + texture_bitangents[1] = bt; // HACK Just copy the values from the first entry. + texture_tangents[1] = tg; +#endif + + Mdl_state state; + + float4 texture_results[NUM_TEXTURE_RESULTS]; + + state.normal = ns; + state.geom_normal = ns; + state.position = thePrd->pos; + state.animation_time = 0.0f; + state.text_coords = texture_coordinates; + state.tangent_u = texture_tangents; + state.tangent_v = texture_bitangents; + state.text_results = texture_results; + state.ro_data_segment = nullptr; + state.world_to_object = worldToObject; + state.object_to_world = objectToWorld; + state.object_id = theData.ids.z; + state.meters_per_scene_unit = 1.0f; + + const MaterialDefinitionMDL& material = sysData.materialDefinitionsMDL[theData.ids.x]; + + mi::neuraylib::Resource_data res_data = { nullptr, material.texture_handler }; + + const DeviceShaderConfiguration& shaderConfiguration = sysData.shaderConfigurations[material.indexShader]; + + // This is always present, even if it just returns. + optixDirectCall(shaderConfiguration.idxCallInit, &state, &res_data, nullptr, material.arg_block); + + // Start fresh with the next BSDF sample. + // Save the current path throughput for the direct lighting contribution. + // The path throughput will be modulated with the BSDF sampling results before that. + const float3 throughput = thePrd->throughput; + // The pdf of the previous event was needed for the emission calculation above. + thePrd->pdf = 0.0f; + + // Importance sample the hair BSDF. + if (0 <= shaderConfiguration.idxCallHairSample) + { + mi::neuraylib::Bsdf_sample_data sample_data; + + int idx = thePrd->idxStack; + + // FIXME The MDL-SDK libbsdf_hair.h ignores these ior values and only uses the ior value from the chiang_hair_bsdf structure! + sample_data.ior1 = thePrd->stack[idx].ior; // From surrounding medium ior + sample_data.ior2.x = MI_NEURAYLIB_BSDF_USE_MATERIAL_IOR; // to material ior. + sample_data.k1 = thePrd->wo; // == -optixGetWorldRayDirection() + sample_data.xi = rng4(thePrd->seed); + + optixDirectCall(shaderConfiguration.idxCallHairSample, &sample_data, &state, &res_data, nullptr, material.arg_block); + + thePrd->wi = sample_data.k2; // Continuation direction. + thePrd->throughput *= sample_data.bsdf_over_pdf; // Adjust the path throughput for all following incident lighting. + thePrd->pdf = sample_data.pdf; // Specular events return pdf == 0.0f! + thePrd->eventType = sample_data.event_type; // This replaces the previous PRD flags. + } + else + { + // If there is no valid scattering BSDF, it's the black bsdf() which ends the path. + // This is usually happening with arbitrary mesh lights which only specify emission. + thePrd->eventType = mi::neuraylib::BSDF_EVENT_ABSORB; + // None of the following code will have any effect in that case. + return; + } + + // Direct lighting if the sampled BSDF was diffuse and any light is in the scene. + const int numLights = sysData.numLights; + + if (sysData.directLighting && 0 < numLights && (thePrd->eventType & (mi::neuraylib::BSDF_EVENT_DIFFUSE | mi::neuraylib::BSDF_EVENT_GLOSSY))) + { + // Sample one of many lights. + // The caller picks the light to sample. Make sure the index stays in the bounds of the sysData.lightDefinitions array. + const int indexLight = (1 < numLights) ? clamp(static_cast(floorf(rng(thePrd->seed) * numLights)), 0, numLights - 1) : 0; + + const LightDefinition& light = sysData.lightDefinitions[indexLight]; + + LightSample lightSample = optixDirectCall(NUM_LENS_TYPES + light.typeLight, light, thePrd); + + if (0.0f < lightSample.pdf && 0 <= shaderConfiguration.idxCallHairEval) // Useful light sample and valid shader? + { + mi::neuraylib::Bsdf_evaluate_data eval_data; + + int idx = thePrd->idxStack; + + // DAR FIXME The MDL-SDK libbsdf_hair.h ignores these values and only uses the ior value from the chiang_hair-bsdf structure! + eval_data.ior1 = thePrd->stack[idx].ior; // From surrounding medium ior + eval_data.ior2.x = MI_NEURAYLIB_BSDF_USE_MATERIAL_IOR; // to material ior. + eval_data.k1 = thePrd->wo; + eval_data.k2 = lightSample.direction; + + optixDirectCall(shaderConfiguration.idxCallHairEval, &eval_data, &state, &res_data, nullptr, material.arg_block); + + // DAR DEBUG This already contains the fabsf(dot(lightSample.direction, state.normal)) factor! + // For a white Lambert material, the bxdf components match the eval_data.pdf + const float3 bxdf = eval_data.bsdf_diffuse + eval_data.bsdf_glossy; + + if (0.0f < eval_data.pdf && isNotNull(bxdf)) + { + // Pass the current payload registers through to the shadow ray. + unsigned int p0 = optixGetPayload_0(); + unsigned int p1 = optixGetPayload_1(); + + thePrd->flags &= ~FLAG_SHADOW; // Clear the shadow flag. + + // Note that the sysData.sceneEpsilon is applied on both sides of the shadow ray [t_min, t_max] interval + // to prevent self-intersections with the actual light geometry in the scene. + optixTrace(sysData.topObject, + thePrd->pos, lightSample.direction, // origin, direction + sysData.sceneEpsilon, lightSample.distance - sysData.sceneEpsilon, 0.0f, // tmin, tmax, time + OptixVisibilityMask(0xFF), OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT, // The shadow ray type only uses anyhit programs. + TYPE_RAY_SHADOW, NUM_RAY_TYPES, TYPE_RAY_SHADOW, + p0, p1); // Pass through thePrd to the shadow ray. + + if ((thePrd->flags & FLAG_SHADOW) == 0) // Shadow flag not set? + { + const float weightMIS = (TYPE_LIGHT_POINT <= light.typeLight) ? 1.0f : balanceHeuristic(lightSample.pdf, eval_data.pdf); + + // The sampled emission needs to be scaled by the inverse probability to have selected this light, + // Selecting one of many lights means the inverse of 1.0f / numLights. + // This is using the path throughput before the sampling modulated it above. + thePrd->radiance += throughput * bxdf * lightSample.radiance_over_pdf * (float(numLights) * weightMIS); + } + } + } + } +} diff --git a/apps/MDL_renderer/shaders/shader_configuration.h b/apps/MDL_renderer/shaders/shader_configuration.h index a8cf72b7..61bfb02f 100644 --- a/apps/MDL_renderer/shaders/shader_configuration.h +++ b/apps/MDL_renderer/shaders/shader_configuration.h @@ -71,6 +71,9 @@ struct DeviceShaderConfiguration int idxCallGeometryCutoutOpacity; + int idxCallHairSample; + int idxCallHairEval; + // The constant expression values: //bool thin_walled; // Stored inside flags. float3 surface_intensity; diff --git a/apps/MDL_renderer/src/Application.cpp b/apps/MDL_renderer/src/Application.cpp index caf8b3a0..85faa0ca 100644 --- a/apps/MDL_renderer/src/Application.cpp +++ b/apps/MDL_renderer/src/Application.cpp @@ -1843,6 +1843,45 @@ bool Application::loadSceneDescription(const std::string& filename) appendInstance(m_scene, geometry, cur.matrix, indexMaterial, -1); } + else if (token == "hair") // model hair scale material filename + { + // This scales the thickness defined inside the hair files. + tokenType = parser.getNextToken(token); + MY_ASSERT(tokenType == PTT_VAL); + const float scale = float(atof(token.c_str())); + + std::string nameMaterialReference; + tokenType = parser.getNextToken(nameMaterialReference); + + std::string filename; + tokenType = parser.getNextToken(filename); + MY_ASSERT(tokenType == PTT_STRING); + convertPath(filename); + + std::ostringstream keyGeometry; + keyGeometry << "hair_" << scale << "_" << filename; + + std::shared_ptr geometry; + + std::map::const_iterator itg = m_mapGeometries.find(keyGeometry.str()); + if (itg == m_mapGeometries.end()) + { + m_mapGeometries[keyGeometry.str()] = m_idGeometry; + + geometry = std::make_shared(m_idGeometry++); + geometry->createHair(filename, scale); + + m_geometries.push_back(geometry); + } + else + { + geometry = std::dynamic_pointer_cast(m_geometries[itg->second]); + } + + const int indexMaterial = findMaterial(nameMaterialReference); + + appendInstance(m_scene, geometry, cur.matrix, indexMaterial, -1); + } else if (token == "assimp") { std::string filenameModel; @@ -2162,6 +2201,15 @@ void Application::traverseGraph(std::shared_ptr node, InstanceData& in } } break; + + case sg::NodeType::NT_CURVES: + { + std::shared_ptr geometry = std::dynamic_pointer_cast(node); + + instanceData.idGeometry = geometry->getId(); + instanceData.idLight = -1; // No support for emissive curves. + } + break; } } @@ -2557,14 +2605,14 @@ void Application::calculateTangents(std::vector& attributes, normal.y = attributes[i].normal.y; normal.z = attributes[i].normal.z; - if (0.001f < 1.0f - fabsf(dot(normal, tangent))) + if (DENOMINATOR_EPSILON < 1.0f - fabsf(dot(normal, tangent))) { bitangent = normalize(cross(normal, tangent)); tangent = normalize(cross(bitangent, normal)); } else // Normal and tangent direction too collinear. { - MY_ASSERT(0.001f < 1.0f - fabsf(dot(bitangent, normal))); + MY_ASSERT(DENOMINATOR_EPSILON < 1.0f - fabsf(dot(bitangent, normal))); tangent = normalize(cross(bitangent, normal)); //bitangent = normalize(cross(normal, tangent)); } diff --git a/apps/MDL_renderer/src/Assimp.cpp b/apps/MDL_renderer/src/Assimp.cpp index dce54172..4e8d52b9 100644 --- a/apps/MDL_renderer/src/Assimp.cpp +++ b/apps/MDL_renderer/src/Assimp.cpp @@ -173,7 +173,7 @@ std::shared_ptr Application::createASSIMP(const std::string& filename if (mesh->HasTangentsAndBitangents()) { const aiVector3D& t = mesh->mTangents[iVertex]; - attrib.tangent = make_float3(t.x, t.y, t.z); + attrib.tangent = normalize(make_float3(t.x, t.y, t.z)); } else { @@ -184,7 +184,8 @@ std::shared_ptr Application::createASSIMP(const std::string& filename if (mesh->HasNormals()) { const aiVector3D& n = mesh->mNormals[iVertex]; - attrib.normal = make_float3(n.x, n.y, n.z); + // There exist OBJ files with unnormalized normals! + attrib.normal = normalize(make_float3(n.x, n.y, n.z)); } else { diff --git a/apps/MDL_renderer/src/Curves.cpp b/apps/MDL_renderer/src/Curves.cpp new file mode 100644 index 00000000..ab1194db --- /dev/null +++ b/apps/MDL_renderer/src/Curves.cpp @@ -0,0 +1,342 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +#include "inc/SceneGraph.h" +#include "inc/MyAssert.h" + +#include "inc/Hair.h" + +#include "dp/math/math.h" + +#include "shaders/vector_math.h" +#include "shaders/shader_common.h" + + +#include +#include +#include +#include + + +namespace sg +{ + +static float3 cubeProjection(const float3 r) +{ + // Spherical projection was pretty unpredictable with hair models. + // Try a less distorted cubemap projection instead. + // Note that each of the faces is the same 2D texture though. + + // See OpenGL 4.6 specs chapter "8.13 Cube Map Texture Selection". + const float x = fabsf(r.x); + const float y = fabsf(r.y); + const float z = fabsf(r.z); + + float ma = 0.0f; + float sc = 0.0f; + float tc = 0.0f; + + if (x >= y && x >= z) + { + ma = x; // major axis rx + if (r.x >= 0.0f) + { + // major axis +rx + sc = -r.z; + tc = -r.y; + } + else + { + // major axis -rx + sc = r.z; + tc = -r.y; + } + } + else if (y >= z) + { + ma = y; // major axis ry + if (r.y >= 0.0f) + { + // major axis +ry + sc = r.x; + tc = r.z; + } + else + { + // major axis -ry + sc = r.x; + tc = -r.z; + } + } + else + { + ma = z; // major axis rz + if (r.z >= 0.0f) + { + // major axis +rz + sc = r.x; + tc = -r.y; + } + else + { + // major axis -rz + sc = -r.x; + tc = -r.y; + } + } + + const float s = 0.5f * (sc / ma + 1.0f); + const float t = 0.5f * (tc / ma + 1.0f); + + return make_float3(s, t, 0.0f); +} + + +bool Curves::createHair(std::string const& filename, const float scale) +{ + // Note that hair files usually have a z-up coordinate system and they seem to be defined in centimeters. + // That can be adjusted inside the scene description with a scale and rotate transform. + // The "scale" parameter coming from the "model hair scale material filename" option + // is modulating the thickness parameter defined inside the the hair file. + // Use scale == 1.0 to get the original thickness. + + // push + // scale 0.01 0.01 0.01 + // rotate 1 0 0 -90 + // model hair 1.0 bsdf_hair "file.hair" + // pop + + Hair hair; + + if (!hair.load(filename)) + { + return false; + } + + // Iterate over all strands and build the curve attributes for the scene graph. + const unsigned int numStrands = hair.getNumStrands(); + + // Calculate a texture coordinate for each strand. + std::vector texcoords; + texcoords.resize(numStrands); + + // Simply do some spherical mapping to the center of the root points. + // Calculate the center of the root points. + float3 center = make_float3(0.0f); + + // This variable is always the running index over the segments. + unsigned int idx = 0; // Set to root index of first strand. + + for (unsigned int strand = 0; strand < numStrands; ++strand) + { + const unsigned short numSegments = hair.getNumSegments(strand); + + if (numSegments == 0) + { + continue; + } + + center += hair.getPoint(idx); + + idx += numSegments + 1; // Advance to next strand's root point. + } + + // Center of mass of the root strand points. + center /= float(numStrands); + + idx = 0; // Reset to root index of first strand. + + for (unsigned int strand = 0; strand < numStrands; ++strand) + { + const unsigned short numSegments = hair.getNumSegments(strand); + + if (numSegments == 0) + { + continue; + } + + const float3 r = normalize(hair.getPoint(idx) - center); + + texcoords[strand] = cubeProjection(r); + + idx += numSegments + 1; // Advance to next strand's root point. + } + + // The idxRoot value is always the root point index of the current strand. + unsigned int idxRoot = 0; // Set to root point of the first strand. + + for (unsigned int strand = 0; strand < numStrands; ++strand) + { + const unsigned short numSegments = hair.getNumSegments(strand); + + // If there is ever a strand defintion with zero segments, + // just skip that and don't change any of the indices. + if (numSegments == 0) + { + continue; + } + + // Calculate the length of each strand. + // Linear length along the control points, not the actual cubic curve. + // Needed for uFiber value in state::texture_coordinate(0).x + float lengthStrand = 0.0f; + + // Calculate some fixed reference vector per strand. + // Needed for vFiber value inside state::texture_coordinate(0).y + // It's calculated as a "face normal" of the control points "polygon", which results in something like a + // fixed bitangent to the fiber tangent direction which works OK for usual hair definitions. + float3 reference = make_float3(0.0f, 0.0f, 0.0f); + + idx = idxRoot; // Start local running index over the current strand. + + float3 v0 = hair.getPoint(idx); // Root point of the strand. + float3 v1; + + for (unsigned short segment = 0; segment < numSegments; ++segment) + { + v1 = hair.getPoint(idx + 1); + + lengthStrand += length(v1 - v0); + + // Interpret the hair control points as a polygon and calculate a face normal of that. + // The face normal is proportional to the projected surface of the polygon onto the ortho-normal basis planes. + reference.x += (v0.y - v1.y) * (v0.z + v1.z); + reference.y += (v0.z - v1.z) * (v0.x + v1.x); + reference.z += (v0.x - v1.x) * (v0.y + v1.y); + + // Advance to next segment. + v0 = v1; + ++idx; + } + + // v0 contains the endpoint of the last segment here. + // Close the "polygon" to the root point. + v1 = hair.getPoint(idxRoot); // First point of the strand. + + reference.x += (v0.y - v1.y) * (v0.z + v1.z); + reference.y += (v0.z - v1.z) * (v0.x + v1.x); + reference.z += (v0.x - v1.x) * (v0.y + v1.y); + + // If the Control points are not building some polygon face (maybe a straight line with no area), + // just pick some orthogonal vector to a strand "tangent", the vector from root to tip control point. + if (reference.x == 0.0f && reference.y == 0.0f && reference.z == 0.0f) + { + float3 tangent = hair.getPoint(idxRoot + hair.getNumSegments(strand) + 1) - hair.getPoint(idxRoot); + + // Generate an orthogonal vector to the reference tangent. + reference = (fabsf(tangent.z) < fabsf(tangent.x)) + ? make_float3(tangent.z, 0.0f, -tangent.x) + : make_float3(0.0f, tangent.z, -tangent.y); + } + reference = normalize(reference); + + // Build the cubic B-spline data. + CurveAttributes attrib; + + // Remember the attribute start index of the B-Spline attributes building this strand. + unsigned int index = static_cast(m_attributes.size()); + + idx = idxRoot; // Start local running index over the current strand again. + + // Start point, radius and fiber interpolant. + float3 p0 = hair.getPoint(idx); // Start point of this curve segment. + float r0 = hair.getThickness(idx) * 0.5f * scale; // radius = thickness * 0.5f. The scale allows modulating modulate the hair thickness in the file. + float u0 = 0.0f; // Interpolant along the hair strand from 0.0f at the root to 1.0 at the tip. + + // Initialize to keep the compiler happy. + float3 p1 = p0; + float r1 = r0; + float u1 = u0; + + for (unsigned short segment = 0; segment < numSegments; ++segment) + { + // End point, radius and fiber interpolant. + p1 = hair.getPoint(idx + 1); + r1 = hair.getThickness(idx + 1) * 0.5f * scale; + u1 = u0 + length(p1 - p0) / lengthStrand; + + if (segment == 0) + { + // Push an additional phantom point before the hair control points + // to let the cubic B-spline start exactly at the first control point. + attrib.vertex = make_float4(p0 + (p0 - p1), std::max(0.0f, r0 + (r0 - r1))); + attrib.reference = make_float4(reference, 0.0f); + attrib.texcoord = make_float4(texcoords[strand], u0 + (u0 - u1)); + m_attributes.push_back(attrib); + } + + // Push the start point of this segment. + attrib.vertex = make_float4(p0, r0); + attrib.reference = make_float4(reference, 0.0f); + attrib.texcoord = make_float4(texcoords[strand], u0); + m_attributes.push_back(attrib); + + // The last segment will store the strand endpoint and append another phantom control point. + if (segment + 1 == numSegments) + { + // Push the end point of the last segment. + attrib.vertex = make_float4(p1, r1); + attrib.reference = make_float4(reference, 0.0f); + attrib.texcoord = make_float4(texcoords[strand], u1); + m_attributes.push_back(attrib); + + // Push an additional phantom point after the hair control points + // to let the cubic B-spline end exactly at the last control point. + attrib.vertex = make_float4(p1 + (p1 - p0), std::max(0.0f, r1 + (r1 - r0))); + attrib.reference = make_float4(reference, 0.0f); + attrib.texcoord = make_float4(texcoords[strand], u1 + (u1 - u0)); + m_attributes.push_back(attrib); + } + + // Let the end point become the new start point, except for the last segment + p0 = p1; + r0 = r1; + u0 = u1; + + ++idx; // Advance to the next segment's start point. + } + + // Generate the indices for this strand. + const unsigned int indexEnd = static_cast(m_attributes.size()) - 3; // Cubic B-spline curve uses 4 control points for each primitive. + + while (index < indexEnd) + { + m_indices.push_back(index); + ++index; + } + + // Done with one strand's control points. + // (When we reach this numSegments != 0.) + + idxRoot += numSegments + 1; // Skip over all control points of the current strand. + } + + return true; +} + +} // namespace sg diff --git a/apps/MDL_renderer/src/Device.cpp b/apps/MDL_renderer/src/Device.cpp index 9d4965d4..380fe5ac 100644 --- a/apps/MDL_renderer/src/Device.cpp +++ b/apps/MDL_renderer/src/Device.cpp @@ -593,7 +593,6 @@ void Device::initDeviceAttributes() CU_CHECK( cuDeviceGetAttribute(&m_deviceAttribute.concurrentManagedAccess, CU_DEVICE_ATTRIBUTE_CONCURRENT_MANAGED_ACCESS, m_ordinal) ); CU_CHECK( cuDeviceGetAttribute(&m_deviceAttribute.computePreemptionSupported, CU_DEVICE_ATTRIBUTE_COMPUTE_PREEMPTION_SUPPORTED, m_ordinal) ); CU_CHECK( cuDeviceGetAttribute(&m_deviceAttribute.canUseHostPointerForRegisteredMem, CU_DEVICE_ATTRIBUTE_CAN_USE_HOST_POINTER_FOR_REGISTERED_MEM, m_ordinal) ); - CU_CHECK( cuDeviceGetAttribute(&m_deviceAttribute.canUseStreamMemOps, CU_DEVICE_ATTRIBUTE_CAN_USE_STREAM_MEM_OPS, m_ordinal) ); CU_CHECK( cuDeviceGetAttribute(&m_deviceAttribute.canUse64BitStreamMemOps, CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS, m_ordinal) ); CU_CHECK( cuDeviceGetAttribute(&m_deviceAttribute.canUseStreamWaitValueNor, CU_DEVICE_ATTRIBUTE_CAN_USE_STREAM_WAIT_VALUE_NOR, m_ordinal) ); CU_CHECK( cuDeviceGetAttribute(&m_deviceAttribute.cooperativeLaunch, CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH, m_ordinal) ); @@ -683,6 +682,16 @@ void Device::initPipeline() OPTIX_CHECK( m_api.optixModuleCreateFromPTX(m_optixContext, &m_mco, &m_pco, programData.data(), programData.size(), nullptr, nullptr, &modules[i]) ); } + // Get the OptiX internal module with the intersection program for cubic B-spline curves; + OptixBuiltinISOptions builtinISOptions = {}; + + builtinISOptions.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE; + builtinISOptions.usesMotionBlur = 0; + + OptixModule moduleIntersectionCubicCurves; + + OPTIX_CHECK( m_api.optixBuiltinISModuleGet(m_optixContext, &m_mco, &m_pco, &builtinISOptions, &moduleIntersectionCubicCurves) ); + std::vector programGroupDescriptions(NUM_PROGRAM_GROUP_IDS); memset(programGroupDescriptions.data(), 0, sizeof(OptixProgramGroupDesc) * programGroupDescriptions.size()); @@ -789,6 +798,22 @@ void Device::initPipeline() pgd->hitgroup.moduleAH = modules[MODULE_ID_HIT]; pgd->hitgroup.entryFunctionNameAH = "__anyhit__shadow_cutout"; + pgd = &programGroupDescriptions[PGID_HIT_CURVES]; + pgd->kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; + pgd->flags = OPTIX_PROGRAM_GROUP_FLAGS_NONE; + pgd->hitgroup.moduleCH = modules[MODULE_ID_HIT]; + pgd->hitgroup.entryFunctionNameCH = "__closesthit__curves"; + pgd->hitgroup.moduleIS = moduleIntersectionCubicCurves; + pgd->hitgroup.entryFunctionNameIS = nullptr; // Uses built-in IS for cubic curves. + + pgd = &programGroupDescriptions[PGID_HIT_CURVES_SHADOW]; + pgd->kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; + pgd->flags = OPTIX_PROGRAM_GROUP_FLAGS_NONE; + pgd->hitgroup.moduleAH = modules[MODULE_ID_HIT]; + pgd->hitgroup.entryFunctionNameAH = "__anyhit__shadow"; + pgd->hitgroup.moduleIS = moduleIntersectionCubicCurves; + pgd->hitgroup.entryFunctionNameIS = nullptr; // Uses built-in IS for cubic curves. + // CALLABLES // Lens Shader pgd = &programGroupDescriptions[PGID_LENS_PINHOLE]; @@ -918,7 +943,7 @@ void Device::initPipeline() m_sbt.hitgroupRecordBase = m_d_sbtRecordHeaders + sizeof(SbtRecordHeader) * PGID_HIT_RADIANCE_0; m_sbt.hitgroupRecordStrideInBytes = (unsigned int) sizeof(SbtRecordHeader); - m_sbt.hitgroupRecordCount = NUM_RAY_TYPES * 4; // Four hitRecords: (no emission, emission) x (no cutout, cutout) + m_sbt.hitgroupRecordCount = NUM_RAY_TYPES * 5; // Five hitRecords: 0 to 3 == (no emission, emission) x (no cutout, cutout), and 4 = opaque cubic curves. m_sbt.callablesRecordBase = m_d_sbtRecordHeaders + sizeof(SbtRecordHeader) * PGID_LENS_PINHOLE; // The pinhole lens shader is the first callable. m_sbt.callablesRecordStrideInBytes = (unsigned int) sizeof(SbtRecordHeader); @@ -1359,6 +1384,122 @@ GeometryData Device::createGeometry(std::shared_ptr geometry) } +GeometryData Device::createGeometry(std::shared_ptr geometry) +{ + activateContext(); + synchronizeStream(); + + GeometryData data; + + data.primitiveType = PT_CURVES; + data.owner = m_index; + + const std::vector& attributes = geometry->getAttributes(); + const std::vector& indices = geometry->getIndices(); + + const size_t attributesSizeInBytes = sizeof(CurveAttributes) * attributes.size(); + + data.d_attributes = memAlloc(attributesSizeInBytes, 16); + data.numAttributes = attributes.size(); + + CU_CHECK( cuMemcpyHtoDAsync(data.d_attributes, attributes.data(), attributesSizeInBytes, m_cudaStream) ); + + CUdeviceptr d_radii = data.d_attributes + sizeof(float3); // Pointer to the radius in the .w component of the float4 vertex attribute + + const size_t indicesSizeInBytes = sizeof(unsigned int) * indices.size(); + + data.d_indices = memAlloc(indicesSizeInBytes, sizeof(unsigned int)); + data.numIndices = indices.size(); + + CU_CHECK( cuMemcpyHtoDAsync(data.d_indices, indices.data(), indicesSizeInBytes, m_cudaStream) ); + + OptixBuildInput buildInput = {}; + + buildInput.type = OPTIX_BUILD_INPUT_TYPE_CURVES; + + buildInput.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE; + buildInput.curveArray.numPrimitives = static_cast(indices.size()); + buildInput.curveArray.vertexBuffers = &data.d_attributes; + buildInput.curveArray.numVertices = static_cast(attributes.size()); + buildInput.curveArray.vertexStrideInBytes = sizeof(CurveAttributes); + buildInput.curveArray.widthBuffers = &d_radii; + buildInput.curveArray.widthStrideInBytes = sizeof(CurveAttributes); + buildInput.curveArray.normalBuffers = nullptr; // Reserved for future use + buildInput.curveArray.normalStrideInBytes = 0; // Reserved for future use + buildInput.curveArray.indexBuffer = data.d_indices; + buildInput.curveArray.indexStrideInBytes = sizeof(unsigned int); + buildInput.curveArray.flag = OPTIX_GEOMETRY_FLAG_NONE; // Only one flag because Curves have only one SBT entry. + buildInput.curveArray.primitiveIndexOffset = 0; + + OptixAccelBuildOptions accelBuildOptions = {}; + + accelBuildOptions.buildFlags = OPTIX_BUILD_FLAG_ALLOW_COMPACTION; + if (m_count == 1) + { + // PERF Enable OPTIX_BUILD_FLAG_PREFER_FAST_TRACE on single-GPU only. + // Note that OPTIX_BUILD_FLAG_PREFER_FAST_TRACE will use more memory, + // which performs worse when sharing across the NVLINK bridge which is much slower than VRAM accesses. + // This means comparisons between single-GPU and multi-GPU are not doing exactly the same! + accelBuildOptions.buildFlags |= OPTIX_BUILD_FLAG_PREFER_FAST_TRACE; + } + accelBuildOptions.operation = OPTIX_BUILD_OPERATION_BUILD; + + OptixAccelBufferSizes accelBufferSizes; + + OPTIX_CHECK( m_api.optixAccelComputeMemoryUsage(m_optixContext, &accelBuildOptions, &buildInput, 1, &accelBufferSizes) ); + + data.d_gas = memAlloc(accelBufferSizes.outputSizeInBytes, OPTIX_ACCEL_BUFFER_BYTE_ALIGNMENT, cuda::USAGE_TEMP); // This is a temporary buffer. The Compaction will be the static one! + + CUdeviceptr d_tmp = memAlloc(accelBufferSizes.tempSizeInBytes, OPTIX_ACCEL_BUFFER_BYTE_ALIGNMENT, cuda::USAGE_TEMP); + + OptixAccelEmitDesc accelEmit = {}; + + accelEmit.result = memAlloc(sizeof(size_t), sizeof(size_t), cuda::USAGE_TEMP); + accelEmit.type = OPTIX_PROPERTY_TYPE_COMPACTED_SIZE; + + OPTIX_CHECK( m_api.optixAccelBuild(m_optixContext, m_cudaStream, + &accelBuildOptions, &buildInput, 1, + d_tmp, accelBufferSizes.tempSizeInBytes, + data.d_gas, accelBufferSizes.outputSizeInBytes, + &data.traversable, &accelEmit, 1) ); + + size_t sizeCompact; + + CU_CHECK( cuMemcpyDtoHAsync(&sizeCompact, accelEmit.result, sizeof(size_t), m_cudaStream) ); + + synchronizeStream(); + + memFree(accelEmit.result); + memFree(d_tmp); + + // Compact the AS only when possible. This can save more than half the memory on RTX boards. + if (sizeCompact < accelBufferSizes.outputSizeInBytes) + { + CUdeviceptr d_gasCompact = memAlloc(sizeCompact, OPTIX_ACCEL_BUFFER_BYTE_ALIGNMENT); // This is the static GAS allocation! + + OPTIX_CHECK( m_api.optixAccelCompact(m_optixContext, m_cudaStream, data.traversable, d_gasCompact, sizeCompact, &data.traversable) ); + + synchronizeStream(); // Must finish accessing data.d_gas source before it can be freed and overridden. + + memFree(data.d_gas); + + data.d_gas = d_gasCompact; + + //std::cout << "Compaction saved " << accelBufferSizes.outputSizeInBytes - sizeCompact << '\n'; // DEBUG + accelBufferSizes.outputSizeInBytes = sizeCompact; // DEBUG for the std::cout below. + } + + // Return the relocation info for this GAS traversable handle from this device's OptiX context. + // It's used to assert that the GAS is compatible across devices which means NVLINK peer-to-peer sharing is allowed. + // (This is more meant as example code, because in NVLINK islands the GPU configuration must be homogeneous and addresses are unique with UVA.) + OPTIX_CHECK( m_api.optixAccelGetRelocationInfo(m_optixContext, data.traversable, &data.info) ); + + //std::cout << "createGeometry() device index = " << m_index << ": attributes = " << attributesSizeInBytes << ", indices = " << indicesSizeInBytes << ", GAS = " << accelBufferSizes.outputSizeInBytes << "\n"; // DEBUG + + return data; +} + + void Device::destroyGeometry(GeometryData& data) { memFree(data.d_gas); @@ -1422,10 +1563,19 @@ void Device::createInstance(const GeometryData& geometryData, const InstanceData instance.visibilityMask = 255; // PERF Determine which hit record to use. - // The matrix is Triangles with (no emission, emission) x (no cutout, cutout) shader. + // Triangles: Hit records 0 to 3. (no emission, emission) x (no cutout, cutout) + // Cubic B-splines: Hit record 4. + unsigned int hitRecord = 0; - hitRecord |= ((m_deviceShaderConfigurations[indexShader].flags & USE_EMISSION ) == 0) ? 0 : 1; // no emission, emission - hitRecord |= ((m_deviceShaderConfigurations[indexShader].flags & USE_CUTOUT_OPACITY) == 0) ? 0 : 2; // no cutout , cutout + if (geometryData.primitiveType == PT_CURVES) + { + hitRecord = 4; // Cubic B-spline curves. + } + else + { + hitRecord |= ((m_deviceShaderConfigurations[indexShader].flags & USE_EMISSION ) == 0) ? 0 : 1; // no emission, emission + hitRecord |= ((m_deviceShaderConfigurations[indexShader].flags & USE_CUTOUT_OPACITY) == 0) ? 0 : 2; // no cutout , cutout + } instance.sbtOffset = NUM_RAY_TYPES * hitRecord; instance.flags = OPTIX_INSTANCE_FLAG_NONE; @@ -1955,6 +2105,9 @@ void Device::initMaterialMDL(mi::neuraylib::ITransaction* transaction, dsc.idxCallGeometryCutoutOpacity = -1; + dsc.idxCallHairSample = -1; + dsc.idxCallHairEval = -1; + // Simplify the conditions by translating all constants unconditionally. if (config.thin_walled) { @@ -2104,7 +2257,13 @@ void Device::initMaterialMDL(mi::neuraylib::ITransaction* transaction, } } - // material.hair not implemented. + if (config.is_hair_bsdf_valid) + { + const std::string name = std::string("__direct_callable__hair") + suffix; + + dsc.idxCallHairSample = appendProgramGroupMDL(indexShader, name + std::string("_sample")); + dsc.idxCallHairEval = appendProgramGroupMDL(indexShader, name + std::string("_evaluate")); + } m_deviceShaderConfigurations.push_back(dsc); diff --git a/apps/MDL_renderer/src/Hair.cpp b/apps/MDL_renderer/src/Hair.cpp new file mode 100644 index 00000000..e23bbb80 --- /dev/null +++ b/apps/MDL_renderer/src/Hair.cpp @@ -0,0 +1,481 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +#include "inc/Hair.h" + +#include +#include +#include +#include + +#include "inc/MyAssert.h" + +#include "shaders/vector_math.h" + +// Hair file format and models courtesy of Cem Yuksel: http://www.cemyuksel.com/research/hairmodels/ + +Hair::Hair() +{ + memset(&m_header, 0, sizeof(m_header)); + + m_header.signature[0] = 'H'; + m_header.signature[1] = 'A'; + m_header.signature[2] = 'I'; + m_header.signature[3] = 'R'; +} + +//Hair::~Hair() +//{ +//} + +void Hair::setNumStrands(const unsigned int num) +{ + m_header.numStrands = num; +} + +unsigned int Hair::getNumStrands() const +{ + return m_header.numStrands; +} + +// Used when all strands have the same number of segments. +void Hair::setNumSegments(const unsigned int num) +{ + m_header.numSegments = num; +} + +// Used when each hair strand can have a different number of segments. +void Hair::setSegmentsArray(const std::vector& segments) +{ + // segments arrays can either be empty or must match the number of strands. + // That means numStrands needs to be set first! + MY_ASSERT(segments.empty() || m_header.numStrands == static_cast(segments.size())); + m_segmentsArray = segments; +} + +unsigned int Hair::getNumSegments(const unsigned int idxStrand) const +{ + MY_ASSERT(idxStrand < m_header.numStrands); + if (!m_segmentsArray.empty()) + { + return m_segmentsArray[idxStrand]; + } + return m_header.numSegments; +} + +void Hair::setPointsArray(const std::vector& points) +{ + MY_ASSERT(!points.empty()); // Points array cannot be empty. + m_pointsArray = points; + + // Changing numPoints potentially breaks the thickness, transparency, or color arrays temporarily. + // These should always be set after the points array when present. + m_header.numPoints = static_cast(m_pointsArray.size()); +} + +float3 Hair::getPoint(const unsigned int idx) const +{ + MY_ASSERT(idx < m_header.numPoints); + return m_pointsArray[idx]; +} + +void Hair::setThickness(const float thickness) +{ + m_header.thickness = thickness; +} + +void Hair::setThicknessArray(const std::vector& thickness) +{ + MY_ASSERT(thickness.empty() || m_header.numPoints == static_cast(thickness.size())); + m_thicknessArray = thickness; +} + +float Hair::getThickness(const unsigned int idx) const +{ + MY_ASSERT(idx < m_header.numPoints); + if (!m_thicknessArray.empty()) + { + return m_thicknessArray[idx]; + } + return m_header.thickness; +} + +void Hair::setTransparency(const float transparency) +{ + m_header.transparency = transparency; +} + +void Hair::setTransparencyArray(const std::vector& transparency) +{ + MY_ASSERT(transparency.empty() || m_header.numPoints == static_cast(transparency.size())); + m_transparencyArray = transparency; +} + +float Hair::getTransparency(const unsigned int idx) const +{ + MY_ASSERT(idx < m_header.numPoints); + if (!m_transparencyArray.empty()) + { + return m_transparencyArray[idx]; + } + return m_header.transparency; +} + +void Hair::setColor(const float3 color) +{ + m_header.color = color; +} + +void Hair::setColorArray(const std::vector& color) +{ + MY_ASSERT(color.empty() || m_header.numPoints == static_cast(color.size())); + m_colorArray = color; +} + +float3 Hair::getColor(const unsigned int idx) const +{ + MY_ASSERT(idx < m_header.numPoints); + if (!m_colorArray.empty()) + { + return m_colorArray[idx]; + } + return m_header.color; +} + +bool Hair::load(const std::string& filename) +{ + // This is a binary file format for 3D hair models. + FILE *file = fopen(filename.c_str(), "rb"); + if (!file) + { + std::cerr << "ERROR: Hair::load(): fopen(" << filename << ") failed.\n"; + return false; + } + + // A HAIR file begins with a 128-Byte long header. + size_t count = fread(&m_header, sizeof(Header), 1, file); + if (count < 1) + { + std::cerr << "ERROR: Hair::load(): Reading the header of " << filename << " failed.\n"; + fclose(file); + return false; + } + + if (strncmp(m_header.signature, "HAIR", 4) != 0) + { + std::cerr << "ERROR: Hair::load(): Header signature in " << filename << " not matching \"HAIR\"\n"; + fclose(file); + return false; + } + + // A HAIR file must have a points array, but all the other arrays are optional. + // When an array does not exist, corresponding default value from the file header is used instead of the missing array. + if ((m_header.bits & HAS_POINTS_ARRAY) == 0) + { + std::cerr << "ERROR: Hair::load(): Header bits in " << filename << " has no points array\n"; + fclose(file); + return false; + } + +#if 0 // DEBUG + std::cout << "Header of " << filename << " :\n"; + std::cout << "Number of hair strands = " << m_header.numStrands << "\n"; + std::cout << "Total number of points = " << m_header.numPoints; + if (m_header.bits & HAS_POINTS_ARRAY) // Put this first because this must always exist. + { + std::cout << "POINTS_ARRAY"; + } + if (m_header.bits & HAS_SEGMENTS_ARRAY) + { + std::cout << " | SEGMENTS_ARRAY"; + } + if (m_header.bits & HAS_THICKNESS_ARRAY) + { + std::cout << " | THICKNESS_ARRAY"; + } + if (m_header.bits & HAS_TRANSPARENCY_ARRAY) + { + std::cout << " | TRANSPARENCY_ARRAY"; + } + if (m_header.bits & HAS_COLOR_ARRAY) + { + std::cout << " | COLOR_ARRAY"; + } + std::cout << std::endl; + + if ((m_header.bits & HAS_SEGMENTS_ARRAY) == 0) + { + std::cout << "Default number of segments = " << m_header.numSegments << '\n'; + } + if ((m_header.bits & HAS_THICKNESS_ARRAY) == 0) + { + std::cout << "Default thickness = " << m_header.thickness << '\n'; + } + if ((m_header.bits & HAS_TRANSPARENCY_ARRAY) == 0) + { + std::cout << "Default transparency = " << m_header.transparency << '\n'; + } + if ((m_header.bits & HAS_COLOR_ARRAY) == 0) + { + std::cout << "Default color = " << m_header.color[0] << ", " << m_header.color[1] << ", " << m_header.color[2] << '\n'; + } +#endif + + // The 3D hair model consists of strands, each one of which is represented by a number of line segments. + + // Segments array (unsigned short) + // This array keeps the number of segments of each hair strand. + // Each entry is a 16-bit unsigned integer, therefore each hair strand can have up to 65536 segments. + if (m_header.bits & HAS_SEGMENTS_ARRAY) + { + m_segmentsArray.resize(m_header.numStrands); + + size_t count = fread(m_segmentsArray.data(), sizeof(unsigned short), m_header.numStrands, file); + if (count < m_header.numStrands) + { + std::cerr << "ERROR: Hair::load(): Failed to read segments array of " << filename << '\n'; + fclose(file); + return false; + } + } + + // Points array (float) + // This array keeps the 3D positions each of hair strand point. + // These points are not shared by different hair strands; each point belongs to a particular hair strand only. + // Line segments of a hair strand connects consecutive points. + // The points in this array are ordered by strand and from root to tip; + // such that it begins with the root point of the first hair strand, + // continues with the next point of the first hair strand until the tip of the first hair strand, + // and then comes the points of the next hair strands. + // Each entry is a 32-bit floating point number, and each point is defined by 3 consecutive numbers + // that correspond to x, y, and z coordinates. + if (m_header.bits & HAS_POINTS_ARRAY) + { + m_pointsArray.resize(m_header.numPoints); + + size_t count = fread(m_pointsArray.data(), sizeof(float3), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::load(): Failed to read points array of " << filename << '\n'; + fclose(file); + return false; + } + } + + // Thickness array (float) + // This array keeps the thickness of hair strands at point locations, + // therefore the size of this array is equal to the number of points. + // Each entry is a 32-bit floating point number. + if (m_header.bits & HAS_THICKNESS_ARRAY) + { + m_thicknessArray.resize(m_header.numPoints); + + size_t count = fread(m_thicknessArray.data(), sizeof(float), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::load(): Failed to read thickness array of " << filename << '\n'; + fclose(file); + return false; + } + } + + // Transparency array (float) + // This array keeps the transparency of hair strands at point locations, + // therefore the size of this array is equal to the number of points. + // Each entry is a 32-bit floating point number. + if (m_header.bits & HAS_TRANSPARENCY_ARRAY) + { + m_transparencyArray.resize(m_header.numPoints); + + size_t count = fread(m_transparencyArray.data(), sizeof(float), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::load(): Failed to read transparency array of " << filename << '\n'; + fclose(file); + return false; + } + } + + // Color array (float) + // This array keeps the color of hair strands at point locations, + // therefore the size of this array is three times the number of points. + // Each entry is a 32-bit floating point number, and each color is defined by 3 consecutive numbers + // that correspond to red, green, and blue components. + if (m_header.bits & HAS_COLOR_ARRAY) + { + m_colorArray.resize(m_header.numPoints); + + size_t count = fread(m_colorArray.data(), sizeof(float3), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::load(): Failed to read color array of " << filename << '\n'; + fclose(file); + return false; + } + } + + fclose(file); + + return true; +} + +bool Hair::save(const std::string& filename) +{ + // The points array must always be present and defines the sizes for thickness, transparency and color arrays when those are present. + if (m_pointsArray.empty()) + { + std::cerr << "ERROR: Hair::save(): Points array for " << filename << " is empty.\n"; + return false; + } + m_header.bits = HAS_POINTS_ARRAY; + + if (!m_segmentsArray.empty()) + { + m_header.bits |= HAS_SEGMENTS_ARRAY; + + if (m_header.numStrands != static_cast(m_segmentsArray.size())) + { + std::cerr << "ERROR: Hair::save(): Segments arrays for " << filename << " has incorrect size " << m_segmentsArray.size() << ", numStrands = "<< m_header.numStrands << '\n'; + return false; + } + else if (m_header.numPoints != m_header.numStrands * (m_header.numSegments + 1)) + { + std::cerr << "ERROR: Hair::save(): Number of points " << m_header.numPoints << " for " << filename << " not matching expected count " << m_header.numStrands * (m_header.numSegments + 1) << '\n'; + return false; + } + } + if (!m_thicknessArray.empty()) + { + m_header.bits |= HAS_THICKNESS_ARRAY; + + if (m_header.numPoints != static_cast(m_thicknessArray.size())) + { + std::cerr << "ERROR: Hair::save(): Thickness array for " << filename << " has incorrect size " << m_thicknessArray.size() << ", numPoints = "<< m_header.numPoints << '\n'; + return false; + } + } + if (!m_transparencyArray.empty()) + { + m_header.bits |= HAS_TRANSPARENCY_ARRAY; + + if (m_header.numPoints != static_cast(m_transparencyArray.size())) + { + std::cerr << "ERROR: Hair::save(): Transparency array for " << filename << " has incorrect size " << m_transparencyArray.size() << ", numPoints = "<< m_header.numPoints << '\n'; + return false; + } + } + if (!m_colorArray.empty()) + { + m_header.bits |= HAS_COLOR_ARRAY; + + if (m_header.numPoints != static_cast(m_colorArray.size())) + { + std::cerr << "ERROR: Hair::save(): Color array for " << filename << " has incorrect size " << m_colorArray.size() << ", numPoints = "<< m_header.numPoints << '\n'; + return false; + } + } + + // All other fields inside the Hair header are supposed to be filled by the caller. + + FILE *file = fopen(filename.c_str(), "wb"); + if (!file) + { + std::cerr << "ERROR: Hair::save(): fopen(" << filename << ") failed.\n"; + return false; + } + + // A HAIR file begins with a 128-Byte long header. + size_t count = fwrite(&m_header, sizeof(Header), 1, file); + if (count < 1) + { + std::cerr << "ERROR: Hair::save(): Writing the header of " << filename << " failed.\n"; + fclose(file); + return false; + } + + // The 3D hair model consists of strands, each one of which is represented by a number of line segments. + + if (m_header.bits & HAS_SEGMENTS_ARRAY) + { + size_t count = fwrite(m_segmentsArray.data(), sizeof(unsigned short), m_header.numStrands, file); + if (count < m_header.numStrands) + { + std::cerr << "ERROR: Hair::save(): Failed to write segments array of " << filename << '\n'; + fclose(file); + return false; + } + } + + if (m_header.bits & HAS_POINTS_ARRAY) + { + size_t count = fwrite(m_pointsArray.data(), sizeof(float3), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::save(): Failed to write points array of " << filename << '\n'; + fclose(file); + return false; + } + } + + if (m_header.bits & HAS_THICKNESS_ARRAY) + { + size_t count = fread(m_thicknessArray.data(), sizeof(float), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::save(): Failed to write thickness array of " << filename << '\n'; + fclose(file); + return false; + } + } + + if (m_header.bits & HAS_TRANSPARENCY_ARRAY) + { + size_t count = fwrite(m_transparencyArray.data(), sizeof(float), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::save(): Failed to write transparency array of " << filename << '\n'; + fclose(file); + return false; + } + } + + if (m_header.bits & HAS_COLOR_ARRAY) + { + size_t count = fwrite(m_colorArray.data(), sizeof(float3), m_header.numPoints, file); + if (count < m_header.numPoints) + { + std::cerr << "ERROR: Hair::save(): Failed to write color array of " << filename << '\n'; + fclose(file); + return false; + } + } + + fclose(file); + + return true; +} diff --git a/apps/MDL_renderer/src/Raytracer.cpp b/apps/MDL_renderer/src/Raytracer.cpp index a210aa89..1161a2fe 100644 --- a/apps/MDL_renderer/src/Raytracer.cpp +++ b/apps/MDL_renderer/src/Raytracer.cpp @@ -31,6 +31,8 @@ #include "inc/CheckMacros.h" +#include "shaders/config.h" + #include #include #include @@ -1037,6 +1039,70 @@ void Raytracer::traverseNode(std::shared_ptr node, InstanceData instan } } break; + + case sg::NodeType::NT_CURVES: + { + std::shared_ptr geometry = std::dynamic_pointer_cast(node); + + instanceData.idGeometry = geometry->getId(); + + const bool allowSharingGas = ((m_peerToPeer & P2P_GAS) != 0); + + if (allowSharingGas) + { + const unsigned int numIslands = static_cast(m_islands.size()); + + for (unsigned int indexIsland = 0; indexIsland < numIslands; ++indexIsland) + { + const auto& island = m_islands[indexIsland]; // Vector of device indices. + + const int deviceHome = getDeviceHome(island); + + // GeometryData is always shared and tracked per island. + GeometryData& geometryData = m_geometryData[instanceData.idGeometry * numIslands + indexIsland]; + + if (geometryData.traversable == 0) // If there is no traversable handle for this geometry in this island, try to create one on the home device. + { + geometryData = m_devicesActive[deviceHome]->createGeometry(geometry); + } + else + { + std::cout << "traverseNode() Geometry " << instanceData.idGeometry << " reused\n"; // DEBUG + } + + m_devicesActive[deviceHome]->createInstance(geometryData, instanceData, matrix); + + // Now share the GeometryData on the other devices in this island. + for (const auto device : island) + { + if (device != deviceHome) + { + // Create the instance referencing the shared GAS traversable on the peer device in this island. + // This is only host data. The IAS is created after gathering all flattened instances in the scene. + m_devicesActive[device]->createInstance(geometryData, instanceData, matrix); + } + } + } + } + else + { + const unsigned int numDevices = static_cast(m_devicesActive.size()); + + for (unsigned int device = 0; device < numDevices; ++device) + { + GeometryData& geometryData = m_geometryData[instanceData.idGeometry * numDevices + device]; + + if (geometryData.traversable == 0) // If there is no traversable handle for this geometry on this device, try to create one. + { + geometryData = m_devicesActive[device]->createGeometry(geometry); + } + + m_devicesActive[device]->createInstance(geometryData, instanceData, matrix); + } + } + } + break; + } } @@ -1214,12 +1280,13 @@ class Default_logger: public mi::base::Interface_implement severity = "WARN: "; break; case mi::base::MESSAGE_SEVERITY_INFO: + //return; // DEBUG No info messages. severity = "INFO: "; break; case mi::base::MESSAGE_SEVERITY_VERBOSE: - return; // DAR DEBUG No verbose messages. + return; // DEBUG No verbose messages. case mi::base::MESSAGE_SEVERITY_DEBUG: - return; // DAR DEBUG No debug messages. + return; // DEBUG No debug messages. case mi::base::MESSAGE_SEVERITY_FORCE_32_BIT: return; } @@ -1601,22 +1668,17 @@ bool Raytracer::initMDL(const std::vector& searchPaths) m_mdl_backend = mdl_backend_api->get_backend(mi::neuraylib::IMdl_backend_api::MB_CUDA_PTX); // Hardcoded values! - - // The hair BSDF would require two texture coordinates. Not supported in this example renderer version. - const unsigned int num_texture_spaces = 1; - - // FIXME PERF What is a good number here? - // If the number is smaller than there are results required by a shader, - // then the code calculating the additional results will happen inside the sample/evaluate/pdf functions! - // Meaning if this is too small, the code size and the runtime of a materal shader will increase. - const unsigned int num_texture_results = 16; + MY_STATIC_ASSERT(NUM_TEXTURE_SPACES == 1 || NUM_TEXTURE_SPACES == 2); + // The renderer only supports one or two texture spaces. + // The hair BSDF requires two texture coordinates! + // If you do not use the hair BSDF, NUM_TEXTURE_SPACES should be set to 1 for performance reasons. - if (m_mdl_backend->set_option("num_texture_spaces", std::to_string(num_texture_spaces).c_str()) != 0) + if (m_mdl_backend->set_option("num_texture_spaces", std::to_string(NUM_TEXTURE_SPACES).c_str()) != 0) { return false; } - if (m_mdl_backend->set_option("num_texture_results", std::to_string(num_texture_results).c_str()) != 0) + if (m_mdl_backend->set_option("num_texture_results", std::to_string(NUM_TEXTURE_RESULTS).c_str()) != 0) { return false; } @@ -1965,9 +2027,9 @@ void Raytracer::determineShaderConfiguration(const Compile_result& res, ShaderCo config.is_cutout_opacity_constant = res.compiled_material->get_cutout_opacity(&config.cutout_opacity); // This sets cutout opacity to -1.0 when it's not constant! config.use_cutout_opacity = !config.is_cutout_opacity_constant || config.cutout_opacity < 1.0f; - // material.hair is not supported by this renderer. - // That requires support for two texture coordinate and tangent space sets which would make everything else slower. - // to be done in a separate example. + mi::base::Handle hair_bsdf_expr(res.compiled_material->lookup_sub_expression("hair")); + + config.is_hair_bsdf_valid = isValidDistribution(hair_bsdf_expr.get()); // True if hair != hair_bsdf(). } @@ -2209,6 +2271,7 @@ void Raytracer::initMaterialMDL(MaterialMDL* materialMDL) std::string name_volume_scattering_coefficient = "__direct_callable__volume_scattering_coefficient" + suffix; std::string name_volume_directional_bias = "__direct_callable__volume_directional_bias" + suffix; std::string name_geometry_cutout_opacity = "__direct_callable__geometry_cutout_opacity" + suffix; + std::string name_hair_bsdf = "__direct_callable__hair" + suffix; // Centralize the init functions in a single material init(). // This will only save time when there would have been multiple init functions inside the shader. @@ -2291,6 +2354,10 @@ void Raytracer::initMaterialMDL(MaterialMDL* materialMDL) { descs.push_back(mi::neuraylib::Target_function_description("geometry.cutout_opacity", name_geometry_cutout_opacity.c_str())); } + if (config.is_hair_bsdf_valid) + { + descs.push_back(mi::neuraylib::Target_function_description("hair", name_hair_bsdf.c_str())); + } // Generate target code for the compiled material. mi::base::Handle link_unit(m_mdl_backend->create_link_unit(transaction, context.get())); diff --git a/apps/MDL_renderer/src/SceneGraph.cpp b/apps/MDL_renderer/src/SceneGraph.cpp index fa938fe8..58082227 100644 --- a/apps/MDL_renderer/src/SceneGraph.cpp +++ b/apps/MDL_renderer/src/SceneGraph.cpp @@ -236,5 +236,43 @@ namespace sg lightGUI.area = areaSurface; } + + // ========== Curves + Curves::Curves(const unsigned int id) + : Node(id) + { + } + + //Curves::~Curves() + //{ + //} + + sg::NodeType Curves::getType() const + { + return NT_CURVES; + } + + void Curves::setAttributes(std::vector const& attributes) + { + m_attributes.resize(attributes.size()); + memcpy(m_attributes.data(), attributes.data(), sizeof(CurveAttributes) * attributes.size()); + } + + std::vector const& Curves::getAttributes() const + { + return m_attributes; + } + + void Curves::setIndices(std::vector const& indices) + { + m_indices.resize(indices.size()); + memcpy(m_indices.data(), indices.data(), sizeof(unsigned int) * indices.size()); + } + + std::vector const& Curves::getIndices() const + { + return m_indices; + } + } // namespace sg diff --git a/data/hair/fur.hair b/data/hair/fur.hair new file mode 100644 index 00000000..32d499ae Binary files /dev/null and b/data/hair/fur.hair differ diff --git a/data/mdl/bsdf_diffuse_reflection_tex.mdl b/data/mdl/bsdf_diffuse_reflection_tex.mdl index 674416ee..d2f775c4 100644 --- a/data/mdl/bsdf_diffuse_reflection_tex.mdl +++ b/data/mdl/bsdf_diffuse_reflection_tex.mdl @@ -7,7 +7,7 @@ export material bsdf_diffuse_reflection_tex( uniform color parDiffuseTint = color(0.980392, 0.729412, 0.470588), uniform texture_2d parTexture = texture_2d("./textures/logo_512.png", tex::gamma_srgb), uniform float3 parRotation = float3(0.0), - uniform float3 parTranslation = float3(-0.25, 0.0, 0.0), + uniform float3 parTranslation = float3(0.0, 0.0, 0.0), uniform float3 parScaling = float3(1.0) ) = let diff --git a/data/mdl/bsdf_glossy_reflect_transmit.mdl b/data/mdl/bsdf_glossy_reflect_transmit.mdl index 6e6d7031..cbb553d7 100644 --- a/data/mdl/bsdf_glossy_reflect_transmit.mdl +++ b/data/mdl/bsdf_glossy_reflect_transmit.mdl @@ -1,8 +1,9 @@ mdl 1.7; import ::df::*; +import ::anno::*; export material bsdf_glossy_reflect_transmit( - uniform float parIor = 1.5, + uniform float parIor = 1.5 [[anno::hard_range(0.0, 5.0)]], uniform float parGlossyRoughnessU = 0.1, uniform float parGlossyRoughnessV = 0.1, uniform color parGlossyTint = color(0.980392, 0.729412, 0.470588) diff --git a/data/mdl/bsdf_glossy_transmit.mdl b/data/mdl/bsdf_glossy_transmit.mdl index 7b5e90b7..0493a444 100644 --- a/data/mdl/bsdf_glossy_transmit.mdl +++ b/data/mdl/bsdf_glossy_transmit.mdl @@ -1,8 +1,9 @@ mdl 1.7; import ::df::*; +import ::anno::*; export material bsdf_glossy_transmit( - uniform float parIor = 1.5, + uniform float parIor = 1.5 [[anno::hard_range(0.0, 5.0)]], uniform float parGlossyRoughnessU = 0.1, uniform float parGlossyRoughnessV = 0.1, uniform color parGlossyTint = color(0.980392, 0.729412, 0.470588) diff --git a/data/mdl/bsdf_hair.mdl b/data/mdl/bsdf_hair.mdl new file mode 100644 index 00000000..f22b8e54 --- /dev/null +++ b/data/mdl/bsdf_hair.mdl @@ -0,0 +1,31 @@ +mdl 1.7; +import ::df::*; +import ::math::*; +import ::state::*; +import ::tex::*; +import ::anno::*; + +export material bsdf_hair( + float parDiffuseReflectionWeight = 0.1 + , color parDiffuseReflectionTint = color(0.980392, 0.729412, 0.470588) + , float2 parRoughness_R = float2(0.1) + , float2 parRoughness_TT = float2(0.2) + , float2 parRoughness_TRT = float2(0.3) + , float parCuticleAngleDeg = 3.0 [[anno::hard_range(-90.0, 90.0)]] // Reasonable values would be [0.0, 6.0] + , color parAbsorptionColor = color(0.980392, 0.729412, 0.470588) // This gets inverted to produce the absorption coefficient. + , float parIor = 1.55 [[anno::hard_range(0.0, 5.0)]] +) += +material( + hair: df::chiang_hair_bsdf( + diffuse_reflection_weight: parDiffuseReflectionWeight, + diffuse_reflection_tint: parDiffuseReflectionTint, + roughness_R: parRoughness_R, + roughness_TT: parRoughness_TT, + roughness_TRT: parRoughness_TRT, + cuticle_angle: math::radians(parCuticleAngleDeg), + // Invert to make color selection easier. Also make sure this never drops to zero since there is a log() on that. + absorption_coefficient: color(1.001) - parAbsorptionColor, + ior: parIor + ) +); diff --git a/data/mdl/bsdf_hair_uv.mdl b/data/mdl/bsdf_hair_uv.mdl new file mode 100644 index 00000000..521f1513 --- /dev/null +++ b/data/mdl/bsdf_hair_uv.mdl @@ -0,0 +1,41 @@ +mdl 1.7; +import ::df::*; +import ::math::*; +import ::state::*; +import ::tex::*; +import ::anno::*; + +export material bsdf_hair_uv( + uniform texture_2d parTexture = texture_2d("./textures/arrows.png", tex::gamma_linear) + , float parDiffuseReflectionWeight = 0.1 + , color parDiffuseReflectionTint = color(0.980392, 0.729412, 0.470588) + , float2 parRoughness_R = float2(0.1) + , float2 parRoughness_TT = float2(0.2) + , float2 parRoughness_TRT = float2(0.3) + , float parCuticleAngleDeg = 3.0 [[anno::hard_range(-90.0, 90.0)]] // Reasonable values would be [0.0, 6.0] + , color parAbsorptionColor = color(0.980392, 0.729412, 0.470588) // This gets inverted to produce the absorption coefficient. + , float parIor = 1.55 [[anno::hard_range(0.0, 5.0)]] +) += let +{ + + float3 tc = state::texture_coordinate(0); + +} in material( + hair: df::chiang_hair_bsdf( + diffuse_reflection_weight: parDiffuseReflectionWeight, + diffuse_reflection_tint: parDiffuseReflectionTint, + //diffuse_reflection_tint: math::lerp(parDiffuseReflectionTint, color(1.0, 0.0, 0.0), state::texture_coordinate(0).x), // Debug uFiber + //diffuse_reflection_tint: math::lerp(parDiffuseReflectionTint, color(1.0, 0.0, 0.0), state::texture_coordinate(0).y), // Debug vFiber + roughness_R: parRoughness_R, + roughness_TT: parRoughness_TT, + roughness_TRT: parRoughness_TRT, + cuticle_angle: math::radians(parCuticleAngleDeg), + // Invert to make color selection easier. Also make sure this never drops to zero since there is a log() on that. + // The fiber is (u,v) parameterized, which means a 2D texture can be applied to it's surface. + // Add some tiny arrows which modulate the absoption. Zoom in on individual hairs to see them! + absorption_coefficient: color(1.001) - (parAbsorptionColor * tex::lookup_color(parTexture, float2(tc.x * 8.0, tc.y * 4.0f))), + ior: parIor + ) +); + diff --git a/data/mdl/bsdf_specular_reflect_transmit.mdl b/data/mdl/bsdf_specular_reflect_transmit.mdl index 84e969c8..44b6d487 100644 --- a/data/mdl/bsdf_specular_reflect_transmit.mdl +++ b/data/mdl/bsdf_specular_reflect_transmit.mdl @@ -1,8 +1,9 @@ mdl 1.7; import ::df::*; +import ::anno::*; export material bsdf_specular_reflect_transmit( - uniform float parIor = 1.5, + uniform float parIor = 1.5 [[anno::hard_range(0.0, 5.0)]], uniform color parSpecularTint = color(0.980392, 0.729412, 0.470588) ) = diff --git a/data/mdl/bsdf_specular_transmit.mdl b/data/mdl/bsdf_specular_transmit.mdl index 653bbd69..ba84bbce 100644 --- a/data/mdl/bsdf_specular_transmit.mdl +++ b/data/mdl/bsdf_specular_transmit.mdl @@ -3,7 +3,7 @@ import ::df::*; import ::anno::*; export material bsdf_specular_transmit( - uniform float parIor = 1.5, + uniform float parIor = 1.5 [[anno::hard_range(0.0, 5.0)]], uniform color parSpecularTint = color(0.980392, 0.729412, 0.470588) ) = diff --git a/data/mdl/diffuse_light.mdl b/data/mdl/diffuse_light.mdl new file mode 100644 index 00000000..f2858261 --- /dev/null +++ b/data/mdl/diffuse_light.mdl @@ -0,0 +1,18 @@ +mdl 1.7; +import ::df::*; + +// This is used inside scene_mdl_hair.txt. +export material diffuse_light( + uniform color parIntensityTint = color(1.0), + uniform float parIntensity = 31.41592 +) += +material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: parIntensityTint * parIntensity, + mode: intensity_radiant_exitance + ) + ) +); diff --git a/data/mdl/edf_spot.mdl b/data/mdl/edf_spot.mdl index 1e75df9c..12d09a74 100644 --- a/data/mdl/edf_spot.mdl +++ b/data/mdl/edf_spot.mdl @@ -5,7 +5,7 @@ export material edf_spot( uniform color parIntensityTint = color(1.0), uniform float parIntensity = 3.141592, uniform float parSpotExponent = 1.0, - uniform bool parGlobalDistribution = false + uniform bool parGlobalDistribution = false // The MDL SDK does not generate code for true! ) = material( diff --git a/data/mdl/edf_spot_power.mdl b/data/mdl/edf_spot_power.mdl index 11e48617..62a29322 100644 --- a/data/mdl/edf_spot_power.mdl +++ b/data/mdl/edf_spot_power.mdl @@ -5,7 +5,7 @@ export material edf_spot_power( uniform color parIntensityTint = color(1.0), uniform float parIntensity = 3.141592, uniform float parSpotExponent = 1.0, - uniform bool parGlobalDistribution = false + uniform bool parGlobalDistribution = false // The MDL SDK does not generate code for true! ) = material( diff --git a/data/mdl/layer_custom_curve.mdl b/data/mdl/layer_custom_curve.mdl index 0feff125..477ac2ed 100644 --- a/data/mdl/layer_custom_curve.mdl +++ b/data/mdl/layer_custom_curve.mdl @@ -1,11 +1,12 @@ mdl 1.7; import ::df::*; +import ::anno::*; export material layer_custom_curve( - uniform float parLayerNormalReflectivity = 0.0, - uniform float parLayerGrazingReflectivity = 1.0, + uniform float parLayerNormalReflectivity = 0.0 [[anno::hard_range(0.0, 1.0)]], + uniform float parLayerGrazingReflectivity = 1.0 [[anno::hard_range(0.0, 1.0)]], uniform float parLayerExponent = 5.0, - uniform float parLayerWeight = 1.0, + uniform float parLayerWeight = 1.0 [[anno::hard_range(0.0, 1.0)]], uniform color parSpecularTint = color(1.0, 1.0, 1.0), uniform color parDiffuseTint = color(0.980392, 0.729412, 0.470588) // , uniform float parDiffuseRoughness = 0.0 diff --git a/data/mdl/layer_custom_curve_glossy.mdl b/data/mdl/layer_custom_curve_glossy.mdl index 3dd736a0..dd8b03b4 100644 --- a/data/mdl/layer_custom_curve_glossy.mdl +++ b/data/mdl/layer_custom_curve_glossy.mdl @@ -1,11 +1,12 @@ mdl 1.7; import ::df::*; +import ::anno::*; export material layer_custom_curve_glossy( - uniform float parLayerNormalReflectivity = 0.0, - uniform float parLayerGrazingReflectivity = 1.0, + uniform float parLayerNormalReflectivity = 0.0 [[anno::hard_range(0.0, 1.0)]], + uniform float parLayerGrazingReflectivity = 1.0 [[anno::hard_range(0.0, 1.0)]], uniform float parLayerExponent = 5.0, - uniform float parLayerWeight = 1.0, + uniform float parLayerWeight = 1.0 [[anno::hard_range(0.0, 1.0)]], uniform float parGlossyRoughnessU = 0.1, uniform float parGlossyRoughnessV = 0.1, uniform color parGlossyTint = color(1.0), diff --git a/data/mdl/layer_fresnel.mdl b/data/mdl/layer_fresnel.mdl index a5df3bf0..dbd43f47 100644 --- a/data/mdl/layer_fresnel.mdl +++ b/data/mdl/layer_fresnel.mdl @@ -1,8 +1,9 @@ mdl 1.7; import ::df::*; +import ::anno::*; export material layer_fresnel( - uniform float parLayerIor = 1.5, + uniform float parLayerIor = 1.5 [[anno::hard_range(0.0, 5.0)]], uniform float parLayerWeight = 1.0, uniform color parSpecularTint = color(1.0, 1.0, 1.0), uniform color parDiffuseTint = color(0.980392, 0.729412, 0.470588) diff --git a/data/mdl/layer_fresnel_color.mdl b/data/mdl/layer_fresnel_color.mdl index a5b69ab8..f9f355cf 100644 --- a/data/mdl/layer_fresnel_color.mdl +++ b/data/mdl/layer_fresnel_color.mdl @@ -1,10 +1,11 @@ mdl 1.4; import ::df::*; +import ::anno::*; export material layer_fresnel_color( - uniform float parLayerIorR = 1.5, // The IOR cannot be a color on the GUI side or 1.0 would be the maximum! - uniform float parLayerIorG = 1.5, - uniform float parLayerIorB = 1.5, + uniform float parLayerIorR = 1.5 [[anno::hard_range(0.0, 5.0)]], + uniform float parLayerIorG = 1.5 [[anno::hard_range(0.0, 5.0)]], + uniform float parLayerIorB = 1.5 [[anno::hard_range(0.0, 5.0)]], uniform color parLayerWeight = color(1.0), uniform color parSpecularTint = color(1.0), uniform color parDiffuseTint = color(0.980392, 0.729412, 0.470588) diff --git a/data/mdl/modifier_directional_factor.mdl b/data/mdl/modifier_directional_factor.mdl index 137f8fea..dec82297 100644 --- a/data/mdl/modifier_directional_factor.mdl +++ b/data/mdl/modifier_directional_factor.mdl @@ -5,10 +5,9 @@ export material modifier_directional_factor( uniform color parNormalTint = color(1.0, 0.0, 0.0), uniform color parGrazingTint = color(0.0, 1.0, 0.0), uniform float parExponent = 1.0, - //uniform color parSpecularTint = color(0.980392, 0.729412, 0.470588) uniform float parGlossyRoughnessU = 0.1, uniform float parGlossyRoughnessV = 0.1, - uniform color parGlossyTint = color(0.980392, 0.729412, 0.470588) + uniform color parGlossyTint = color(1.0) ) = material( @@ -17,10 +16,6 @@ material( normal_tint: parNormalTint, grazing_tint: parGrazingTint, exponent: parExponent, - //base: df::specular_bsdf( - // tint: parSpecularTint, - // mode: df::scatter_reflect - //) base: df::simple_glossy_bsdf( roughness_u: parGlossyRoughnessU, roughness_v: parGlossyRoughnessV, diff --git a/data/mdl/modifier_fresnel_factor.mdl b/data/mdl/modifier_fresnel_factor.mdl index 0c86ce83..3a192dd5 100644 --- a/data/mdl/modifier_fresnel_factor.mdl +++ b/data/mdl/modifier_fresnel_factor.mdl @@ -1,14 +1,15 @@ mdl 1.4; import ::df::*; import ::math::*; +import ::anno::*; export material modifier_fresnel_factor( - uniform float parIorR = 0.2, - uniform float parIorG = 0.3, - uniform float parIorB = 0.4, - uniform float parExtinctionR = 0.4, - uniform float parExtinctionG = 0.5, - uniform float parExtinctionB = 0.6, + uniform float parIorR = 2.7 [[anno::hard_range(0.0, 10.0)]], + uniform float parIorG = 3.0 [[anno::hard_range(0.0, 10.0)]], + uniform float parIorB = 3.3 [[anno::hard_range(0.0, 10.0)]], + uniform float parExtinctionR = 0.7, + uniform float parExtinctionG = 0.8, + uniform float parExtinctionB = 0.9, uniform float parGlossyRoughnessU = 0.1, uniform float parGlossyRoughnessV = 0.1, uniform color parGlossyTint = color(1.0) diff --git a/data/mdl/modifier_thin_film.mdl b/data/mdl/modifier_thin_film.mdl index 3b671af1..61733bc7 100644 --- a/data/mdl/modifier_thin_film.mdl +++ b/data/mdl/modifier_thin_film.mdl @@ -1,21 +1,24 @@ mdl 1.7; import ::df::*; +import ::anno::*; + export material modifier_thin_film( - uniform float parThickness = 500.0, - uniform float parIor = 1.5, - uniform color parDiffuseTint = color(0.980392, 0.729412, 0.470588) - // , uniform float parDiffuseRoughness = 0.0 + uniform float parThickness = 500.0, + uniform float parIor = 1.5 [[anno::hard_range(0.0, 5.0)]], + uniform color parSpecularTint = color(1.0) ) = material( + thin_walled: true, surface: material_surface( scattering: df::thin_film( thickness: parThickness, ior: color(parIor), - base: df::diffuse_reflection_bsdf( - tint: parDiffuseTint - // , roughness: parDiffuseRoughness + // The thin-film modifier only affects BSDFs with a Fresnel term or scatter_reflect_transmit mode! + base: df::specular_bsdf( + tint: parSpecularTint, + mode: df::scatter_reflect_transmit ) ) ) diff --git a/data/mdl/textures/arrows.png b/data/mdl/textures/arrows.png new file mode 100644 index 00000000..4fcb98fa Binary files /dev/null and b/data/mdl/textures/arrows.png differ diff --git a/data/mdl/vdf_anisotropic.mdl b/data/mdl/vdf_anisotropic.mdl index adf3791f..a63ed30d 100644 --- a/data/mdl/vdf_anisotropic.mdl +++ b/data/mdl/vdf_anisotropic.mdl @@ -4,11 +4,11 @@ import ::math::*; import ::anno::*; export material vdf_anisotropic( - uniform float parIor = 1.0 [[ ::anno::hard_range(1.0f, 3.0f) ]], + uniform float parIor = 1.0 [[ ::anno::hard_range(1.0, 5.0) ]], uniform color parSpecularTint = color(1.0, 1.0, 1.0), uniform color parVolumeAbsorption = color(1.0), // color(0.980392, 0.729412, 0.470588), uniform color parVolumeScattering = color(0.9, 0.7, 0.5), - uniform float parDirectionalBias = 0.0 [[ ::anno::hard_range(-1.0f, 1.0f) ]], + uniform float parDirectionalBias = 0.0 [[ ::anno::hard_range(-1.0, 1.0) ]], // The thickness of the material at full absorption tint. uniform float parDistanceScale = 0.1 ) diff --git a/data/scene_mdl_demo.txt b/data/scene_mdl_demo.txt index 5991deef..7d98df68 100644 --- a/data/scene_mdl_demo.txt +++ b/data/scene_mdl_demo.txt @@ -10,7 +10,7 @@ lensShader 0 # Camera center of interest. # Absolute x, y, z coordinates in scene units (meters) -center 25 2.24 -1.25 +center 25 2.24 -4 # Camera orientation relative to center of interest and projection # phi [0.0f, 1.0f] # 0.75 is positive z-axis @@ -103,6 +103,7 @@ mdl modifier_tint modifier_tint "mdl/modifier_tint.mdl" mdl modifier_directional_factor modifier_directional_factor "mdl/modifier_directional_factor.mdl" mdl modifier_fresnel_factor modifier_fresnel_factor "mdl/modifier_fresnel_factor.mdl" mdl modifier_measured_curve_factor modifier_measured_curve_factor "mdl/modifier_measured_curve_factor.mdl" +mdl modifier_thin_film modifier_thin_film "mdl/modifier_thin_film.mdl" mdl bsdf_sheen bsdf_sheen "mdl/bsdf_sheen.mdl" @@ -406,6 +407,11 @@ translate 35 1 -7.5 model sphere 180 90 1 modifier_measured_curve_factor pop +push +translate 35 1 -10 +model sphere 180 90 1 modifier_thin_film +pop + # Sheen BSDF push translate 37.5 1 0 diff --git a/data/scene_mdl_hair.txt b/data/scene_mdl_hair.txt new file mode 100644 index 00000000..8dd6264f --- /dev/null +++ b/data/scene_mdl_hair.txt @@ -0,0 +1,116 @@ +# ========== CAMERA + +# Lens shader callable program. +# 0 = pinhole +# 1 = full format fisheye +# 2 = spherical projection + +lensShader 0 + +# Camera center of interest. +# Absolute x, y, z coordinates in scene units (meters) + +center 0 0 0 + +# Camera orientation relative to center of interest and projection +# phi [0.0f, 1.0f] # 0.75 is positive z-axis +# theta [0.0f, 1.0f] # 0 is south pole, 0.5 is equator, 1.0 is north pole +# yfov in degrees [1, 179] +# distance from center of interest [0.0f, inf] in meters + +camera 0.75 0.5 45 4 + + +# ========== TONEMAPPER +# Neutral tonemapper settings. +#gamma 1 +#colorBalance 1 1 1 +#whitePoint 1 +#burnHighlights 1 +#crushBlacks 0 +#saturation 1 +#brightness 1 + +# Usual tonemapper settings: +gamma 2.2 +colorBalance 1 1 1 +whitePoint 1 +burnHighlights 0.8 +crushBlacks 0.2 +saturation 1.2 +brightness 1 + +# ========== MATERIALS + +# Default material. The special material name "default" is used if the referenced material name is not found. + +mdl default bsdf_diffuse_reflection "mdl/bsdf_diffuse_reflection.mdl" + +mdl diffuse_light diffuse_light "mdl/diffuse_light.mdl" + +mdl bsdf_hair bsdf_hair "mdl/bsdf_hair.mdl" + +# A hair material with (u, v) texture placement of tiny arrows from root to tip on each fiber. Zoom in closely to see them. +mdl bsdf_hair_uv bsdf_hair_uv "mdl/bsdf_hair_uv.mdl" + +# ========== LIGHTS + +# IMPORTANT: The one environment light inside the scene must always be defined as first light! +# Environment lights are hardcoded to light definition index 0 inside the renderer. +# Not defining an environment light will result in a black environment. + +push +# DAR FIXME Lights are currently defined outside the MDL shader handling. +# For arbitrary mesh lights this can be changed in the future. +# The following parameters only affect lights. +emission 1 1 1 +emissionMultiplier 1 +emissionTexture "NV_Default_HDR_3000x1500.hdr" +# The transformations are taken from the current state, means the spherical environment texture can be rotated arbitrarily here. +# rotate 0 1 0 90 +#light env +pop + +# ========== GEOMETRY + +# Place an area light above the hair model. +push +scale 2 2 2 +rotate 1 0 0 90 +translate 0 4 2 +model plane 1 1 2 diffuse_light +pop + + +# The hair below is generated around a unit sphere. +# Try adjusting the default parDiffuseColor parameter. +push +model sphere 180 90 1 default +pop + + +push +# Syntax for loading *.hair models: +# model hair "" + +# The fur.hair model had been generated with a seprate program using the Hair class inside the renderer. It's very simple. +model hair 1.0 bsdf_hair "./hair/fur.hair" + +# Same mode with a material with (u, v) texture placement of tiny arrows from root to tip on each fiber. Zoom in closely to see them. +#model hair 1.0 bsdf_hair_uv "./hair/fur.hair" +pop + + + +## The hair models on http://www.cemyuksel.com/research/hairmodels/ +## are constructed with positive z-axis up and positive x-axis front. +## Use the following transforms to make them y-up and face the positive z-axis. +## They are also modeled in centimeters. Change the light setup when using this. +#push +## Convert right-handed z-up to y-up. +#rotate 1 0 0 -90 +## Rotate face to point to to positive z-axis. +#rotate 0 1 0 -90 +## Assumes the *.hair files are downloaded into the "hair" folder. +#model hair 1.0 bsdf_hair "./hair/wStraight.hair" +#pop diff --git a/data/system_mdl_hair.txt b/data/system_mdl_hair.txt new file mode 100644 index 00000000..d1440147 --- /dev/null +++ b/data/system_mdl_hair.txt @@ -0,0 +1,170 @@ +# This system options file handles multiple settings of the same option, the last one wins! +# (Current system settings can be written to a file by hitting Key S.) + +# Define the raytracer's rendering strategy +# This setting is only used in rtigo3. nvlink_shared, rtigo9 and rtigo10 use the strategy 3 when there are more then one devices active. +# 0 = Interactive Single-GPU, with or without OpenGL interop. +# Full frame accumulation in local memory, read to host buffer when needed. +# 1 = Interactive Multi-GPU Zero Copy, no OpenGL interop. +# Tiled rendering with tileSize blocks in a checkered pattern distributed to all enabled GPUs directly to pinned memory on the host. +# Works with any number of enabled devices. +# 2 = Interactive Multi-GPU Peer Access +# Tiled rendering with tileSize blocks in a checkered pattern evenly distributed to all enabled GPUs. +# The full image is allocated only on the first device, the peer devices directly render into the shared buffer. +# This is not going to work with more than one island in the active devices. +# 3 = Interactive Multi-GPU rendering into local GPU buffers of roughly 1/activeDevices size. +# Tiled rendering with tileSize blocks in a checkered pattern evenly distributed to all enabled GPUs. +# The full image is composited on the first device resp. the OpenGL interop device. +# The local data from other devices (not full resolution) is copied to that main device and composited by a native CUDA kernel. + +# strategy 3 + +# The devicesMask indicates which devices should be used in a 32-bit bitfield. +# The default is 255 which means 8 bits set so all boards in an RTX server. +# The application will only use the boards actually visible. +# This uses the first visble GPU device only: + +devicesMask 1 + +# The default size for Arena allocations in mebi-bytes. +# Default is 64 MiB when not set. Minimum is 1 MiB. + +arenaSize 128 + +# Use different strategies to update the OpenGL display texture. +# If the system configuration is not running an NVIDIA OpenGL implementation, set this to 0. +# The performance effect of interop 2 is only really visible interactive rendering (-m 0) and present 1. +# 0 = Use host buffers to transfer the result into the OpenGL display texture (slowest). +# 1 = Register the texture image with CUDA and copy into the array directly (fewest copies). +# 2 = Register the pixel buffer for direct rendering in single GPU or as staging buffer in multi-GPU (needs more memory than interop 1). +# Not available with multi-GPU zero copy strategy because the buffer resides in host memory then. +# For multi-GPU peer access the renderer cannot directly render with peer-to-peer into the OpenGL PBO and needs a separate shared buffer for rendering. +# Benchmarking is faster with interop off! + +interop 1 + +# Controls if every rendered image or final tile should be displayed (1) or only once per second (0) to save PCI-E bandwidth. +# Can be toggled inside the GUI. +# 0 = present only once per second (except for the first half second which accumulates) +# 1 = present every rendered image + +present 0 + +# Controls resource sharing across devices via CUDA peer-to-peer access when possible. +# Bit 0 = Allow sharing via PCI-E bus. Only share across NVLINK bridges when off (default off) +# Bit 1 = Allow texture CUarray or CUmipmappedArray data sharing via P2P (fast) (default on) +# Bit 2 = Allow geometry acceleration and vertex attribute sharing via P2P (slower) (default on) +# Bit 3 = Allow sharing of spherical environment light texture and CDFs (slow) (default off) + +peerToPeer 0 + +# Rendering resolution is independent of the window client size. +# The display of the texture is centered in the client window. +# If the image fits, the surrounding is black. +# If it's shrunk to fit, the surrounding pixels are dark red. + +resolution 512 512 + +# When usiong multi-GPU the workload per sub-frame is distributed in small tiless of this defined tile sizer. +# Default is tileSize 8 8 +# Values must be power-of-two and shouldn't be narrower than 8 or smaller than 32 pixels due to the warp size! + +tileSize 16 16 + +# The integer samplesSqrt is the sqrt(samples per pixel). Default is 1. +# The camera samples are distributed with a fixed rotated grid. +# Final frame rendering algorithms need the samples per pixels anyway. + +samplesSqrt 16 + +# Path lengths minimum and maximum. +# Minimum path length before Russian Roulette kicks in. +# Maximum path length before termination. +# Set min >= max to disable Russian Rouelette. +# pathLengths in range [0, 100] + +pathLengths 2 6 + +# Volume scattering random walk length after which the maximum distance is selected to potentially exist the volume. +# Minimum and default 1 for single scattering +# The maximum path length includes random walk steps, so the maximum path length should always be bigger than the walkLength. + +walkLength 2 + +# Scene dependent epsilon factor scaled by 1.0e-7. +# The renderer works in meters for the absorption, that means epsilonFactor 1000 is a scene epsilon of 1e-4 which is a tenth of a millimeter. +# Used for cheap self intersection avoidance by changing ray t_min (and t_max for visibility checks) +# epsilonFactor in range [0.0f, 10000.0f] (because of the GUI). + +epsilonFactor 1000 + +# Time vizualization clock factor scaled by 1.0e-9. +# Time view is a compile time options and this has not effect when that is not enabled. +# Means with 1000 all values >1.0 in the time view output (alpha channel) have taken a million clocks or more. + +clockFactor 1000 + +# Lens shader callable program. +# Can be overriden inside the scene description! +# 0 = pinhole +# 1 = full format fisheye +# 2 = spherical projection + +lensShader 0 + +# Camera center of interest. +# Can be overriden inside the scene description! +# Absolute x, y, z coordinates in scene units (meters) + +center 0 0 0 + +# Camera orientation relative to center of interest and projection, looking at the center of interest. +# Can be overriden inside the scene description! +# phi [0.0f, 1.0f] # 0.75 is positive z-axis +# theta [0.0f, 1.0f] # 0.0 is south pole, 0.5 is equator, 1.0 is north pole. +# yfov in degrees [1, 179] +# distance from center of interest [0.0f, inf] in meters + +camera 0.75 0.5 45 4 + +# Path with a folder and optional partial filename prefix which should receive the screenshots. +# (Key P for tonemapped *.png, Key H for linear *.hdr image.) +# If this is just a folder, end it with '/' +# The folder must already exist or saving screenshot fails! + +prefixScreenshot "./MDL_renderer" + +# The search paths used to look for MDL files and resources referenced inside the scene description materials. +# Multiple searchPath can be given here, which get strored into a vector in the given order. + +# Note that the two default MDL vMaterials search paths from the +# environment variables MDL_SYSTEM_PATH and MDL_USER_PATH are automatically added by the application in Raytracer::initMDL(). +# Their defaults are "C:\ProgramData\NVIDIA Corporation\mdl\" and "C:\Users\\Documents\mdl\" under Windows. +# Any additional search paths to other *.mdl files and their resources can be added with multiple searchPath options here. +# If there are any collisions of vMaterials files with your user defined files, comment out code linse +# m_mdl_config->add_mdl_system_paths(); +# m_mdl_config->add_mdl_user_paths(); +# and set the searchPaths in the desired order here. + +searchPath "./" + +# Tonemapper settings. +# Can be adjusted in the GUI. +# Can be overriden inside the scene description! + +# Neutral tonemapper GUI settings showing the linear image: +# gamma 1 +# whitePoint 1 +# burnHighlights 1 +# crushBlacks 0 +# saturation 1 +# brightness 1 + +# Standard tonemapper settings: +gamma 2.2 +colorBalance 1 1 1 +whitePoint 1 +burnHighlights 0.8 +crushBlacks 0.2 +saturation 1.2 +brightness 0.8