Skip to content

Commit d55fc17

Browse files
mortaciousmortacious
authored
Merge pull request #2 from keckj/master
Implement dynamic material example and dynamic geometry example
2 parents f021fca + 61c93d8 commit d55fc17

29 files changed

Lines changed: 2332 additions & 39 deletions

‎.gitignore‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,12 @@
11
.idea/
22
__pycache__/
33
build/
4+
dist/
5+
imgui.ini
46
*.egg-info/
57
.*/
68
*.so
79
*.html
810
*.cpp
911
*.c
10-
.*
12+
.*

‎examples/cuda/dynamic_geometry.cu‎

Lines changed: 146 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
1+
//
2+
// Copyright (c) 2021, NVIDIA CORPORATION. All rights reserved.
3+
//
4+
// Redistribution and use in source and binary forms, with or without
5+
// modification, are permitted provided that the following conditions
6+
// are met:
7+
// * Redistributions of source code must retain the above copyright
8+
// notice, this list of conditions and the following disclaimer.
9+
// * Redistributions in binary form must reproduce the above copyright
10+
// notice, this list of conditions and the following disclaimer in the
11+
// documentation and/or other materials provided with the distribution.
12+
// * Neither the name of NVIDIA CORPORATION nor the names of its
13+
// contributors may be used to endorse or promote products derived
14+
// from this software without specific prior written permission.
15+
//
16+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
17+
// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
19+
// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
20+
// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
21+
// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
22+
// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
23+
// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
24+
// OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25+
// (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26+
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27+
//
28+
29+
#include <optix.h>
30+
31+
#include "dynamic_geometry.h"
32+
#include "vec_math.h"
33+
#include "helpers.h"
34+
35+
extern "C" {
36+
__constant__ Params params;
37+
}
38+
39+
40+
static __forceinline__ __device__ void trace(
41+
OptixTraversableHandle handle,
42+
float3 ray_origin,
43+
float3 ray_direction,
44+
float tmin,
45+
float tmax,
46+
float3* prd
47+
)
48+
{
49+
unsigned int p0, p1, p2;
50+
p0 = float_as_int( prd->x );
51+
p1 = float_as_int( prd->y );
52+
p2 = float_as_int( prd->z );
53+
optixTrace(
54+
handle,
55+
ray_origin,
56+
ray_direction,
57+
tmin,
58+
tmax,
59+
0.0f, // rayTime
60+
OptixVisibilityMask( 1 ),
61+
OPTIX_RAY_FLAG_NONE,
62+
0, // SBT offset
63+
0, // SBT stride
64+
0, // missSBTIndex
65+
p0, p1, p2 );
66+
prd->x = int_as_float( p0 );
67+
prd->y = int_as_float( p1 );
68+
prd->z = int_as_float( p2 );
69+
}
70+
71+
72+
static __forceinline__ __device__ void setPayload( float3 p )
73+
{
74+
optixSetPayload_0( float_as_int( p.x ) );
75+
optixSetPayload_1( float_as_int( p.y ) );
76+
optixSetPayload_2( float_as_int( p.z ) );
77+
}
78+
79+
80+
static __forceinline__ __device__ float3 getPayload()
81+
{
82+
return make_float3(
83+
int_as_float( optixGetPayload_0() ),
84+
int_as_float( optixGetPayload_1() ),
85+
int_as_float( optixGetPayload_2() )
86+
);
87+
}
88+
89+
90+
extern "C" __global__ void __raygen__rg()
91+
{
92+
const uint3 idx = optixGetLaunchIndex();
93+
const uint3 dim = optixGetLaunchDimensions();
94+
95+
const float3 eye = params.eye;
96+
const float3 U = params.U;
97+
const float3 V = params.V;
98+
const float3 W = params.W;
99+
const float2 d = 2.0f * make_float2(
100+
static_cast< float >( idx.x ) / static_cast< float >( dim.x ),
101+
static_cast< float >( idx.y ) / static_cast< float >( dim.y )
102+
) - 1.0f;
103+
104+
const float3 direction = normalize( d.x * U + d.y * V + W );
105+
float3 payload_rgb = make_float3( 0.5f, 0.5f, 0.5f );
106+
107+
trace( params.trav_handle,
108+
eye,
109+
direction,
110+
0.00f, // tmin
111+
1e16f, // tmax
112+
&payload_rgb );
113+
114+
params.frame_buffer[idx.y * params.width + idx.x] = make_color( payload_rgb );
115+
}
116+
117+
118+
extern "C" __global__ void __miss__ms()
119+
{
120+
MissData* rt_data = reinterpret_cast< MissData* >( optixGetSbtDataPointer() );
121+
float3 payload = getPayload();
122+
setPayload( make_float3( rt_data->bg_color.x, rt_data->bg_color.y, rt_data->bg_color.z ) );
123+
}
124+
125+
126+
extern "C" __global__ void __closesthit__ch()
127+
{
128+
HitGroupData* rt_data = reinterpret_cast< HitGroupData* >( optixGetSbtDataPointer() );
129+
130+
// fetch current triangle vertices
131+
float3 data[3];
132+
optixGetTriangleVertexData( optixGetGASTraversableHandle(), optixGetPrimitiveIndex(), optixGetSbtGASIndex(),
133+
optixGetRayTime(), data );
134+
135+
// compute triangle normal
136+
data[1] -= data[0];
137+
data[2] -= data[0];
138+
float3 normal = make_float3(
139+
data[1].y*data[2].z - data[1].z*data[2].y,
140+
data[1].z*data[2].x - data[1].x*data[2].z,
141+
data[1].x*data[2].y - data[1].y*data[2].x );
142+
const float s = 0.5f / sqrtf( normal.x*normal.x + normal.y*normal.y + normal.z*normal.z );
143+
144+
// convert normal to color and store in payload
145+
setPayload( (normal*s + make_float3( 0.5 )) * rt_data->color );
146+
}

