Implement dynamic material example and dynamic geometry example

.gitignore

@@ -1,10 +1,12 @@
 .idea/
 __pycache__/
 build/
+dist/
+imgui.ini
 *.egg-info/
 .*/
 *.so
 *.html
 *.cpp
 *.c
-.*
+.*

examples/cuda/dynamic_geometry.cu

@@ -0,0 +1,146 @@
+//
+// Copyright (c) 2021, 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 <optix.h>
+
+#include "dynamic_geometry.h"
+#include "vec_math.h"
+#include "helpers.h"
+
+extern "C" {
+    __constant__ Params params;
+}
+
+
+static __forceinline__ __device__ void trace(
+    OptixTraversableHandle handle,
+    float3                 ray_origin,
+    float3                 ray_direction,
+    float                  tmin,
+    float                  tmax,
+    float3*                prd
+)
+{
+    unsigned int p0, p1, p2;
+    p0 = float_as_int( prd->x );
+    p1 = float_as_int( prd->y );
+    p2 = float_as_int( prd->z );
+    optixTrace(
+        handle,
+        ray_origin,
+        ray_direction,
+        tmin,
+        tmax,
+        0.0f,                // rayTime
+        OptixVisibilityMask( 1 ),
+        OPTIX_RAY_FLAG_NONE,
+        0,                   // SBT offset
+        0,                   // SBT stride
+        0,                   // missSBTIndex
+        p0, p1, p2 );
+    prd->x = int_as_float( p0 );
+    prd->y = int_as_float( p1 );
+    prd->z = int_as_float( p2 );
+}
+
+
+static __forceinline__ __device__ void setPayload( float3 p )
+{
+    optixSetPayload_0( float_as_int( p.x ) );
+    optixSetPayload_1( float_as_int( p.y ) );
+    optixSetPayload_2( float_as_int( p.z ) );
+}
+
+
+static __forceinline__ __device__ float3 getPayload()
+{
+    return make_float3(
+        int_as_float( optixGetPayload_0() ),
+        int_as_float( optixGetPayload_1() ),
+        int_as_float( optixGetPayload_2() )
+    );
+}
+
+
+extern "C" __global__ void __raygen__rg()
+{
+    const uint3 idx = optixGetLaunchIndex();
+    const uint3 dim = optixGetLaunchDimensions();
+
+    const float3 eye = params.eye;
+    const float3 U = params.U;
+    const float3 V = params.V;
+    const float3 W = params.W;
+    const float2      d = 2.0f * make_float2(
+        static_cast< float >( idx.x ) / static_cast< float >( dim.x ),
+        static_cast< float >( idx.y ) / static_cast< float >( dim.y )
+    ) - 1.0f;
+
+    const float3 direction = normalize( d.x * U + d.y * V + W );
+    float3       payload_rgb = make_float3( 0.5f, 0.5f, 0.5f );
+
+    trace( params.trav_handle,
+        eye,
+        direction,
+        0.00f,  // tmin
+        1e16f,  // tmax
+        &payload_rgb );
+
+    params.frame_buffer[idx.y * params.width + idx.x] = make_color( payload_rgb );
+}
+
+
+extern "C" __global__ void __miss__ms()
+{
+    MissData* rt_data = reinterpret_cast< MissData* >( optixGetSbtDataPointer() );
+    float3    payload = getPayload();
+    setPayload( make_float3( rt_data->bg_color.x, rt_data->bg_color.y, rt_data->bg_color.z ) );
+}
+
+
+extern "C" __global__ void __closesthit__ch()
+{
+    HitGroupData* rt_data = reinterpret_cast< HitGroupData* >( optixGetSbtDataPointer() );
+
+    // fetch current triangle vertices
+    float3 data[3];
+    optixGetTriangleVertexData( optixGetGASTraversableHandle(), optixGetPrimitiveIndex(), optixGetSbtGASIndex(),
+        optixGetRayTime(), data );
+
+    // compute triangle normal
+    data[1] -= data[0];
+    data[2] -= data[0];
+    float3 normal = make_float3(
+        data[1].y*data[2].z - data[1].z*data[2].y,
+        data[1].z*data[2].x - data[1].x*data[2].z,
+        data[1].x*data[2].y - data[1].y*data[2].x );
+    const float s = 0.5f / sqrtf( normal.x*normal.x + normal.y*normal.y + normal.z*normal.z );
+
+    // convert normal to color and store in payload
+    setPayload( (normal*s + make_float3( 0.5 )) * rt_data->color );
+}

examples/cuda/dynamic_geometry.h

