29#if !defined( __OPTIX_INCLUDE_INTERNAL_HEADERS__ )
30#error("optix_device_impl_transformations.h is an internal header file and must not be used directly. Please use optix_device.h or optix.h instead.")
33#ifndef OPTIX_OPTIX_DEVICE_IMPL_TRANSFORMATIONS_H
34#define OPTIX_OPTIX_DEVICE_IMPL_TRANSFORMATIONS_H
38static __forceinline__ __device__ float4
optixAddFloat4(
const float4& a,
const float4& b )
40 return make_float4( a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w );
43static __forceinline__ __device__ float4
optixMulFloat4(
const float4& a,
float b )
45 return make_float4( a.x * b, a.y * b, a.z * b, a.w * b );
48static __forceinline__ __device__ uint4
optixLdg(
unsigned long long addr )
51 asm volatile(
"cvta.to.global.u64 %0, %1;" :
"=l"( ptr ) :
"l"( addr ) );
53 asm volatile(
"ld.global.v4.u32 {%0,%1,%2,%3}, [%4];"
54 :
"=r"( ret.x ),
"=r"( ret.y ),
"=r"( ret.z ),
"=r"( ret.w )
65 for(
int ofs = 0; ofs <
sizeof( T ); ofs += 16 )
66 *(uint4*)( (
char*)&v + ofs ) =
optixLdg( (
unsigned long long)( (
char*)ptr + ofs ) );
71static __forceinline__ __device__ float4
optixMultiplyRowMatrix(
const float4 vec,
const float4 m0,
const float4 m1,
const float4 m2 )
75 result.x = vec.x * m0.x + vec.y * m1.x + vec.z * m2.x;
76 result.y = vec.x * m0.y + vec.y * m1.y + vec.z * m2.y;
77 result.z = vec.x * m0.z + vec.y * m1.z + vec.z * m2.z;
78 result.w = vec.x * m0.w + vec.y * m1.w + vec.z * m2.w + vec.w;
86 const float4 q = {srt.
qx, srt.
qy, srt.
qz, srt.
qw};
89 const float inv_sql = 1.f / ( srt.
qx * srt.
qx + srt.
qy * srt.
qy + srt.
qz * srt.
qz + srt.
qw * srt.
qw );
92 const float sqw = q.w * nq.w;
93 const float sqx = q.x * nq.x;
94 const float sqy = q.y * nq.y;
95 const float sqz = q.z * nq.z;
97 const float xy = q.x * nq.y;
98 const float zw = q.z * nq.w;
99 const float xz = q.x * nq.z;
100 const float yw = q.y * nq.w;
101 const float yz = q.y * nq.z;
102 const float xw = q.x * nq.w;
104 m0.x = ( sqx - sqy - sqz + sqw );
105 m0.y = 2.0f * ( xy - zw );
106 m0.z = 2.0f * ( xz + yw );
108 m1.x = 2.0f * ( xy + zw );
109 m1.y = ( -sqx + sqy - sqz + sqw );
110 m1.z = 2.0f * ( yz - xw );
112 m2.x = 2.0f * ( xz - yw );
113 m2.y = 2.0f * ( yz + xw );
114 m2.z = ( -sqx - sqy + sqz + sqw );
116 m0.w = m0.x * srt.
pvx + m0.y * srt.
pvy + m0.z * srt.
pvz + srt.
tx;
117 m1.w = m1.x * srt.
pvx + m1.y * srt.
pvy + m1.z * srt.
pvz + srt.
ty;
118 m2.w = m2.x * srt.
pvx + m2.y * srt.
pvy + m2.z * srt.
pvz + srt.
tz;
120 m0.z = m0.x * srt.
b + m0.y * srt.
c + m0.z * srt.
sz;
121 m1.z = m1.x * srt.
b + m1.y * srt.
c + m1.z * srt.
sz;
122 m2.z = m2.x * srt.
b + m2.y * srt.
c + m2.z * srt.
sz;
124 m0.y = m0.x * srt.
a + m0.y * srt.
sy;
125 m1.y = m1.x * srt.
a + m1.y * srt.
sy;
126 m2.y = m2.x * srt.
a + m2.y * srt.
sy;
128 m0.x = m0.x * srt.
sx;
129 m1.x = m1.x * srt.
sx;
130 m2.x = m2.x * srt.
sx;
137 m0.x * ( m1.y * m2.z - m1.z * m2.y ) - m0.y * ( m1.x * m2.z - m1.z * m2.x ) + m0.z * ( m1.x * m2.y - m1.y * m2.x );
139 const float inv_det3 = 1.0f / det3;
142 inv3[0][0] = inv_det3 * ( m1.y * m2.z - m2.y * m1.z );
143 inv3[0][1] = inv_det3 * ( m0.z * m2.y - m2.z * m0.y );
144 inv3[0][2] = inv_det3 * ( m0.y * m1.z - m1.y * m0.z );
146 inv3[1][0] = inv_det3 * ( m1.z * m2.x - m2.z * m1.x );
147 inv3[1][1] = inv_det3 * ( m0.x * m2.z - m2.x * m0.z );
148 inv3[1][2] = inv_det3 * ( m0.z * m1.x - m1.z * m0.x );
150 inv3[2][0] = inv_det3 * ( m1.x * m2.y - m2.x * m1.y );
151 inv3[2][1] = inv_det3 * ( m0.y * m2.x - m2.y * m0.x );
152 inv3[2][2] = inv_det3 * ( m0.x * m1.y - m1.x * m0.y );
154 const float b[3] = {m0.w, m1.w, m2.w};
159 m0.w = -inv3[0][0] * b[0] - inv3[0][1] * b[1] - inv3[0][2] * b[2];
164 m1.w = -inv3[1][0] * b[0] - inv3[1][1] * b[1] - inv3[1][2] * b[2];
169 m2.w = -inv3[2][0] * b[0] - inv3[2][1] * b[1] - inv3[2][2] * b[2];
181 const float t0 = 1.0f - t1;
203 const float t0 = 1.0f - t1;
209 float inv_length = 1.f / sqrt( srt2.y * srt2.y + srt2.z * srt2.z + srt2.w * srt2.w + srt3.x * srt3.x );
210 srt2.y *= inv_length;
211 srt2.z *= inv_length;
212 srt2.w *= inv_length;
213 srt3.x *= inv_length;
219 const float timeBegin = options.
timeBegin;
220 const float timeEnd = options.
timeEnd;
221 const float numIntervals = (float)( options.
numKeys - 1 );
227 float time = max( 0.f, min( numIntervals, numIntervals * __fdividef( globalt - timeBegin, timeEnd - timeBegin ) ) );
233 const float fltKey = fminf( floorf(time), numIntervals - 1 );
235 localt = time - fltKey;
252 const float4* transform = (
const float4*)( &transformData->
transform[key][0] );
271 const float4* dataPtr =
reinterpret_cast<const float4*
>( &transformData->
srtData[key] );
277 OptixSRTData srt = {data[0].x, data[0].y, data[0].z, data[0].w, data[1].x, data[1].y, data[1].z, data[1].w,
278 data[2].x, data[2].y, data[2].z, data[2].w, data[3].x, data[3].y, data[3].z, data[3].w};
290 const bool objectToWorld )
312 const float4* transform;
322 transform = (
const float4*)( ( objectToWorld ) ? traversable->
transform : traversable->
invTransform );
331 trf0 = {1.0f, 0.0f, 0.0f, 0.0f};
332 trf1 = {0.0f, 1.0f, 0.0f, 0.0f};
333 trf2 = {0.0f, 0.0f, 1.0f, 0.0f};
344 for(
unsigned int i = 0; i < size; ++i )
348 float4 trf0, trf1, trf2;
360 float4 tmp0 = m0, tmp1 = m1, tmp2 = m2;
375 for(
int i = size - 1; i >= 0; --i )
379 float4 trf0, trf1, trf2;
391 float4 tmp0 = m0, tmp1 = m1, tmp2 = m2;
400static __forceinline__ __device__ float3
optixTransformPoint(
const float4& m0,
const float4& m1,
const float4& m2,
const float3& p )
403 result.x = m0.x * p.x + m0.y * p.y + m0.z * p.z + m0.w;
404 result.y = m1.x * p.x + m1.y * p.y + m1.z * p.z + m1.w;
405 result.z = m2.x * p.x + m2.y * p.y + m2.z * p.z + m2.w;
410static __forceinline__ __device__ float3
optixTransformVector(
const float4& m0,
const float4& m1,
const float4& m2,
const float3& v )
413 result.x = m0.x * v.x + m0.y * v.y + m0.z * v.z;
414 result.y = m1.x * v.x + m1.y * v.y + m1.z * v.z;
415 result.z = m2.x * v.x + m2.y * v.y + m2.z * v.z;
421static __forceinline__ __device__ float3
optixTransformNormal(
const float4& m0,
const float4& m1,
const float4& m2,
const float3& n )
424 result.x = m0.x * n.x + m1.x * n.y + m2.x * n.z;
425 result.y = m0.y * n.x + m1.y * n.y + m2.y * n.z;
426 result.z = m0.z * n.x + m1.z * n.y + m2.z * n.z;
OptixTransformType
Transform.
Definition: optix_types.h:1855
unsigned long long OptixTraversableHandle
Traversable handle.
Definition: optix_types.h:77
@ OPTIX_TRANSFORM_TYPE_SRT_MOTION_TRANSFORM
Definition: optix_types.h:1859
@ OPTIX_TRANSFORM_TYPE_STATIC_TRANSFORM
Definition: optix_types.h:1857
@ OPTIX_TRANSFORM_TYPE_INSTANCE
Definition: optix_types.h:1860
@ OPTIX_TRANSFORM_TYPE_MATRIX_MOTION_TRANSFORM
Definition: optix_types.h:1858
Definition: optix_device_impl_transformations.h:36
static __forceinline__ __device__ void optixGetMatrixFromSrt(float4 &m0, float4 &m1, float4 &m2, const OptixSRTData &srt)
Definition: optix_device_impl_transformations.h:84
static __forceinline__ __device__ void optixResolveMotionKey(float &localt, int &key, const OptixMotionOptions &options, const float globalt)
Definition: optix_device_impl_transformations.h:217
static __forceinline__ __device__ float4 optixAddFloat4(const float4 &a, const float4 &b)
Definition: optix_device_impl_transformations.h:38
static __forceinline__ __device__ float4 optixMulFloat4(const float4 &a, float b)
Definition: optix_device_impl_transformations.h:43
static __forceinline__ __device__ T optixLoadReadOnlyAlign16(const T *ptr)
Definition: optix_device_impl_transformations.h:60
static __forceinline__ __device__ void optixLoadInterpolatedMatrixKey(float4 &m0, float4 &m1, float4 &m2, const float4 *matrix, const float t1)
Definition: optix_device_impl_transformations.h:172
static __forceinline__ __device__ void optixGetObjectToWorldTransformMatrix(float4 &m0, float4 &m1, float4 &m2)
Definition: optix_device_impl_transformations.h:369
static __forceinline__ __device__ float3 optixTransformPoint(const float4 &m0, const float4 &m1, const float4 &m2, const float3 &p)
Definition: optix_device_impl_transformations.h:400
static __forceinline__ __device__ void optixInvertMatrix(float4 &m0, float4 &m1, float4 &m2)
Definition: optix_device_impl_transformations.h:134
static __forceinline__ __device__ float3 optixTransformVector(const float4 &m0, const float4 &m1, const float4 &m2, const float3 &v)
Definition: optix_device_impl_transformations.h:410
static __forceinline__ __device__ void optixLoadInterpolatedSrtKey(float4 &srt0, float4 &srt1, float4 &srt2, float4 &srt3, const float4 *srt, const float t1)
Definition: optix_device_impl_transformations.h:188
static __forceinline__ __device__ float3 optixTransformNormal(const float4 &m0, const float4 &m1, const float4 &m2, const float3 &n)
Definition: optix_device_impl_transformations.h:421
static __forceinline__ __device__ uint4 optixLdg(unsigned long long addr)
Definition: optix_device_impl_transformations.h:48
static __forceinline__ __device__ float4 optixMultiplyRowMatrix(const float4 vec, const float4 m0, const float4 m1, const float4 m2)
Definition: optix_device_impl_transformations.h:71
static __forceinline__ __device__ void optixGetInterpolatedTransformationFromHandle(float4 &trf0, float4 &trf1, float4 &trf2, const OptixTraversableHandle handle, const float time, const bool objectToWorld)
Definition: optix_device_impl_transformations.h:285
static __forceinline__ __device__ void optixGetInterpolatedTransformation(float4 &trf0, float4 &trf1, float4 &trf2, const OptixMatrixMotionTransform *transformData, const float time)
Definition: optix_device_impl_transformations.h:240
static __forceinline__ __device__ void optixGetWorldToObjectTransformMatrix(float4 &m0, float4 &m1, float4 &m2)
Definition: optix_device_impl_transformations.h:338
static __forceinline__ __device__ const OptixStaticTransform * optixGetStaticTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1660
static __forceinline__ __device__ const OptixMatrixMotionTransform * optixGetMatrixMotionTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1674
static __forceinline__ __device__ OptixTraversableHandle optixGetTransformListHandle(unsigned int index)
Definition: optix_device_impl.h:1646
static __forceinline__ __device__ const float4 * optixGetInstanceTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1695
static __forceinline__ __device__ unsigned int optixGetTransformListSize()
Definition: optix_device_impl.h:1639
static __forceinline__ __device__ const OptixSRTMotionTransform * optixGetSRTMotionTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1667
static __forceinline__ __device__ float optixGetRayTime()
Definition: optix_device_impl.h:1298
static __forceinline__ __device__ OptixTransformType optixGetTransformTypeFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1653
static __forceinline__ __device__ const float4 * optixGetInstanceInverseTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1702
Motion options.
Definition: optix_types.h:1317
unsigned short numKeys
If numKeys > 1, motion is enabled. timeBegin, timeEnd and flags are all ignored when motion is disabl...
Definition: optix_types.h:1320
float timeBegin
Point in time where motion starts. Must be lesser than timeEnd.
Definition: optix_types.h:1326
float timeEnd
Point in time where motion ends. Must be greater than timeBegin.
Definition: optix_types.h:1329
Represents an SRT transformation.
Definition: optix_types.h:1500
float qz
Definition: optix_types.h:1503
float b
Definition: optix_types.h:1503
float tz
Definition: optix_types.h:1503
float qx
Definition: optix_types.h:1503
float sz
Definition: optix_types.h:1503
float c
Definition: optix_types.h:1503
float sy
Definition: optix_types.h:1503
float pvz
Definition: optix_types.h:1503
float qw
Definition: optix_types.h:1503
float ty
Definition: optix_types.h:1503
float pvy
Definition: optix_types.h:1503
float pvx
Definition: optix_types.h:1503
float sx
Definition: optix_types.h:1503
float a
Definition: optix_types.h:1503
float qy
Definition: optix_types.h:1503
float tx
Definition: optix_types.h:1503