‎examples/cuda/dynamic_geometry.h‎

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
//
2+
// Copyright (c) 2021, NVIDIA CORPORATION. All rights reserved.
3+
//
4+
// Redistribution and use in source and binary forms, with or without
5+
// modification, are permitted provided that the following conditions
6+
// are met:
7+
// * Redistributions of source code must retain the above copyright
8+
// notice, this list of conditions and the following disclaimer.
9+
// * Redistributions in binary form must reproduce the above copyright
10+
// notice, this list of conditions and the following disclaimer in the
11+
// documentation and/or other materials provided with the distribution.
12+
// * Neither the name of NVIDIA CORPORATION nor the names of its
13+
// contributors may be used to endorse or promote products derived
14+
// from this software without specific prior written permission.
15+
//
16+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
17+
// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
19+
// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
20+
// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
21+
// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
22+
// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
23+
// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
24+
// OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25+
// (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26+
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27+
//
28+
//
29+
30+
struct Params
31+
{
32+
uchar4* frame_buffer;
33+
unsigned int width;
34+
unsigned int height;
35+
float3 eye, U, V, W;
36+
OptixTraversableHandle trav_handle;
37+
int subframe_index;
38+
};
39+
40+
struct RayGenData
41+
{
42+
float3 cam_eye;
43+
float3 camera_u, camera_v, camera_w;
44+
};
45+
46+
47+
struct MissData
48+
{
49+
float4 bg_color;
50+
};
51+
52+
53+
struct HitGroupData
54+
{
55+
float3 color;
56+
};

‎examples/cuda/dynamic_geometry_vertex_generation.cu‎