@@ -0,0 +1,56 @@
+//
+// Copyright (c) 2021, 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.
+//
+//
+
+struct Params
+{
+    uchar4*                frame_buffer;
+    unsigned int           width;
+    unsigned int           height;
+    float3                 eye, U, V, W;
+    OptixTraversableHandle trav_handle;
+    int                    subframe_index;
+};
+
+struct RayGenData
+{
+    float3 cam_eye;
+    float3 camera_u, camera_v, camera_w;
+};
+
+
+struct MissData
+{
+    float4 bg_color;
+};
+
+
+struct HitGroupData
+{
+    float3 color;
+};

examples/cuda/dynamic_geometry_vertex_generation.cu

@@ -0,0 +1,108 @@
+//
+// Copyright (c) 2021, 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 "vec_math.h"
+
+
+enum struct AnimationMode: int
+{
+    NONE = 0,
+    DEFORM = 1,
+    EXPLODE = 2,
+};
+
+
+__forceinline__ __device__ float triangle_wave( float x, float shift = 0.f, float period = 2.f * M_PIf, float amplitude = 1.f )
+{
+    return fabsf( fmodf( ( 4.f / period ) * ( x - shift ), 4.f * amplitude ) - 2.f * amplitude ) - amplitude;
+}
+
+__forceinline__ __device__ void write_animated_triangle( float3* out_vertices, int tidx, float3 v0, float3 v1, float3 v2, AnimationMode mode, float time )
+{
+    float3 v = make_float3( 0 );
+
+    if( mode == AnimationMode::EXPLODE )
+    {
+        // Generate displacement vector from triangle index
+        const float theta = ( (float)M_PIf * ( ( tidx + 1 ) * ( 13 / M_PIf ) ) );
+        const float phi   = ( (float)( 2.0 * M_PIf ) * ( ( tidx + 1 ) * ( 97 / M_PIf ) ) );
+
+        // Apply displacement to the sphere triangles
+        v = make_float3( triangle_wave( phi ) * triangle_wave( theta, M_PIf / 2.f ),
+            triangle_wave( phi, M_PIf / 2.f ) * triangle_wave( theta, M_PIf / 2.f ), triangle_wave( theta ) )
+            * triangle_wave( time, M_PIf / 2.f ) * 2.f;
+    }
+
+    out_vertices[tidx * 3 + 0] = v0 + v;
+    out_vertices[tidx * 3 + 1] = v1 + v;
+    out_vertices[tidx * 3 + 2] = v2 + v;
+}
+
+__forceinline__ __device__ float3 deform_vertex( const float3& c, AnimationMode mode, float time )
+{
+    // Apply sine wave to the y coordinate of the sphere vertices
+    if( mode == AnimationMode::DEFORM )
+        return make_float3( c.x, c.y * ( 0.5f + 0.4f * cosf( 4 * ( c.x + time ) ) ), c.z );
+    return c;
+}
+
+extern "C" __global__ void generate_vertices(float3* out_vertices, AnimationMode mode, float time, int width, int height)
+{
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    if( idx < width * height )
+    {
+        // generate a single patch (two unindexed triangles) of a tessellated sphere
+
+        int x = idx % width;
+        int y = idx / width;
+
+        const float theta0 = ( ( float )M_PIf * ( y + 0 ) ) / height;
+        const float theta1 = ( ( float )M_PIf * ( y + 1 ) ) / height;
+        const float phi0 = ( ( float )( 2.0 * M_PIf ) * ( x + 0 ) ) / width;
+        const float phi1 = ( ( float )( 2.0 * M_PIf ) * ( x + 1 ) ) / width;
+
+        const float ct0 = cosf( theta0 );
+        const float st0 = sinf( theta0 );
+        const float ct1 = cosf( theta1 );
+        const float st1 = sinf( theta1 );
+
+        const float cp0 = cosf( phi0 );
+        const float sp0 = sinf( phi0 );
+        const float cp1 = cosf( phi1 );
+        const float sp1 = sinf( phi1 );
+
+        const float3 v00 = deform_vertex( make_float3( cp0 * st0, sp0 * st0, ct0 ), mode, time );
+        const float3 v10 = deform_vertex( make_float3( cp0 * st1, sp0 * st1, ct1 ), mode, time );
+        const float3 v01 = deform_vertex( make_float3( cp1 * st0, sp1 * st0, ct0 ), mode, time );
+        const float3 v11 = deform_vertex( make_float3( cp1 * st1, sp1 * st1, ct1 ), mode, time );
+
+        write_animated_triangle( out_vertices, idx * 2 + 0, v00, v10, v11, mode, time );
+        write_animated_triangle( out_vertices, idx * 2 + 1, v00, v11, v01, mode, time );
+    }
+}