Lines changed: 108 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,108 @@
1+
//
2+
// Copyright (c) 2021, NVIDIA CORPORATION. All rights reserved.
3+
//
4+
// Redistribution and use in source and binary forms, with or without
5+
// modification, are permitted provided that the following conditions
6+
// are met:
7+
// * Redistributions of source code must retain the above copyright
8+
// notice, this list of conditions and the following disclaimer.
9+
// * Redistributions in binary form must reproduce the above copyright
10+
// notice, this list of conditions and the following disclaimer in the
11+
// documentation and/or other materials provided with the distribution.
12+
// * Neither the name of NVIDIA CORPORATION nor the names of its
13+
// contributors may be used to endorse or promote products derived
14+
// from this software without specific prior written permission.
15+
//
16+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
17+
// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
19+
// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
20+
// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
21+
// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
22+
// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
23+
// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
24+
// OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25+
// (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26+
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27+
//
28+
29+
#include "vec_math.h"
30+
31+
32+
enum struct AnimationMode: int
33+
{
34+
NONE = 0,
35+
DEFORM = 1,
36+
EXPLODE = 2,
37+
};
38+
39+
40+
__forceinline__ __device__ float triangle_wave( float x, float shift = 0.f, float period = 2.f * M_PIf, float amplitude = 1.f )
41+
{
42+
return fabsf( fmodf( ( 4.f / period ) * ( x - shift ), 4.f * amplitude ) - 2.f * amplitude ) - amplitude;
43+
}
44+
45+
__forceinline__ __device__ void write_animated_triangle( float3* out_vertices, int tidx, float3 v0, float3 v1, float3 v2, AnimationMode mode, float time )
46+
{
47+
float3 v = make_float3( 0 );
48+
49+
if( mode == AnimationMode::EXPLODE )
50+
{
51+
// Generate displacement vector from triangle index
52+
const float theta = ( (float)M_PIf * ( ( tidx + 1 ) * ( 13 / M_PIf ) ) );
53+
const float phi = ( (float)( 2.0 * M_PIf ) * ( ( tidx + 1 ) * ( 97 / M_PIf ) ) );
54+
55+
// Apply displacement to the sphere triangles
56+
v = make_float3( triangle_wave( phi ) * triangle_wave( theta, M_PIf / 2.f ),
57+
triangle_wave( phi, M_PIf / 2.f ) * triangle_wave( theta, M_PIf / 2.f ), triangle_wave( theta ) )
58+
* triangle_wave( time, M_PIf / 2.f ) * 2.f;
59+
}
60+
61+
out_vertices[tidx * 3 + 0] = v0 + v;
62+
out_vertices[tidx * 3 + 1] = v1 + v;
63+
out_vertices[tidx * 3 + 2] = v2 + v;
64+
}
65+
66+
__forceinline__ __device__ float3 deform_vertex( const float3& c, AnimationMode mode, float time )
67+
{
68+
// Apply sine wave to the y coordinate of the sphere vertices
69+
if( mode == AnimationMode::DEFORM )
70+
return make_float3( c.x, c.y * ( 0.5f + 0.4f * cosf( 4 * ( c.x + time ) ) ), c.z );
71+
return c;
72+
}
73+
74+
extern "C" __global__ void generate_vertices(float3* out_vertices, AnimationMode mode, float time, int width, int height)
75+
{
76+
int idx = blockIdx.x * blockDim.x + threadIdx.x;
77+
78+
if( idx < width * height )
79+
{
80+
// generate a single patch (two unindexed triangles) of a tessellated sphere
81+
82+
int x = idx % width;
83+
int y = idx / width;
84+
85+
const float theta0 = ( ( float )M_PIf * ( y + 0 ) ) / height;
86+
const float theta1 = ( ( float )M_PIf * ( y + 1 ) ) / height;
87+
const float phi0 = ( ( float )( 2.0 * M_PIf ) * ( x + 0 ) ) / width;
88+
const float phi1 = ( ( float )( 2.0 * M_PIf ) * ( x + 1 ) ) / width;
89+
90+
const float ct0 = cosf( theta0 );
91+
const float st0 = sinf( theta0 );
92+
const float ct1 = cosf( theta1 );
93+
const float st1 = sinf( theta1 );
94+
95+
const float cp0 = cosf( phi0 );
96+
const float sp0 = sinf( phi0 );
97+
const float cp1 = cosf( phi1 );
98+
const float sp1 = sinf( phi1 );
99+
100+
const float3 v00 = deform_vertex( make_float3( cp0 * st0, sp0 * st0, ct0 ), mode, time );
101+
const float3 v10 = deform_vertex( make_float3( cp0 * st1, sp0 * st1, ct1 ), mode, time );
102+
const float3 v01 = deform_vertex( make_float3( cp1 * st0, sp1 * st0, ct0 ), mode, time );
103+
const float3 v11 = deform_vertex( make_float3( cp1 * st1, sp1 * st1, ct1 ), mode, time );
104+
105+
write_animated_triangle( out_vertices, idx * 2 + 0, v00, v10, v11, mode, time );
106+
write_animated_triangle( out_vertices, idx * 2 + 1, v00, v11, v01, mode, time );
107+
}
108+
}

0 commit comments

Comments
 (0)