NVIDIA OptiX 8.0 nvidia_logo_transpbg.gif Up
optix_device_impl.h
Go to the documentation of this file.
1/*
2* Copyright (c) 2023 NVIDIA Corporation. All rights reserved.
3*
4* NVIDIA Corporation and its licensors retain all intellectual property and proprietary
5* rights in and to this software, related documentation and any modifications thereto.
6* Any use, reproduction, disclosure or distribution of this software and related
7* documentation without an express license agreement from NVIDIA Corporation is strictly
8* prohibited.
9*
10* TO THE MAXIMUM EXTENT PERMITTED BY APPLICABLE LAW, THIS SOFTWARE IS PROVIDED *AS IS*
11* AND NVIDIA AND ITS SUPPLIERS DISCLAIM ALL WARRANTIES, EITHER EXPRESS OR IMPLIED,
12* INCLUDING, BUT NOT LIMITED TO, IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
13* PARTICULAR PURPOSE. IN NO EVENT SHALL NVIDIA OR ITS SUPPLIERS BE LIABLE FOR ANY
14* SPECIAL, INCIDENTAL, INDIRECT, OR CONSEQUENTIAL DAMAGES WHATSOEVER (INCLUDING, WITHOUT
15* LIMITATION, DAMAGES FOR LOSS OF BUSINESS PROFITS, BUSINESS INTERRUPTION, LOSS OF
16* BUSINESS INFORMATION, OR ANY OTHER PECUNIARY LOSS) ARISING OUT OF THE USE OF OR
17* INABILITY TO USE THIS SOFTWARE, EVEN IF NVIDIA HAS BEEN ADVISED OF THE POSSIBILITY OF
18* SUCH DAMAGES
19*/
20
29#if !defined( __OPTIX_INCLUDE_INTERNAL_HEADERS__ )
30#error("optix_device_impl.h is an internal header file and must not be used directly. Please use optix_device.h or optix.h instead.")
31#endif
32
33#ifndef OPTIX_OPTIX_DEVICE_IMPL_H
34#define OPTIX_OPTIX_DEVICE_IMPL_H
35
36#include "internal/optix_device_impl_transformations.h"
37
38#ifndef __CUDACC_RTC__
39#include <initializer_list>
40#include <type_traits>
41#endif
42
43namespace optix_internal {
44template <typename...>
45struct TypePack{};
46} // namespace optix_internal
47
48template <typename... Payload>
49static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
50 float3 rayOrigin,
51 float3 rayDirection,
52 float tmin,
53 float tmax,
54 float rayTime,
55 OptixVisibilityMask visibilityMask,
56 unsigned int rayFlags,
57 unsigned int SBToffset,
58 unsigned int SBTstride,
59 unsigned int missSBTIndex,
60 Payload&... payload )
61{
62 static_assert( sizeof...( Payload ) <= 32, "Only up to 32 payload values are allowed." );
63 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
64 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
65 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
66#ifndef __CUDACC_RTC__
67 static_assert( std::is_same<optix_internal::TypePack<unsigned int, Payload...>, optix_internal::TypePack<Payload..., unsigned int>>::value,
68 "All payload parameters need to be unsigned int." );
69#endif
70
71 OptixPayloadTypeID type = OPTIX_PAYLOAD_TYPE_DEFAULT;
72 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
73 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
74 unsigned int p[33] = { 0, payload... };
75 int payloadSize = (int)sizeof...( Payload );
76 asm volatile(
77 "call"
78 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26,%27,%28,%"
79 "29,%30,%31),"
80 "_optix_trace_typed_32,"
81 "(%32,%33,%34,%35,%36,%37,%38,%39,%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,%50,%51,%52,%53,%54,%55,%56,%57,%58,%"
82 "59,%60,%61,%62,%63,%64,%65,%66,%67,%68,%69,%70,%71,%72,%73,%74,%75,%76,%77,%78,%79,%80);"
83 : "=r"( p[1] ), "=r"( p[2] ), "=r"( p[3] ), "=r"( p[4] ), "=r"( p[5] ), "=r"( p[6] ), "=r"( p[7] ),
84 "=r"( p[8] ), "=r"( p[9] ), "=r"( p[10] ), "=r"( p[11] ), "=r"( p[12] ), "=r"( p[13] ), "=r"( p[14] ),
85 "=r"( p[15] ), "=r"( p[16] ), "=r"( p[17] ), "=r"( p[18] ), "=r"( p[19] ), "=r"( p[20] ), "=r"( p[21] ),
86 "=r"( p[22] ), "=r"( p[23] ), "=r"( p[24] ), "=r"( p[25] ), "=r"( p[26] ), "=r"( p[27] ), "=r"( p[28] ),
87 "=r"( p[29] ), "=r"( p[30] ), "=r"( p[31] ), "=r"( p[32] )
88 : "r"( type ), "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ),
89 "f"( tmax ), "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
90 "r"( missSBTIndex ), "r"( payloadSize ), "r"( p[1] ), "r"( p[2] ), "r"( p[3] ), "r"( p[4] ), "r"( p[5] ),
91 "r"( p[6] ), "r"( p[7] ), "r"( p[8] ), "r"( p[9] ), "r"( p[10] ), "r"( p[11] ), "r"( p[12] ), "r"( p[13] ),
92 "r"( p[14] ), "r"( p[15] ), "r"( p[16] ), "r"( p[17] ), "r"( p[18] ), "r"( p[19] ), "r"( p[20] ),
93 "r"( p[21] ), "r"( p[22] ), "r"( p[23] ), "r"( p[24] ), "r"( p[25] ), "r"( p[26] ), "r"( p[27] ),
94 "r"( p[28] ), "r"( p[29] ), "r"( p[30] ), "r"( p[31] ), "r"( p[32] )
95 : );
96 unsigned int index = 1;
97 (void)std::initializer_list<unsigned int>{index, ( payload = p[index++] )...};
98}
99
100template <typename... Payload>
101static __forceinline__ __device__ void optixTraverse( OptixTraversableHandle handle,
102 float3 rayOrigin,
103 float3 rayDirection,
104 float tmin,
105 float tmax,
106 float rayTime,
107 OptixVisibilityMask visibilityMask,
108 unsigned int rayFlags,
109 unsigned int SBToffset,
110 unsigned int SBTstride,
111 unsigned int missSBTIndex,
112 Payload&... payload )
113{
114 static_assert( sizeof...( Payload ) <= 32, "Only up to 32 payload values are allowed." );
115 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
116 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
117 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
118#ifndef __CUDACC_RTC__
119 static_assert( std::is_same<optix_internal::TypePack<unsigned int, Payload...>, optix_internal::TypePack<Payload..., unsigned int>>::value,
120 "All payload parameters need to be unsigned int." );
121#endif
122
123 OptixPayloadTypeID type = OPTIX_PAYLOAD_TYPE_DEFAULT;
124 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
125 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
126 unsigned int p[33] = {0, payload...};
127 int payloadSize = (int)sizeof...( Payload );
128 asm volatile(
129 "call"
130 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26,%27,%28,%"
131 "29,%30,%31),"
132 "_optix_hitobject_traverse,"
133 "(%32,%33,%34,%35,%36,%37,%38,%39,%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,%50,%51,%52,%53,%54,%55,%56,%57,%58,%"
134 "59,%60,%61,%62,%63,%64,%65,%66,%67,%68,%69,%70,%71,%72,%73,%74,%75,%76,%77,%78,%79,%80);"
135 : "=r"( p[1] ), "=r"( p[2] ), "=r"( p[3] ), "=r"( p[4] ), "=r"( p[5] ), "=r"( p[6] ), "=r"( p[7] ),
136 "=r"( p[8] ), "=r"( p[9] ), "=r"( p[10] ), "=r"( p[11] ), "=r"( p[12] ), "=r"( p[13] ), "=r"( p[14] ),
137 "=r"( p[15] ), "=r"( p[16] ), "=r"( p[17] ), "=r"( p[18] ), "=r"( p[19] ), "=r"( p[20] ), "=r"( p[21] ),
138 "=r"( p[22] ), "=r"( p[23] ), "=r"( p[24] ), "=r"( p[25] ), "=r"( p[26] ), "=r"( p[27] ), "=r"( p[28] ),
139 "=r"( p[29] ), "=r"( p[30] ), "=r"( p[31] ), "=r"( p[32] )
140 : "r"( type ), "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ),
141 "f"( tmax ), "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
142 "r"( missSBTIndex ), "r"( payloadSize ), "r"( p[1] ), "r"( p[2] ), "r"( p[3] ), "r"( p[4] ), "r"( p[5] ),
143 "r"( p[6] ), "r"( p[7] ), "r"( p[8] ), "r"( p[9] ), "r"( p[10] ), "r"( p[11] ), "r"( p[12] ), "r"( p[13] ),
144 "r"( p[14] ), "r"( p[15] ), "r"( p[16] ), "r"( p[17] ), "r"( p[18] ), "r"( p[19] ), "r"( p[20] ),
145 "r"( p[21] ), "r"( p[22] ), "r"( p[23] ), "r"( p[24] ), "r"( p[25] ), "r"( p[26] ), "r"( p[27] ),
146 "r"( p[28] ), "r"( p[29] ), "r"( p[30] ), "r"( p[31] ), "r"( p[32] )
147 : );
148 unsigned int index = 1;
149 (void)std::initializer_list<unsigned int>{index, ( payload = p[index++] )...};
150}
151
152template <typename... Payload>
153static __forceinline__ __device__ void optixTrace( OptixPayloadTypeID type,
154 OptixTraversableHandle handle,
155 float3 rayOrigin,
156 float3 rayDirection,
157 float tmin,
158 float tmax,
159 float rayTime,
160 OptixVisibilityMask visibilityMask,
161 unsigned int rayFlags,
162 unsigned int SBToffset,
163 unsigned int SBTstride,
164 unsigned int missSBTIndex,
165 Payload&... payload )
166{
167 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
168 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
169 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
170 static_assert( sizeof...( Payload ) <= 32, "Only up to 32 payload values are allowed." );
171#ifndef __CUDACC_RTC__
172 static_assert( std::is_same<optix_internal::TypePack<unsigned int, Payload...>, optix_internal::TypePack<Payload..., unsigned int>>::value,
173 "All payload parameters need to be unsigned int." );
174#endif
175
176 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
177 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
178 unsigned int p[33] = {0, payload...};
179 int payloadSize = (int)sizeof...( Payload );
180
181 asm volatile(
182 "call"
183 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26,%27,%28,%"
184 "29,%30,%31),"
185 "_optix_trace_typed_32,"
186 "(%32,%33,%34,%35,%36,%37,%38,%39,%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,%50,%51,%52,%53,%54,%55,%56,%57,%58,%"
187 "59,%60,%61,%62,%63,%64,%65,%66,%67,%68,%69,%70,%71,%72,%73,%74,%75,%76,%77,%78,%79,%80);"
188 : "=r"( p[1] ), "=r"( p[2] ), "=r"( p[3] ), "=r"( p[4] ), "=r"( p[5] ), "=r"( p[6] ), "=r"( p[7] ),
189 "=r"( p[8] ), "=r"( p[9] ), "=r"( p[10] ), "=r"( p[11] ), "=r"( p[12] ), "=r"( p[13] ), "=r"( p[14] ),
190 "=r"( p[15] ), "=r"( p[16] ), "=r"( p[17] ), "=r"( p[18] ), "=r"( p[19] ), "=r"( p[20] ), "=r"( p[21] ),
191 "=r"( p[22] ), "=r"( p[23] ), "=r"( p[24] ), "=r"( p[25] ), "=r"( p[26] ), "=r"( p[27] ), "=r"( p[28] ),
192 "=r"( p[29] ), "=r"( p[30] ), "=r"( p[31] ), "=r"( p[32] )
193 : "r"( type ), "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ),
194 "f"( tmax ), "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
195 "r"( missSBTIndex ), "r"( payloadSize ), "r"( p[1] ), "r"( p[2] ), "r"( p[3] ), "r"( p[4] ), "r"( p[5] ),
196 "r"( p[6] ), "r"( p[7] ), "r"( p[8] ), "r"( p[9] ), "r"( p[10] ), "r"( p[11] ), "r"( p[12] ), "r"( p[13] ),
197 "r"( p[14] ), "r"( p[15] ), "r"( p[16] ), "r"( p[17] ), "r"( p[18] ), "r"( p[19] ), "r"( p[20] ),
198 "r"( p[21] ), "r"( p[22] ), "r"( p[23] ), "r"( p[24] ), "r"( p[25] ), "r"( p[26] ), "r"( p[27] ),
199 "r"( p[28] ), "r"( p[29] ), "r"( p[30] ), "r"( p[31] ), "r"( p[32] )
200 : );
201 unsigned int index = 1;
202 (void)std::initializer_list<unsigned int>{index, ( payload = p[index++] )...};
203}
204
205template <typename... Payload>
206static __forceinline__ __device__ void optixTraverse( OptixPayloadTypeID type,
207 OptixTraversableHandle handle,
208 float3 rayOrigin,
209 float3 rayDirection,
210 float tmin,
211 float tmax,
212 float rayTime,
213 OptixVisibilityMask visibilityMask,
214 unsigned int rayFlags,
215 unsigned int SBToffset,
216 unsigned int SBTstride,
217 unsigned int missSBTIndex,
218 Payload&... payload )
219{
220 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
221 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
222 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
223 static_assert( sizeof...( Payload ) <= 32, "Only up to 32 payload values are allowed." );
224#ifndef __CUDACC_RTC__
225 static_assert( std::is_same<optix_internal::TypePack<unsigned int, Payload...>, optix_internal::TypePack<Payload..., unsigned int>>::value,
226 "All payload parameters need to be unsigned int." );
227#endif
228
229 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
230 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
231 unsigned int p[33] = {0, payload...};
232 int payloadSize = (int)sizeof...( Payload );
233 asm volatile(
234 "call"
235 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26,%27,%28,%"
236 "29,%30,%31),"
237 "_optix_hitobject_traverse,"
238 "(%32,%33,%34,%35,%36,%37,%38,%39,%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,%50,%51,%52,%53,%54,%55,%56,%57,%58,%"
239 "59,%60,%61,%62,%63,%64,%65,%66,%67,%68,%69,%70,%71,%72,%73,%74,%75,%76,%77,%78,%79,%80);"
240 : "=r"( p[1] ), "=r"( p[2] ), "=r"( p[3] ), "=r"( p[4] ), "=r"( p[5] ), "=r"( p[6] ), "=r"( p[7] ),
241 "=r"( p[8] ), "=r"( p[9] ), "=r"( p[10] ), "=r"( p[11] ), "=r"( p[12] ), "=r"( p[13] ), "=r"( p[14] ),
242 "=r"( p[15] ), "=r"( p[16] ), "=r"( p[17] ), "=r"( p[18] ), "=r"( p[19] ), "=r"( p[20] ), "=r"( p[21] ),
243 "=r"( p[22] ), "=r"( p[23] ), "=r"( p[24] ), "=r"( p[25] ), "=r"( p[26] ), "=r"( p[27] ), "=r"( p[28] ),
244 "=r"( p[29] ), "=r"( p[30] ), "=r"( p[31] ), "=r"( p[32] )
245 : "r"( type ), "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ),
246 "f"( tmax ), "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
247 "r"( missSBTIndex ), "r"( payloadSize ), "r"( p[1] ), "r"( p[2] ), "r"( p[3] ), "r"( p[4] ), "r"( p[5] ),
248 "r"( p[6] ), "r"( p[7] ), "r"( p[8] ), "r"( p[9] ), "r"( p[10] ), "r"( p[11] ), "r"( p[12] ), "r"( p[13] ),
249 "r"( p[14] ), "r"( p[15] ), "r"( p[16] ), "r"( p[17] ), "r"( p[18] ), "r"( p[19] ), "r"( p[20] ),
250 "r"( p[21] ), "r"( p[22] ), "r"( p[23] ), "r"( p[24] ), "r"( p[25] ), "r"( p[26] ), "r"( p[27] ),
251 "r"( p[28] ), "r"( p[29] ), "r"( p[30] ), "r"( p[31] ), "r"( p[32] )
252 : );
253 unsigned int index = 1;
254 (void)std::initializer_list<unsigned int>{index, ( payload = p[index++] )...};
255}
256
257static __forceinline__ __device__ void optixReorder( unsigned int coherenceHint, unsigned int numCoherenceHintBits )
258{
259 asm volatile(
260 "call"
261 "(),"
262 "_optix_hitobject_reorder,"
263 "(%0,%1);"
264 :
265 : "r"( coherenceHint ), "r"( numCoherenceHintBits )
266 : );
267}
268
269static __forceinline__ __device__ void optixReorder()
270{
271 unsigned int coherenceHint = 0;
272 unsigned int numCoherenceHintBits = 0;
273 asm volatile(
274 "call"
275 "(),"
276 "_optix_hitobject_reorder,"
277 "(%0,%1);"
278 :
279 : "r"( coherenceHint ), "r"( numCoherenceHintBits )
280 : );
281}
282
283template <typename... Payload>
284static __forceinline__ __device__ void optixInvoke( OptixPayloadTypeID type, Payload&... payload )
285{
286 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
287 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
288 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
289 static_assert( sizeof...( Payload ) <= 32, "Only up to 32 payload values are allowed." );
290#ifndef __CUDACC_RTC__
291 static_assert( std::is_same<optix_internal::TypePack<unsigned int, Payload...>, optix_internal::TypePack<Payload..., unsigned int>>::value,
292 "All payload parameters need to be unsigned int." );
293#endif
294
295 unsigned int p[33] = {0, payload...};
296 int payloadSize = (int)sizeof...( Payload );
297
298 asm volatile(
299 "call"
300 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26,%27,%28,%"
301 "29,%30,%31),"
302 "_optix_hitobject_invoke,"
303 "(%32,%33,%34,%35,%36,%37,%38,%39,%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,%50,%51,%52,%53,%54,%55,%56,%57,%58,%"
304 "59,%60,%61,%62,%63,%64,%65);"
305 : "=r"( p[1] ), "=r"( p[2] ), "=r"( p[3] ), "=r"( p[4] ), "=r"( p[5] ), "=r"( p[6] ), "=r"( p[7] ),
306 "=r"( p[8] ), "=r"( p[9] ), "=r"( p[10] ), "=r"( p[11] ), "=r"( p[12] ), "=r"( p[13] ), "=r"( p[14] ),
307 "=r"( p[15] ), "=r"( p[16] ), "=r"( p[17] ), "=r"( p[18] ), "=r"( p[19] ), "=r"( p[20] ), "=r"( p[21] ),
308 "=r"( p[22] ), "=r"( p[23] ), "=r"( p[24] ), "=r"( p[25] ), "=r"( p[26] ), "=r"( p[27] ), "=r"( p[28] ),
309 "=r"( p[29] ), "=r"( p[30] ), "=r"( p[31] ), "=r"( p[32] )
310 : "r"( type ), "r"( payloadSize ), "r"( p[1] ), "r"( p[2] ),
311 "r"( p[3] ), "r"( p[4] ), "r"( p[5] ), "r"( p[6] ), "r"( p[7] ), "r"( p[8] ), "r"( p[9] ), "r"( p[10] ),
312 "r"( p[11] ), "r"( p[12] ), "r"( p[13] ), "r"( p[14] ), "r"( p[15] ), "r"( p[16] ), "r"( p[17] ),
313 "r"( p[18] ), "r"( p[19] ), "r"( p[20] ), "r"( p[21] ), "r"( p[22] ), "r"( p[23] ), "r"( p[24] ),
314 "r"( p[25] ), "r"( p[26] ), "r"( p[27] ), "r"( p[28] ), "r"( p[29] ), "r"( p[30] ), "r"( p[31] ), "r"( p[32] )
315 : );
316
317 unsigned int index = 1;
318 (void)std::initializer_list<unsigned int>{index, ( payload = p[index++] )...};
319}
320
321template <typename... Payload>
322static __forceinline__ __device__ void optixInvoke( Payload&... payload )
323{
324 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
325 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
326 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
327 static_assert( sizeof...( Payload ) <= 32, "Only up to 32 payload values are allowed." );
328#ifndef __CUDACC_RTC__
329 static_assert( std::is_same<optix_internal::TypePack<unsigned int, Payload...>, optix_internal::TypePack<Payload..., unsigned int>>::value,
330 "All payload parameters need to be unsigned int." );
331#endif
332
333 OptixPayloadTypeID type = OPTIX_PAYLOAD_TYPE_DEFAULT;
334 unsigned int p[33] = {0, payload...};
335 int payloadSize = (int)sizeof...( Payload );
336
337 asm volatile(
338 "call"
339 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26,%27,%28,%"
340 "29,%30,%31),"
341 "_optix_hitobject_invoke,"
342 "(%32,%33,%34,%35,%36,%37,%38,%39,%40,%41,%42,%43,%44,%45,%46,%47,%48,%49,%50,%51,%52,%53,%54,%55,%56,%57,%58,%"
343 "59,%60,%61,%62,%63,%64,%65);"
344 : "=r"( p[1] ), "=r"( p[2] ), "=r"( p[3] ), "=r"( p[4] ), "=r"( p[5] ), "=r"( p[6] ), "=r"( p[7] ),
345 "=r"( p[8] ), "=r"( p[9] ), "=r"( p[10] ), "=r"( p[11] ), "=r"( p[12] ), "=r"( p[13] ), "=r"( p[14] ),
346 "=r"( p[15] ), "=r"( p[16] ), "=r"( p[17] ), "=r"( p[18] ), "=r"( p[19] ), "=r"( p[20] ), "=r"( p[21] ),
347 "=r"( p[22] ), "=r"( p[23] ), "=r"( p[24] ), "=r"( p[25] ), "=r"( p[26] ), "=r"( p[27] ), "=r"( p[28] ),
348 "=r"( p[29] ), "=r"( p[30] ), "=r"( p[31] ), "=r"( p[32] )
349 : "r"( type ), "r"( payloadSize ), "r"( p[1] ), "r"( p[2] ),
350 "r"( p[3] ), "r"( p[4] ), "r"( p[5] ), "r"( p[6] ), "r"( p[7] ), "r"( p[8] ), "r"( p[9] ), "r"( p[10] ),
351 "r"( p[11] ), "r"( p[12] ), "r"( p[13] ), "r"( p[14] ), "r"( p[15] ), "r"( p[16] ), "r"( p[17] ),
352 "r"( p[18] ), "r"( p[19] ), "r"( p[20] ), "r"( p[21] ), "r"( p[22] ), "r"( p[23] ), "r"( p[24] ),
353 "r"( p[25] ), "r"( p[26] ), "r"( p[27] ), "r"( p[28] ), "r"( p[29] ), "r"( p[30] ), "r"( p[31] ), "r"( p[32] )
354 : );
355
356 unsigned int index = 1;
357 (void)std::initializer_list<unsigned int>{index, ( payload = p[index++] )...};
358}
359
360template <typename... RegAttributes>
361static __forceinline__ __device__ void optixMakeHitObject( OptixTraversableHandle handle,
362 float3 rayOrigin,
363 float3 rayDirection,
364 float tmin,
365 float tmax,
366 float rayTime,
367 unsigned int sbtOffset,
368 unsigned int sbtStride,
369 unsigned int instIdx,
370 unsigned int sbtGASIdx,
371 unsigned int primIdx,
372 unsigned int hitKind,
373 RegAttributes... regAttributes )
374{
375 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
376 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
377 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
378 static_assert( sizeof...( RegAttributes ) <= 8, "Only up to 8 register attribute values are allowed." );
379#ifndef __CUDACC_RTC__
380 static_assert(
381 std::is_same<optix_internal::TypePack<unsigned int, RegAttributes...>, optix_internal::TypePack<RegAttributes..., unsigned int>>::value,
382 "All register attribute parameters need to be unsigned int." );
383#endif
384
385 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
386 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
387 unsigned int a[9] = {0, regAttributes...};
388 int attrSize = (int)sizeof...( RegAttributes );
389
390 OptixTraversableHandle* transforms = nullptr;
391 unsigned int numTransforms = 0;
392
393 asm volatile(
394 "call"
395 "(),"
396 "_optix_hitobject_make_hit,"
397 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26);"
398 :
399 : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
400 "f"( rayTime ), "r"( sbtOffset ), "r"( sbtStride ), "r"( instIdx ), "l"( transforms ), "r"( numTransforms ),
401 "r"( sbtGASIdx ), "r"( primIdx ), "r"( hitKind ), "r"( attrSize ), "r"( a[1] ), "r"( a[2] ), "r"( a[3] ),
402 "r"( a[4] ), "r"( a[5] ), "r"( a[6] ), "r"( a[7] ), "r"( a[8] )
403 : );
404}
405
406template <typename... RegAttributes>
407static __forceinline__ __device__ void optixMakeHitObject( OptixTraversableHandle handle,
408 float3 rayOrigin,
409 float3 rayDirection,
410 float tmin,
411 float tmax,
412 float rayTime,
413 unsigned int sbtOffset,
414 unsigned int sbtStride,
415 unsigned int instIdx,
416 const OptixTraversableHandle* transforms,
417 unsigned int numTransforms,
418 unsigned int sbtGASIdx,
419 unsigned int primIdx,
420 unsigned int hitKind,
421 RegAttributes... regAttributes )
422{
423 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
424 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
425 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
426 static_assert( sizeof...( RegAttributes ) <= 8, "Only up to 8 register attribute values are allowed." );
427#ifndef __CUDACC_RTC__
428 static_assert(
429 std::is_same<optix_internal::TypePack<unsigned int, RegAttributes...>, optix_internal::TypePack<RegAttributes..., unsigned int>>::value,
430 "All register attribute parameters need to be unsigned int." );
431#endif
432
433 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
434 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
435 unsigned int a[9] = {0, regAttributes...};
436 int attrSize = (int)sizeof...( RegAttributes );
437
438 asm volatile(
439 "call"
440 "(),"
441 "_optix_hitobject_make_hit,"
442 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25,%26);"
443 :
444 : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
445 "f"( rayTime ), "r"( sbtOffset ), "r"( sbtStride ), "r"( instIdx ), "l"( transforms ), "r"( numTransforms ),
446 "r"( sbtGASIdx ), "r"( primIdx ), "r"( hitKind ), "r"( attrSize ), "r"( a[1] ), "r"( a[2] ), "r"( a[3] ),
447 "r"( a[4] ), "r"( a[5] ), "r"( a[6] ), "r"( a[7] ), "r"( a[8] )
448 : );
449}
450
451template <typename... RegAttributes>
452static __forceinline__ __device__ void optixMakeHitObjectWithRecord( OptixTraversableHandle handle,
453 float3 rayOrigin,
454 float3 rayDirection,
455 float tmin,
456 float tmax,
457 float rayTime,
458 unsigned int sbtRecordIndex,
459 unsigned int instIdx,
460 const OptixTraversableHandle* transforms,
461 unsigned int numTransforms,
462 unsigned int sbtGASIdx,
463 unsigned int primIdx,
464 unsigned int hitKind,
465 RegAttributes... regAttributes )
466{
467 // std::is_same compares each type in the two TypePacks to make sure that all types are unsigned int.
468 // TypePack 1 unsigned int T0 T1 T2 ... Tn-1 Tn
469 // TypePack 2 T0 T1 T2 T3 ... Tn unsigned int
470 static_assert( sizeof...( RegAttributes ) <= 8, "Only up to 8 register attribute values are allowed." );
471#ifndef __CUDACC_RTC__
472 static_assert(
473 std::is_same<optix_internal::TypePack<unsigned int, RegAttributes...>, optix_internal::TypePack<RegAttributes..., unsigned int>>::value,
474 "All register attribute parameters need to be unsigned int." );
475#endif
476
477 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
478 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
479 unsigned int a[9] = {0, regAttributes...};
480 int attrSize = (int)sizeof...( RegAttributes );
481
482 asm volatile(
483 "call"
484 "(),"
485 "_optix_hitobject_make_hit_with_record,"
486 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9,%10,%11,%12,%13,%14,%15,%16,%17,%18,%19,%20,%21,%22,%23,%24,%25);"
487 :
488 : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
489 "f"( rayTime ), "r"( sbtRecordIndex ), "r"( instIdx ), "l"( transforms ), "r"( numTransforms ),
490 "r"( sbtGASIdx ), "r"( primIdx ), "r"( hitKind ), "r"( attrSize ), "r"( a[1] ), "r"( a[2] ), "r"( a[3] ),
491 "r"( a[4] ), "r"( a[5] ), "r"( a[6] ), "r"( a[7] ), "r"( a[8] )
492 : );
493}
494
495static __forceinline__ __device__ void optixMakeMissHitObject( unsigned int missSBTIndex,
496 float3 rayOrigin,
497 float3 rayDirection,
498 float tmin,
499 float tmax,
500 float rayTime )
501{
502 float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
503 float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
504
505 asm volatile(
506 "call"
507 "(),"
508 "_optix_hitobject_make_miss,"
509 "(%0,%1,%2,%3,%4,%5,%6,%7,%8,%9);"
510 :
511 : "r"( missSBTIndex ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ),
512 "f"( tmax ), "f"( rayTime )
513 : );
514}
515
516static __forceinline__ __device__ void optixMakeNopHitObject()
517{
518 asm volatile(
519 "call"
520 "(),"
521 "_optix_hitobject_make_nop,"
522 "();"
523 :
524 :
525 : );
526}
527
528static __forceinline__ __device__ bool optixHitObjectIsHit()
529{
530 unsigned int result;
531 asm volatile(
532 "call (%0), _optix_hitobject_is_hit,"
533 "();"
534 : "=r"( result )
535 :
536 : );
537 return result;
538}
539
540static __forceinline__ __device__ bool optixHitObjectIsMiss()
541{
542 unsigned int result;
543 asm volatile(
544 "call (%0), _optix_hitobject_is_miss,"
545 "();"
546 : "=r"( result )
547 :
548 : );
549 return result;
550}
551
552static __forceinline__ __device__ bool optixHitObjectIsNop()
553{
554 unsigned int result;
555 asm volatile(
556 "call (%0), _optix_hitobject_is_nop,"
557 "();"
558 : "=r"( result )
559 :
560 : );
561 return result;
562}
563
564static __forceinline__ __device__ unsigned int optixHitObjectGetInstanceId()
565{
566 unsigned int result;
567 asm volatile(
568 "call (%0), _optix_hitobject_get_instance_id,"
569 "();"
570 : "=r"( result )
571 :
572 : );
573 return result;
574}
575
576static __forceinline__ __device__ unsigned int optixHitObjectGetInstanceIndex()
577{
578 unsigned int result;
579 asm volatile(
580 "call (%0), _optix_hitobject_get_instance_idx,"
581 "();"
582 : "=r"( result )
583 :
584 : );
585 return result;
586}
587
588static __forceinline__ __device__ unsigned int optixHitObjectGetPrimitiveIndex()
589{
590 unsigned int result;
591 asm volatile(
592 "call (%0), _optix_hitobject_get_primitive_idx,"
593 "();"
594 : "=r"( result )
595 :
596 : );
597 return result;
598}
599
600static __forceinline__ __device__ unsigned int optixHitObjectGetTransformListSize()
601{
602 unsigned int result;
603 asm volatile(
604 "call (%0), _optix_hitobject_get_transform_list_size,"
605 "();"
606 : "=r"( result )
607 :
608 : );
609 return result;
610}
611
612static __forceinline__ __device__ OptixTraversableHandle optixHitObjectGetTransformListHandle( unsigned int index )
613{
614 unsigned long long result;
615 asm volatile(
616 "call (%0), _optix_hitobject_get_transform_list_handle,"
617 "(%1);"
618 : "=l"( result )
619 : "r"( index )
620 : );
621 return result;
622}
623
624static __forceinline__ __device__ unsigned int optixHitObjectGetSbtGASIndex()
625{
626 unsigned int result;
627 asm volatile(
628 "call (%0), _optix_hitobject_get_sbt_gas_idx,"
629 "();"
630 : "=r"( result )
631 :
632 : );
633 return result;
634}
635
636static __forceinline__ __device__ unsigned int optixHitObjectGetHitKind()
637{
638 unsigned int result;
639 asm volatile(
640 "call (%0), _optix_hitobject_get_hitkind,"
641 "();"
642 : "=r"( result )
643 :
644 : );
645 return result;
646}
647
648static __forceinline__ __device__ float3 optixHitObjectGetWorldRayOrigin()
649{
650 float x, y, z;
651 asm volatile(
652 "call (%0), _optix_hitobject_get_world_ray_origin_x,"
653 "();"
654 : "=f"( x )
655 :
656 : );
657 asm volatile(
658 "call (%0), _optix_hitobject_get_world_ray_origin_y,"
659 "();"
660 : "=f"( y )
661 :
662 : );
663 asm volatile(
664 "call (%0), _optix_hitobject_get_world_ray_origin_z,"
665 "();"
666 : "=f"( z )
667 :
668 : );
669 return make_float3( x, y, z );
670}
671
672static __forceinline__ __device__ float3 optixHitObjectGetWorldRayDirection()
673{
674 float x, y, z;
675 asm volatile(
676 "call (%0), _optix_hitobject_get_world_ray_direction_x,"
677 "();"
678 : "=f"( x )
679 :
680 : );
681 asm volatile(
682 "call (%0), _optix_hitobject_get_world_ray_direction_y,"
683 "();"
684 : "=f"( y )
685 :
686 : );
687 asm volatile(
688 "call (%0), _optix_hitobject_get_world_ray_direction_z,"
689 "();"
690 : "=f"( z )
691 :
692 : );
693 return make_float3( x, y, z );
694}
695
696static __forceinline__ __device__ float optixHitObjectGetRayTmin()
697{
698 float result;
699 asm volatile(
700 "call (%0), _optix_hitobject_get_ray_tmin,"
701 "();"
702 : "=f"( result )
703 :
704 : );
705 return result;
706}
707
708static __forceinline__ __device__ float optixHitObjectGetRayTmax()
709{
710 float result;
711 asm volatile(
712 "call (%0), _optix_hitobject_get_ray_tmax,"
713 "();"
714 : "=f"( result )
715 :
716 : );
717 return result;
718}
719
720static __forceinline__ __device__ float optixHitObjectGetRayTime()
721{
722 float result;
723 asm volatile(
724 "call (%0), _optix_hitobject_get_ray_time,"
725 "();"
726 : "=f"( result )
727 :
728 : );
729 return result;
730}
731
732static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_0()
733{
734 unsigned int ret;
735 asm volatile(
736 "call (%0), _optix_hitobject_get_attribute,"
737 "(%1);"
738 : "=r"( ret )
739 : "r"( 0 )
740 : );
741 return ret;
742}
743
744static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_1()
745{
746 unsigned int ret;
747 asm volatile(
748 "call (%0), _optix_hitobject_get_attribute,"
749 "(%1);"
750 : "=r"( ret )
751 : "r"( 1 )
752 : );
753 return ret;
754}
755
756static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_2()
757{
758 unsigned int ret;
759 asm volatile(
760 "call (%0), _optix_hitobject_get_attribute,"
761 "(%1);"
762 : "=r"( ret )
763 : "r"( 2 )
764 : );
765 return ret;
766}
767
768static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_3()
769{
770 unsigned int ret;
771 asm volatile(
772 "call (%0), _optix_hitobject_get_attribute,"
773 "(%1);"
774 : "=r"( ret )
775 : "r"( 3 )
776 : );
777 return ret;
778}
779
780static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_4()
781{
782 unsigned int ret;
783 asm volatile(
784 "call (%0), _optix_hitobject_get_attribute,"
785 "(%1);"
786 : "=r"( ret )
787 : "r"( 4 )
788 : );
789 return ret;
790}
791
792static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_5()
793{
794 unsigned int ret;
795 asm volatile(
796 "call (%0), _optix_hitobject_get_attribute,"
797 "(%1);"
798 : "=r"( ret )
799 : "r"( 5 )
800 : );
801 return ret;
802}
803
804static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_6()
805{
806 unsigned int ret;
807 asm volatile(
808 "call (%0), _optix_hitobject_get_attribute,"
809 "(%1);"
810 : "=r"( ret )
811 : "r"( 6 )
812 : );
813 return ret;
814}
815
816static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_7()
817{
818 unsigned int ret;
819 asm volatile(
820 "call (%0), _optix_hitobject_get_attribute,"
821 "(%1);"
822 : "=r"( ret )
823 : "r"( 7 )
824 : );
825 return ret;
826}
827
828static __forceinline__ __device__ unsigned int optixHitObjectGetSbtRecordIndex()
829{
830 unsigned int result;
831 asm volatile(
832 "call (%0), _optix_hitobject_get_sbt_record_index,"
833 "();"
834 : "=r"( result )
835 :
836 : );
837 return result;
838}
839
840static __forceinline__ __device__ CUdeviceptr optixHitObjectGetSbtDataPointer()
841{
842 unsigned long long ptr;
843 asm volatile(
844 "call (%0), _optix_hitobject_get_sbt_data_pointer,"
845 "();"
846 : "=l"( ptr )
847 :
848 : );
849 return ptr;
850}
851
852static __forceinline__ __device__ void optixSetPayload_0( unsigned int p )
853{
854 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 0 ), "r"( p ) : );
855}
856
857static __forceinline__ __device__ void optixSetPayload_1( unsigned int p )
858{
859 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 1 ), "r"( p ) : );
860}
861
862static __forceinline__ __device__ void optixSetPayload_2( unsigned int p )
863{
864 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 2 ), "r"( p ) : );
865}
866
867static __forceinline__ __device__ void optixSetPayload_3( unsigned int p )
868{
869 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 3 ), "r"( p ) : );
870}
871
872static __forceinline__ __device__ void optixSetPayload_4( unsigned int p )
873{
874 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 4 ), "r"( p ) : );
875}
876
877static __forceinline__ __device__ void optixSetPayload_5( unsigned int p )
878{
879 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 5 ), "r"( p ) : );
880}
881
882static __forceinline__ __device__ void optixSetPayload_6( unsigned int p )
883{
884 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 6 ), "r"( p ) : );
885}
886
887static __forceinline__ __device__ void optixSetPayload_7( unsigned int p )
888{
889 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 7 ), "r"( p ) : );
890}
891
892static __forceinline__ __device__ void optixSetPayload_8( unsigned int p )
893{
894 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 8 ), "r"( p ) : );
895}
896
897static __forceinline__ __device__ void optixSetPayload_9( unsigned int p )
898{
899 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 9 ), "r"( p ) : );
900}
901
902static __forceinline__ __device__ void optixSetPayload_10( unsigned int p )
903{
904 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 10 ), "r"( p ) : );
905}
906
907static __forceinline__ __device__ void optixSetPayload_11( unsigned int p )
908{
909 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 11 ), "r"( p ) : );
910}
911
912static __forceinline__ __device__ void optixSetPayload_12( unsigned int p )
913{
914 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 12 ), "r"( p ) : );
915}
916
917static __forceinline__ __device__ void optixSetPayload_13( unsigned int p )
918{
919 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 13 ), "r"( p ) : );
920}
921
922static __forceinline__ __device__ void optixSetPayload_14( unsigned int p )
923{
924 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 14 ), "r"( p ) : );
925}
926
927static __forceinline__ __device__ void optixSetPayload_15( unsigned int p )
928{
929 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 15 ), "r"( p ) : );
930}
931
932static __forceinline__ __device__ void optixSetPayload_16( unsigned int p )
933{
934 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 16 ), "r"( p ) : );
935}
936
937static __forceinline__ __device__ void optixSetPayload_17( unsigned int p )
938{
939 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 17 ), "r"( p ) : );
940}
941
942static __forceinline__ __device__ void optixSetPayload_18( unsigned int p )
943{
944 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 18 ), "r"( p ) : );
945}
946
947static __forceinline__ __device__ void optixSetPayload_19( unsigned int p )
948{
949 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 19 ), "r"( p ) : );
950}
951
952static __forceinline__ __device__ void optixSetPayload_20( unsigned int p )
953{
954 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 20 ), "r"( p ) : );
955}
956
957static __forceinline__ __device__ void optixSetPayload_21( unsigned int p )
958{
959 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 21 ), "r"( p ) : );
960}
961
962static __forceinline__ __device__ void optixSetPayload_22( unsigned int p )
963{
964 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 22 ), "r"( p ) : );
965}
966
967static __forceinline__ __device__ void optixSetPayload_23( unsigned int p )
968{
969 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 23 ), "r"( p ) : );
970}
971
972static __forceinline__ __device__ void optixSetPayload_24( unsigned int p )
973{
974 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 24 ), "r"( p ) : );
975}
976
977static __forceinline__ __device__ void optixSetPayload_25( unsigned int p )
978{
979 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 25 ), "r"( p ) : );
980}
981
982static __forceinline__ __device__ void optixSetPayload_26( unsigned int p )
983{
984 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 26 ), "r"( p ) : );
985}
986
987static __forceinline__ __device__ void optixSetPayload_27( unsigned int p )
988{
989 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 27 ), "r"( p ) : );
990}
991
992static __forceinline__ __device__ void optixSetPayload_28( unsigned int p )
993{
994 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 28 ), "r"( p ) : );
995}
996
997static __forceinline__ __device__ void optixSetPayload_29( unsigned int p )
998{
999 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 29 ), "r"( p ) : );
1000}
1001
1002static __forceinline__ __device__ void optixSetPayload_30( unsigned int p )
1003{
1004 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 30 ), "r"( p ) : );
1005}
1006
1007static __forceinline__ __device__ void optixSetPayload_31( unsigned int p )
1008{
1009 asm volatile( "call _optix_set_payload, (%0, %1);" : : "r"( 31 ), "r"( p ) : );
1010}
1011
1012static __forceinline__ __device__ unsigned int optixGetPayload_0()
1013{
1014 unsigned int result;
1015 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 0 ) : );
1016 return result;
1017}
1018
1019static __forceinline__ __device__ unsigned int optixGetPayload_1()
1020{
1021 unsigned int result;
1022 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 1 ) : );
1023 return result;
1024}
1025
1026static __forceinline__ __device__ unsigned int optixGetPayload_2()
1027{
1028 unsigned int result;
1029 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 2 ) : );
1030 return result;
1031}
1032
1033static __forceinline__ __device__ unsigned int optixGetPayload_3()
1034{
1035 unsigned int result;
1036 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 3 ) : );
1037 return result;
1038}
1039
1040static __forceinline__ __device__ unsigned int optixGetPayload_4()
1041{
1042 unsigned int result;
1043 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 4 ) : );
1044 return result;
1045}
1046
1047static __forceinline__ __device__ unsigned int optixGetPayload_5()
1048{
1049 unsigned int result;
1050 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 5 ) : );
1051 return result;
1052}
1053
1054static __forceinline__ __device__ unsigned int optixGetPayload_6()
1055{
1056 unsigned int result;
1057 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 6 ) : );
1058 return result;
1059}
1060
1061static __forceinline__ __device__ unsigned int optixGetPayload_7()
1062{
1063 unsigned int result;
1064 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 7 ) : );
1065 return result;
1066}
1067
1068static __forceinline__ __device__ unsigned int optixGetPayload_8()
1069{
1070 unsigned int result;
1071 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 8 ) : );
1072 return result;
1073}
1074
1075static __forceinline__ __device__ unsigned int optixGetPayload_9()
1076{
1077 unsigned int result;
1078 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 9 ) : );
1079 return result;
1080}
1081
1082static __forceinline__ __device__ unsigned int optixGetPayload_10()
1083{
1084 unsigned int result;
1085 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 10 ) : );
1086 return result;
1087}
1088
1089static __forceinline__ __device__ unsigned int optixGetPayload_11()
1090{
1091 unsigned int result;
1092 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 11 ) : );
1093 return result;
1094}
1095
1096static __forceinline__ __device__ unsigned int optixGetPayload_12()
1097{
1098 unsigned int result;
1099 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 12 ) : );
1100 return result;
1101}
1102
1103static __forceinline__ __device__ unsigned int optixGetPayload_13()
1104{
1105 unsigned int result;
1106 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 13 ) : );
1107 return result;
1108}
1109
1110static __forceinline__ __device__ unsigned int optixGetPayload_14()
1111{
1112 unsigned int result;
1113 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 14 ) : );
1114 return result;
1115}
1116
1117static __forceinline__ __device__ unsigned int optixGetPayload_15()
1118{
1119 unsigned int result;
1120 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 15 ) : );
1121 return result;
1122}
1123
1124static __forceinline__ __device__ unsigned int optixGetPayload_16()
1125{
1126 unsigned int result;
1127 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 16 ) : );
1128 return result;
1129}
1130
1131static __forceinline__ __device__ unsigned int optixGetPayload_17()
1132{
1133 unsigned int result;
1134 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 17 ) : );
1135 return result;
1136}
1137
1138static __forceinline__ __device__ unsigned int optixGetPayload_18()
1139{
1140 unsigned int result;
1141 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 18 ) : );
1142 return result;
1143}
1144
1145static __forceinline__ __device__ unsigned int optixGetPayload_19()
1146{
1147 unsigned int result;
1148 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 19 ) : );
1149 return result;
1150}
1151
1152static __forceinline__ __device__ unsigned int optixGetPayload_20()
1153{
1154 unsigned int result;
1155 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 20 ) : );
1156 return result;
1157}
1158
1159static __forceinline__ __device__ unsigned int optixGetPayload_21()
1160{
1161 unsigned int result;
1162 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 21 ) : );
1163 return result;
1164}
1165
1166static __forceinline__ __device__ unsigned int optixGetPayload_22()
1167{
1168 unsigned int result;
1169 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 22 ) : );
1170 return result;
1171}
1172
1173static __forceinline__ __device__ unsigned int optixGetPayload_23()
1174{
1175 unsigned int result;
1176 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 23 ) : );
1177 return result;
1178}
1179
1180static __forceinline__ __device__ unsigned int optixGetPayload_24()
1181{
1182 unsigned int result;
1183 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 24 ) : );
1184 return result;
1185}
1186
1187static __forceinline__ __device__ unsigned int optixGetPayload_25()
1188{
1189 unsigned int result;
1190 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 25 ) : );
1191 return result;
1192}
1193
1194static __forceinline__ __device__ unsigned int optixGetPayload_26()
1195{
1196 unsigned int result;
1197 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 26 ) : );
1198 return result;
1199}
1200
1201static __forceinline__ __device__ unsigned int optixGetPayload_27()
1202{
1203 unsigned int result;
1204 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 27 ) : );
1205 return result;
1206}
1207
1208static __forceinline__ __device__ unsigned int optixGetPayload_28()
1209{
1210 unsigned int result;
1211 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 28 ) : );
1212 return result;
1213}
1214
1215static __forceinline__ __device__ unsigned int optixGetPayload_29()
1216{
1217 unsigned int result;
1218 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 29 ) : );
1219 return result;
1220}
1221
1222static __forceinline__ __device__ unsigned int optixGetPayload_30()
1223{
1224 unsigned int result;
1225 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 30 ) : );
1226 return result;
1227}
1228
1229static __forceinline__ __device__ unsigned int optixGetPayload_31()
1230{
1231 unsigned int result;
1232 asm volatile( "call (%0), _optix_get_payload, (%1);" : "=r"( result ) : "r"( 31 ) : );
1233 return result;
1234}
1235
1236static __forceinline__ __device__ void optixSetPayloadTypes( unsigned int types )
1237{
1238 asm volatile( "call _optix_set_payload_types, (%0);" : : "r"( types ) : );
1239}
1240
1241static __forceinline__ __device__ unsigned int optixUndefinedValue()
1242{
1243 unsigned int u0;
1244 asm( "call (%0), _optix_undef_value, ();" : "=r"( u0 ) : );
1245 return u0;
1246}
1247
1248static __forceinline__ __device__ float3 optixGetWorldRayOrigin()
1249{
1250 float f0, f1, f2;
1251 asm( "call (%0), _optix_get_world_ray_origin_x, ();" : "=f"( f0 ) : );
1252 asm( "call (%0), _optix_get_world_ray_origin_y, ();" : "=f"( f1 ) : );
1253 asm( "call (%0), _optix_get_world_ray_origin_z, ();" : "=f"( f2 ) : );
1254 return make_float3( f0, f1, f2 );
1255}
1256
1257static __forceinline__ __device__ float3 optixGetWorldRayDirection()
1258{
1259 float f0, f1, f2;
1260 asm( "call (%0), _optix_get_world_ray_direction_x, ();" : "=f"( f0 ) : );
1261 asm( "call (%0), _optix_get_world_ray_direction_y, ();" : "=f"( f1 ) : );
1262 asm( "call (%0), _optix_get_world_ray_direction_z, ();" : "=f"( f2 ) : );
1263 return make_float3( f0, f1, f2 );
1264}
1265
1266static __forceinline__ __device__ float3 optixGetObjectRayOrigin()
1267{
1268 float f0, f1, f2;
1269 asm( "call (%0), _optix_get_object_ray_origin_x, ();" : "=f"( f0 ) : );
1270 asm( "call (%0), _optix_get_object_ray_origin_y, ();" : "=f"( f1 ) : );
1271 asm( "call (%0), _optix_get_object_ray_origin_z, ();" : "=f"( f2 ) : );
1272 return make_float3( f0, f1, f2 );
1273}
1274
1275static __forceinline__ __device__ float3 optixGetObjectRayDirection()
1276{
1277 float f0, f1, f2;
1278 asm( "call (%0), _optix_get_object_ray_direction_x, ();" : "=f"( f0 ) : );
1279 asm( "call (%0), _optix_get_object_ray_direction_y, ();" : "=f"( f1 ) : );
1280 asm( "call (%0), _optix_get_object_ray_direction_z, ();" : "=f"( f2 ) : );
1281 return make_float3( f0, f1, f2 );
1282}
1283
1284static __forceinline__ __device__ float optixGetRayTmin()
1285{
1286 float f0;
1287 asm( "call (%0), _optix_get_ray_tmin, ();" : "=f"( f0 ) : );
1288 return f0;
1289}
1290
1291static __forceinline__ __device__ float optixGetRayTmax()
1292{
1293 float f0;
1294 asm( "call (%0), _optix_get_ray_tmax, ();" : "=f"( f0 ) : );
1295 return f0;
1296}
1297
1298static __forceinline__ __device__ float optixGetRayTime()
1299{
1300 float f0;
1301 asm( "call (%0), _optix_get_ray_time, ();" : "=f"( f0 ) : );
1302 return f0;
1303}
1304
1305static __forceinline__ __device__ unsigned int optixGetRayFlags()
1306{
1307 unsigned int u0;
1308 asm( "call (%0), _optix_get_ray_flags, ();" : "=r"( u0 ) : );
1309 return u0;
1310}
1311
1312static __forceinline__ __device__ unsigned int optixGetRayVisibilityMask()
1313{
1314 unsigned int u0;
1315 asm( "call (%0), _optix_get_ray_visibility_mask, ();" : "=r"( u0 ) : );
1316 return u0;
1317}
1318
1319static __forceinline__ __device__ OptixTraversableHandle optixGetInstanceTraversableFromIAS( OptixTraversableHandle ias,
1320 unsigned int instIdx )
1321{
1322 unsigned long long handle;
1323 asm( "call (%0), _optix_get_instance_traversable_from_ias, (%1, %2);"
1324 : "=l"( handle ) : "l"( ias ), "r"( instIdx ) );
1325 return (OptixTraversableHandle)handle;
1326}
1327
1328
1329static __forceinline__ __device__ void optixGetTriangleVertexData( OptixTraversableHandle gas,
1330 unsigned int primIdx,
1331 unsigned int sbtGASIndex,
1332 float time,
1333 float3 data[3] )
1334{
1335 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8), _optix_get_triangle_vertex_data, "
1336 "(%9, %10, %11, %12);"
1337 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[1].x ), "=f"( data[1].y ),
1338 "=f"( data[1].z ), "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z )
1339 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1340 : );
1341}
1342
1343static __forceinline__ __device__ void optixGetMicroTriangleVertexData( float3 data[3] )
1344{
1345 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8), _optix_get_microtriangle_vertex_data, "
1346 "();"
1347 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[1].x ), "=f"( data[1].y ),
1348 "=f"( data[1].z ), "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z )
1349 : );
1350}
1351static __forceinline__ __device__ void optixGetMicroTriangleBarycentricsData( float2 data[3] )
1352{
1353 asm( "call (%0, %1, %2, %3, %4, %5), _optix_get_microtriangle_barycentrics_data, "
1354 "();"
1355 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[1].x ), "=f"( data[1].y ), "=f"( data[2].x ), "=f"( data[2].y )
1356 : );
1357}
1358
1359static __forceinline__ __device__ void optixGetLinearCurveVertexData( OptixTraversableHandle gas,
1360 unsigned int primIdx,
1361 unsigned int sbtGASIndex,
1362 float time,
1363 float4 data[2] )
1364{
1365 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7), _optix_get_linear_curve_vertex_data, "
1366 "(%8, %9, %10, %11);"
1367 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ),
1368 "=f"( data[1].x ), "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w )
1369 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1370 : );
1371}
1372
1373static __forceinline__ __device__ void optixGetQuadraticBSplineVertexData( OptixTraversableHandle gas,
1374 unsigned int primIdx,
1375 unsigned int sbtGASIndex,
1376 float time,
1377 float4 data[3] )
1378{
1379 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11), _optix_get_quadratic_bspline_vertex_data, "
1380 "(%12, %13, %14, %15);"
1381 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ),
1382 "=f"( data[1].x ), "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w ),
1383 "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z ), "=f"( data[2].w )
1384 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1385 : );
1386}
1387
1388static __forceinline__ __device__ void optixGetCubicBSplineVertexData( OptixTraversableHandle gas,
1389 unsigned int primIdx,
1390 unsigned int sbtGASIndex,
1391 float time,
1392 float4 data[4] )
1393{
1394 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15), "
1395 "_optix_get_cubic_bspline_vertex_data, "
1396 "(%16, %17, %18, %19);"
1397 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ),
1398 "=f"( data[1].x ), "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w ),
1399 "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z ), "=f"( data[2].w ),
1400 "=f"( data[3].x ), "=f"( data[3].y ), "=f"( data[3].z ), "=f"( data[3].w )
1401 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1402 : );
1403}
1404
1405static __forceinline__ __device__ void optixGetCatmullRomVertexData( OptixTraversableHandle gas,
1406 unsigned int primIdx,
1407 unsigned int sbtGASIndex,
1408 float time,
1409 float4 data[4] )
1410{
1411 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15), "
1412 "_optix_get_catmullrom_vertex_data, "
1413 "(%16, %17, %18, %19);"
1414 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ), "=f"( data[1].x ),
1415 "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w ), "=f"( data[2].x ), "=f"( data[2].y ),
1416 "=f"( data[2].z ), "=f"( data[2].w ), "=f"( data[3].x ), "=f"( data[3].y ), "=f"( data[3].z ), "=f"( data[3].w )
1417 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1418 : );
1419}
1420
1421static __forceinline__ __device__ void optixGetCubicBezierVertexData( OptixTraversableHandle gas,
1422 unsigned int primIdx,
1423 unsigned int sbtGASIndex,
1424 float time,
1425 float4 data[4] )
1426{
1427 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15), "
1428 "_optix_get_cubic_bezier_vertex_data, "
1429 "(%16, %17, %18, %19);"
1430 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ), "=f"( data[1].x ),
1431 "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w ), "=f"( data[2].x ), "=f"( data[2].y ),
1432 "=f"( data[2].z ), "=f"( data[2].w ), "=f"( data[3].x ), "=f"( data[3].y ), "=f"( data[3].z ), "=f"( data[3].w )
1433 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1434 : );
1435}
1436
1437static __forceinline__ __device__ void optixGetRibbonVertexData( OptixTraversableHandle gas,
1438 unsigned int primIdx,
1439 unsigned int sbtGASIndex,
1440 float time,
1441 float4 data[3] )
1442{
1443 asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11), _optix_get_ribbon_vertex_data, "
1444 "(%12, %13, %14, %15);"
1445 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ), "=f"( data[1].x ), "=f"( data[1].y ),
1446 "=f"( data[1].z ), "=f"( data[1].w ), "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z ), "=f"( data[2].w )
1447 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1448 : );
1449}
1450
1451static __forceinline__ __device__ float3 optixGetRibbonNormal( OptixTraversableHandle gas,
1452 unsigned int primIdx,
1453 unsigned int sbtGASIndex,
1454 float time,
1455 float2 ribbonParameters )
1456{
1457 float3 normal;
1458 asm( "call (%0, %1, %2), _optix_get_ribbon_normal, "
1459 "(%3, %4, %5, %6, %7, %8);"
1460 : "=f"( normal.x ), "=f"( normal.y ), "=f"( normal.z )
1461 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time ),
1462 "f"( ribbonParameters.x ), "f"( ribbonParameters.y )
1463 : );
1464 return normal;
1465}
1466
1467static __forceinline__ __device__ void optixGetSphereData( OptixTraversableHandle gas,
1468 unsigned int primIdx,
1469 unsigned int sbtGASIndex,
1470 float time,
1471 float4 data[1] )
1472{
1473 asm( "call (%0, %1, %2, %3), "
1474 "_optix_get_sphere_data, "
1475 "(%4, %5, %6, %7);"
1476 : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w )
1477 : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
1478 : );
1479}
1480
1481static __forceinline__ __device__ OptixTraversableHandle optixGetGASTraversableHandle()
1482{
1483 unsigned long long handle;
1484 asm( "call (%0), _optix_get_gas_traversable_handle, ();" : "=l"( handle ) : );
1485 return (OptixTraversableHandle)handle;
1486}
1487
1488static __forceinline__ __device__ float optixGetGASMotionTimeBegin( OptixTraversableHandle handle )
1489{
1490 float f0;
1491 asm( "call (%0), _optix_get_gas_motion_time_begin, (%1);" : "=f"( f0 ) : "l"( handle ) : );
1492 return f0;
1493}
1494
1495static __forceinline__ __device__ float optixGetGASMotionTimeEnd( OptixTraversableHandle handle )
1496{
1497 float f0;
1498 asm( "call (%0), _optix_get_gas_motion_time_end, (%1);" : "=f"( f0 ) : "l"( handle ) : );
1499 return f0;
1500}
1501
1502static __forceinline__ __device__ unsigned int optixGetGASMotionStepCount( OptixTraversableHandle handle )
1503{
1504 unsigned int u0;
1505 asm( "call (%0), _optix_get_gas_motion_step_count, (%1);" : "=r"( u0 ) : "l"( handle ) : );
1506 return u0;
1507}
1508
1509static __forceinline__ __device__ void optixGetWorldToObjectTransformMatrix( float m[12] )
1510{
1511 if( optixGetTransformListSize() == 0 )
1512 {
1513 m[0] = 1.0f;
1514 m[1] = 0.0f;
1515 m[2] = 0.0f;
1516 m[3] = 0.0f;
1517 m[4] = 0.0f;
1518 m[5] = 1.0f;
1519 m[6] = 0.0f;
1520 m[7] = 0.0f;
1521 m[8] = 0.0f;
1522 m[9] = 0.0f;
1523 m[10] = 1.0f;
1524 m[11] = 0.0f;
1525 return;
1526 }
1527
1528 float4 m0, m1, m2;
1529 optix_impl::optixGetWorldToObjectTransformMatrix( m0, m1, m2 );
1530 m[0] = m0.x;
1531 m[1] = m0.y;
1532 m[2] = m0.z;
1533 m[3] = m0.w;
1534 m[4] = m1.x;
1535 m[5] = m1.y;
1536 m[6] = m1.z;
1537 m[7] = m1.w;
1538 m[8] = m2.x;
1539 m[9] = m2.y;
1540 m[10] = m2.z;
1541 m[11] = m2.w;
1542}
1543
1544static __forceinline__ __device__ void optixGetObjectToWorldTransformMatrix( float m[12] )
1545{
1546 if( optixGetTransformListSize() == 0 )
1547 {
1548 m[0] = 1.0f;
1549 m[1] = 0.0f;
1550 m[2] = 0.0f;
1551 m[3] = 0.0f;
1552 m[4] = 0.0f;
1553 m[5] = 1.0f;
1554 m[6] = 0.0f;
1555 m[7] = 0.0f;
1556 m[8] = 0.0f;
1557 m[9] = 0.0f;
1558 m[10] = 1.0f;
1559 m[11] = 0.0f;
1560 return;
1561 }
1562
1563 float4 m0, m1, m2;
1564 optix_impl::optixGetObjectToWorldTransformMatrix( m0, m1, m2 );
1565 m[0] = m0.x;
1566 m[1] = m0.y;
1567 m[2] = m0.z;
1568 m[3] = m0.w;
1569 m[4] = m1.x;
1570 m[5] = m1.y;
1571 m[6] = m1.z;
1572 m[7] = m1.w;
1573 m[8] = m2.x;
1574 m[9] = m2.y;
1575 m[10] = m2.z;
1576 m[11] = m2.w;
1577}
1578
1579static __forceinline__ __device__ float3 optixTransformPointFromWorldToObjectSpace( float3 point )
1580{
1581 if( optixGetTransformListSize() == 0 )
1582 return point;
1583
1584 float4 m0, m1, m2;
1585 optix_impl::optixGetWorldToObjectTransformMatrix( m0, m1, m2 );
1586 return optix_impl::optixTransformPoint( m0, m1, m2, point );
1587}
1588
1589static __forceinline__ __device__ float3 optixTransformVectorFromWorldToObjectSpace( float3 vec )
1590{
1591 if( optixGetTransformListSize() == 0 )
1592 return vec;
1593
1594 float4 m0, m1, m2;
1595 optix_impl::optixGetWorldToObjectTransformMatrix( m0, m1, m2 );
1596 return optix_impl::optixTransformVector( m0, m1, m2, vec );
1597}
1598
1599static __forceinline__ __device__ float3 optixTransformNormalFromWorldToObjectSpace( float3 normal )
1600{
1601 if( optixGetTransformListSize() == 0 )
1602 return normal;
1603
1604 float4 m0, m1, m2;
1605 optix_impl::optixGetObjectToWorldTransformMatrix( m0, m1, m2 ); // inverse of optixGetWorldToObjectTransformMatrix()
1606 return optix_impl::optixTransformNormal( m0, m1, m2, normal );
1607}
1608
1609static __forceinline__ __device__ float3 optixTransformPointFromObjectToWorldSpace( float3 point )
1610{
1611 if( optixGetTransformListSize() == 0 )
1612 return point;
1613
1614 float4 m0, m1, m2;
1615 optix_impl::optixGetObjectToWorldTransformMatrix( m0, m1, m2 );
1616 return optix_impl::optixTransformPoint( m0, m1, m2, point );
1617}
1618
1619static __forceinline__ __device__ float3 optixTransformVectorFromObjectToWorldSpace( float3 vec )
1620{
1621 if( optixGetTransformListSize() == 0 )
1622 return vec;
1623
1624 float4 m0, m1, m2;
1625 optix_impl::optixGetObjectToWorldTransformMatrix( m0, m1, m2 );
1626 return optix_impl::optixTransformVector( m0, m1, m2, vec );
1627}
1628
1629static __forceinline__ __device__ float3 optixTransformNormalFromObjectToWorldSpace( float3 normal )
1630{
1631 if( optixGetTransformListSize() == 0 )
1632 return normal;
1633
1634 float4 m0, m1, m2;
1635 optix_impl::optixGetWorldToObjectTransformMatrix( m0, m1, m2 ); // inverse of optixGetObjectToWorldTransformMatrix()
1636 return optix_impl::optixTransformNormal( m0, m1, m2, normal );
1637}
1638
1639static __forceinline__ __device__ unsigned int optixGetTransformListSize()
1640{
1641 unsigned int u0;
1642 asm( "call (%0), _optix_get_transform_list_size, ();" : "=r"( u0 ) : );
1643 return u0;
1644}
1645
1646static __forceinline__ __device__ OptixTraversableHandle optixGetTransformListHandle( unsigned int index )
1647{
1648 unsigned long long u0;
1649 asm( "call (%0), _optix_get_transform_list_handle, (%1);" : "=l"( u0 ) : "r"( index ) : );
1650 return u0;
1651}
1652
1653static __forceinline__ __device__ OptixTransformType optixGetTransformTypeFromHandle( OptixTraversableHandle handle )
1654{
1655 int i0;
1656 asm( "call (%0), _optix_get_transform_type_from_handle, (%1);" : "=r"( i0 ) : "l"( handle ) : );
1657 return (OptixTransformType)i0;
1658}
1659
1660static __forceinline__ __device__ const OptixStaticTransform* optixGetStaticTransformFromHandle( OptixTraversableHandle handle )
1661{
1662 unsigned long long ptr;
1663 asm( "call (%0), _optix_get_static_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
1664 return (const OptixStaticTransform*)ptr;
1665}
1666
1667static __forceinline__ __device__ const OptixSRTMotionTransform* optixGetSRTMotionTransformFromHandle( OptixTraversableHandle handle )
1668{
1669 unsigned long long ptr;
1670 asm( "call (%0), _optix_get_srt_motion_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
1671 return (const OptixSRTMotionTransform*)ptr;
1672}
1673
1674static __forceinline__ __device__ const OptixMatrixMotionTransform* optixGetMatrixMotionTransformFromHandle( OptixTraversableHandle handle )
1675{
1676 unsigned long long ptr;
1677 asm( "call (%0), _optix_get_matrix_motion_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
1678 return (const OptixMatrixMotionTransform*)ptr;
1679}
1680
1681static __forceinline__ __device__ unsigned int optixGetInstanceIdFromHandle( OptixTraversableHandle handle )
1682{
1683 int i0;
1684 asm( "call (%0), _optix_get_instance_id_from_handle, (%1);" : "=r"( i0 ) : "l"( handle ) : );
1685 return i0;
1686}
1687
1688static __forceinline__ __device__ OptixTraversableHandle optixGetInstanceChildFromHandle( OptixTraversableHandle handle )
1689{
1690 unsigned long long i0;
1691 asm( "call (%0), _optix_get_instance_child_from_handle, (%1);" : "=l"( i0 ) : "l"( handle ) : );
1692 return (OptixTraversableHandle)i0;
1693}
1694
1695static __forceinline__ __device__ const float4* optixGetInstanceTransformFromHandle( OptixTraversableHandle handle )
1696{
1697 unsigned long long ptr;
1698 asm( "call (%0), _optix_get_instance_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
1699 return (const float4*)ptr;
1700}
1701
1702static __forceinline__ __device__ const float4* optixGetInstanceInverseTransformFromHandle( OptixTraversableHandle handle )
1703{
1704 unsigned long long ptr;
1705 asm( "call (%0), _optix_get_instance_inverse_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
1706 return (const float4*)ptr;
1707}
1708
1709static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind )
1710{
1711 int ret;
1712 asm volatile(
1713 "call (%0), _optix_report_intersection_0"
1714 ", (%1, %2);"
1715 : "=r"( ret )
1716 : "f"( hitT ), "r"( hitKind )
1717 : );
1718 return ret;
1719}
1720
1721static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind, unsigned int a0 )
1722{
1723 int ret;
1724 asm volatile(
1725 "call (%0), _optix_report_intersection_1"
1726 ", (%1, %2, %3);"
1727 : "=r"( ret )
1728 : "f"( hitT ), "r"( hitKind ), "r"( a0 )
1729 : );
1730 return ret;
1731}
1732
1733static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind, unsigned int a0, unsigned int a1 )
1734{
1735 int ret;
1736 asm volatile(
1737 "call (%0), _optix_report_intersection_2"
1738 ", (%1, %2, %3, %4);"
1739 : "=r"( ret )
1740 : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 )
1741 : );
1742 return ret;
1743}
1744
1745static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind, unsigned int a0, unsigned int a1, unsigned int a2 )
1746{
1747 int ret;
1748 asm volatile(
1749 "call (%0), _optix_report_intersection_3"
1750 ", (%1, %2, %3, %4, %5);"
1751 : "=r"( ret )
1752 : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 )
1753 : );
1754 return ret;
1755}
1756
1757static __forceinline__ __device__ bool optixReportIntersection( float hitT,
1758 unsigned int hitKind,
1759 unsigned int a0,
1760 unsigned int a1,
1761 unsigned int a2,
1762 unsigned int a3 )
1763{
1764 int ret;
1765 asm volatile(
1766 "call (%0), _optix_report_intersection_4"
1767 ", (%1, %2, %3, %4, %5, %6);"
1768 : "=r"( ret )
1769 : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 )
1770 : );
1771 return ret;
1772}
1773
1774static __forceinline__ __device__ bool optixReportIntersection( float hitT,
1775 unsigned int hitKind,
1776 unsigned int a0,
1777 unsigned int a1,
1778 unsigned int a2,
1779 unsigned int a3,
1780 unsigned int a4 )
1781{
1782 int ret;
1783 asm volatile(
1784 "call (%0), _optix_report_intersection_5"
1785 ", (%1, %2, %3, %4, %5, %6, %7);"
1786 : "=r"( ret )
1787 : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 )
1788 : );
1789 return ret;
1790}
1791
1792static __forceinline__ __device__ bool optixReportIntersection( float hitT,
1793 unsigned int hitKind,
1794 unsigned int a0,
1795 unsigned int a1,
1796 unsigned int a2,
1797 unsigned int a3,
1798 unsigned int a4,
1799 unsigned int a5 )
1800{
1801 int ret;
1802 asm volatile(
1803 "call (%0), _optix_report_intersection_6"
1804 ", (%1, %2, %3, %4, %5, %6, %7, %8);"
1805 : "=r"( ret )
1806 : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 ), "r"( a5 )
1807 : );
1808 return ret;
1809}
1810
1811static __forceinline__ __device__ bool optixReportIntersection( float hitT,
1812 unsigned int hitKind,
1813 unsigned int a0,
1814 unsigned int a1,
1815 unsigned int a2,
1816 unsigned int a3,
1817 unsigned int a4,
1818 unsigned int a5,
1819 unsigned int a6 )
1820{
1821 int ret;
1822 asm volatile(
1823 "call (%0), _optix_report_intersection_7"
1824 ", (%1, %2, %3, %4, %5, %6, %7, %8, %9);"
1825 : "=r"( ret )
1826 : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 ), "r"( a5 ), "r"( a6 )
1827 : );
1828 return ret;
1829}
1830
1831static __forceinline__ __device__ bool optixReportIntersection( float hitT,
1832 unsigned int hitKind,
1833 unsigned int a0,
1834 unsigned int a1,
1835 unsigned int a2,
1836 unsigned int a3,
1837 unsigned int a4,
1838 unsigned int a5,
1839 unsigned int a6,
1840 unsigned int a7 )
1841{
1842 int ret;
1843 asm volatile(
1844 "call (%0), _optix_report_intersection_8"
1845 ", (%1, %2, %3, %4, %5, %6, %7, %8, %9, %10);"
1846 : "=r"( ret )
1847 : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 ), "r"( a5 ), "r"( a6 ), "r"( a7 )
1848 : );
1849 return ret;
1850}
1851
1852#define OPTIX_DEFINE_optixGetAttribute_BODY( which ) \
1853 unsigned int ret; \
1854 asm( "call (%0), _optix_get_attribute_" #which ", ();" : "=r"( ret ) : ); \
1855 return ret;
1856
1857static __forceinline__ __device__ unsigned int optixGetAttribute_0()
1858{
1859 OPTIX_DEFINE_optixGetAttribute_BODY( 0 );
1860}
1861
1862static __forceinline__ __device__ unsigned int optixGetAttribute_1()
1863{
1864 OPTIX_DEFINE_optixGetAttribute_BODY( 1 );
1865}
1866
1867static __forceinline__ __device__ unsigned int optixGetAttribute_2()
1868{
1869 OPTIX_DEFINE_optixGetAttribute_BODY( 2 );
1870}
1871
1872static __forceinline__ __device__ unsigned int optixGetAttribute_3()
1873{
1874 OPTIX_DEFINE_optixGetAttribute_BODY( 3 );
1875}
1876
1877static __forceinline__ __device__ unsigned int optixGetAttribute_4()
1878{
1879 OPTIX_DEFINE_optixGetAttribute_BODY( 4 );
1880}
1881
1882static __forceinline__ __device__ unsigned int optixGetAttribute_5()
1883{
1884 OPTIX_DEFINE_optixGetAttribute_BODY( 5 );
1885}
1886
1887static __forceinline__ __device__ unsigned int optixGetAttribute_6()
1888{
1889 OPTIX_DEFINE_optixGetAttribute_BODY( 6 );
1890}
1891
1892static __forceinline__ __device__ unsigned int optixGetAttribute_7()
1893{
1894 OPTIX_DEFINE_optixGetAttribute_BODY( 7 );
1895}
1896
1897#undef OPTIX_DEFINE_optixGetAttribute_BODY
1898
1899static __forceinline__ __device__ void optixTerminateRay()
1900{
1901 asm volatile( "call _optix_terminate_ray, ();" );
1902}
1903
1904static __forceinline__ __device__ void optixIgnoreIntersection()
1905{
1906 asm volatile( "call _optix_ignore_intersection, ();" );
1907}
1908
1909static __forceinline__ __device__ unsigned int optixGetPrimitiveIndex()
1910{
1911 unsigned int u0;
1912 asm( "call (%0), _optix_read_primitive_idx, ();" : "=r"( u0 ) : );
1913 return u0;
1914}
1915
1916static __forceinline__ __device__ unsigned int optixGetSbtGASIndex()
1917{
1918 unsigned int u0;
1919 asm( "call (%0), _optix_read_sbt_gas_idx, ();" : "=r"( u0 ) : );
1920 return u0;
1921}
1922
1923static __forceinline__ __device__ unsigned int optixGetInstanceId()
1924{
1925 unsigned int u0;
1926 asm( "call (%0), _optix_read_instance_id, ();" : "=r"( u0 ) : );
1927 return u0;
1928}
1929
1930static __forceinline__ __device__ unsigned int optixGetInstanceIndex()
1931{
1932 unsigned int u0;
1933 asm( "call (%0), _optix_read_instance_idx, ();" : "=r"( u0 ) : );
1934 return u0;
1935}
1936
1937static __forceinline__ __device__ unsigned int optixGetHitKind()
1938{
1939 unsigned int u0;
1940 asm( "call (%0), _optix_get_hit_kind, ();" : "=r"( u0 ) : );
1941 return u0;
1942}
1943
1944static __forceinline__ __device__ OptixPrimitiveType optixGetPrimitiveType(unsigned int hitKind)
1945{
1946 unsigned int u0;
1947 asm( "call (%0), _optix_get_primitive_type_from_hit_kind, (%1);" : "=r"( u0 ) : "r"( hitKind ) );
1948 return (OptixPrimitiveType)u0;
1949}
1950
1951static __forceinline__ __device__ bool optixIsBackFaceHit( unsigned int hitKind )
1952{
1953 unsigned int u0;
1954 asm( "call (%0), _optix_get_backface_from_hit_kind, (%1);" : "=r"( u0 ) : "r"( hitKind ) );
1955 return (u0 == 0x1);
1956}
1957
1958static __forceinline__ __device__ bool optixIsFrontFaceHit( unsigned int hitKind )
1959{
1960 return !optixIsBackFaceHit( hitKind );
1961}
1962
1963
1964static __forceinline__ __device__ OptixPrimitiveType optixGetPrimitiveType()
1965{
1966 return optixGetPrimitiveType( optixGetHitKind() );
1967}
1968
1969static __forceinline__ __device__ bool optixIsBackFaceHit()
1970{
1971 return optixIsBackFaceHit( optixGetHitKind() );
1972}
1973
1974static __forceinline__ __device__ bool optixIsFrontFaceHit()
1975{
1976 return optixIsFrontFaceHit( optixGetHitKind() );
1977}
1978
1979static __forceinline__ __device__ bool optixIsTriangleHit()
1980{
1981 return optixIsTriangleFrontFaceHit() || optixIsTriangleBackFaceHit();
1982}
1983
1984static __forceinline__ __device__ bool optixIsTriangleFrontFaceHit()
1985{
1986 return optixGetHitKind() == OPTIX_HIT_KIND_TRIANGLE_FRONT_FACE;
1987}
1988
1989static __forceinline__ __device__ bool optixIsTriangleBackFaceHit()
1990{
1991 return optixGetHitKind() == OPTIX_HIT_KIND_TRIANGLE_BACK_FACE;
1992}
1993
1994static __forceinline__ __device__ bool optixIsDisplacedMicromeshTriangleHit()
1995{
1996 return optixGetPrimitiveType( optixGetHitKind() ) == OPTIX_PRIMITIVE_TYPE_DISPLACED_MICROMESH_TRIANGLE;
1997}
1998
1999static __forceinline__ __device__ bool optixIsDisplacedMicromeshTriangleFrontFaceHit()
2000{
2001 return optixIsDisplacedMicromeshTriangleHit() && optixIsFrontFaceHit();
2002}
2003
2004static __forceinline__ __device__ bool optixIsDisplacedMicromeshTriangleBackFaceHit()
2005{
2006 return optixIsDisplacedMicromeshTriangleHit() && optixIsBackFaceHit();
2007}
2008
2009static __forceinline__ __device__ float optixGetCurveParameter()
2010{
2011 float f0;
2012 asm( "call (%0), _optix_get_curve_parameter, ();" : "=f"(f0) : );
2013 return f0;
2014}
2015
2016static __forceinline__ __device__ float2 optixGetRibbonParameters()
2017{
2018 float f0, f1;
2019 asm( "call (%0, %1), _optix_get_ribbon_parameters, ();" : "=f"( f0 ), "=f"( f1 ) : );
2020 return make_float2( f0, f1 );
2021}
2022
2023static __forceinline__ __device__ float2 optixGetTriangleBarycentrics()
2024{
2025 float f0, f1;
2026 asm( "call (%0, %1), _optix_get_triangle_barycentrics, ();" : "=f"( f0 ), "=f"( f1 ) : );
2027 return make_float2( f0, f1 );
2028}
2029
2030static __forceinline__ __device__ uint3 optixGetLaunchIndex()
2031{
2032 unsigned int u0, u1, u2;
2033 asm( "call (%0), _optix_get_launch_index_x, ();" : "=r"( u0 ) : );
2034 asm( "call (%0), _optix_get_launch_index_y, ();" : "=r"( u1 ) : );
2035 asm( "call (%0), _optix_get_launch_index_z, ();" : "=r"( u2 ) : );
2036 return make_uint3( u0, u1, u2 );
2037}
2038
2039static __forceinline__ __device__ uint3 optixGetLaunchDimensions()
2040{
2041 unsigned int u0, u1, u2;
2042 asm( "call (%0), _optix_get_launch_dimension_x, ();" : "=r"( u0 ) : );
2043 asm( "call (%0), _optix_get_launch_dimension_y, ();" : "=r"( u1 ) : );
2044 asm( "call (%0), _optix_get_launch_dimension_z, ();" : "=r"( u2 ) : );
2045 return make_uint3( u0, u1, u2 );
2046}
2047
2048static __forceinline__ __device__ CUdeviceptr optixGetSbtDataPointer()
2049{
2050 unsigned long long ptr;
2051 asm( "call (%0), _optix_get_sbt_data_ptr_64, ();" : "=l"( ptr ) : );
2052 return (CUdeviceptr)ptr;
2053}
2054
2055static __forceinline__ __device__ void optixThrowException( int exceptionCode )
2056{
2057 asm volatile(
2058 "call _optix_throw_exception_0, (%0);"
2059 : /* no return value */
2060 : "r"( exceptionCode )
2061 : );
2062}
2063
2064static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0 )
2065{
2066 asm volatile(
2067 "call _optix_throw_exception_1, (%0, %1);"
2068 : /* no return value */
2069 : "r"( exceptionCode ), "r"( exceptionDetail0 )
2070 : );
2071}
2072
2073static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1 )
2074{
2075 asm volatile(
2076 "call _optix_throw_exception_2, (%0, %1, %2);"
2077 : /* no return value */
2078 : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 )
2079 : );
2080}
2081
2082static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2 )
2083{
2084 asm volatile(
2085 "call _optix_throw_exception_3, (%0, %1, %2, %3);"
2086 : /* no return value */
2087 : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 )
2088 : );
2089}
2090
2091static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3 )
2092{
2093 asm volatile(
2094 "call _optix_throw_exception_4, (%0, %1, %2, %3, %4);"
2095 : /* no return value */
2096 : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 )
2097 : );
2098}
2099
2100static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3, unsigned int exceptionDetail4 )
2101{
2102 asm volatile(
2103 "call _optix_throw_exception_5, (%0, %1, %2, %3, %4, %5);"
2104 : /* no return value */
2105 : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 )
2106 : );
2107}
2108
2109static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3, unsigned int exceptionDetail4, unsigned int exceptionDetail5 )
2110{
2111 asm volatile(
2112 "call _optix_throw_exception_6, (%0, %1, %2, %3, %4, %5, %6);"
2113 : /* no return value */
2114 : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 ), "r"( exceptionDetail5 )
2115 : );
2116}
2117
2118static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3, unsigned int exceptionDetail4, unsigned int exceptionDetail5, unsigned int exceptionDetail6 )
2119{
2120 asm volatile(
2121 "call _optix_throw_exception_7, (%0, %1, %2, %3, %4, %5, %6, %7);"
2122 : /* no return value */
2123 : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 ), "r"( exceptionDetail5 ), "r"( exceptionDetail6 )
2124 : );
2125}
2126
2127static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3, unsigned int exceptionDetail4, unsigned int exceptionDetail5, unsigned int exceptionDetail6, unsigned int exceptionDetail7 )
2128{
2129 asm volatile(
2130 "call _optix_throw_exception_8, (%0, %1, %2, %3, %4, %5, %6, %7, %8);"
2131 : /* no return value */
2132 : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 ), "r"( exceptionDetail5 ), "r"( exceptionDetail6 ), "r"( exceptionDetail7 )
2133 : );
2134}
2135
2136static __forceinline__ __device__ int optixGetExceptionCode()
2137{
2138 int s0;
2139 asm( "call (%0), _optix_get_exception_code, ();" : "=r"( s0 ) : );
2140 return s0;
2141}
2142
2143#define OPTIX_DEFINE_optixGetExceptionDetail_BODY( which ) \
2144 unsigned int ret; \
2145 asm( "call (%0), _optix_get_exception_detail_" #which ", ();" : "=r"( ret ) : ); \
2146 return ret;
2147
2148static __forceinline__ __device__ unsigned int optixGetExceptionDetail_0()
2149{
2150 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 0 );
2151}
2152
2153static __forceinline__ __device__ unsigned int optixGetExceptionDetail_1()
2154{
2155 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 1 );
2156}
2157
2158static __forceinline__ __device__ unsigned int optixGetExceptionDetail_2()
2159{
2160 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 2 );
2161}
2162
2163static __forceinline__ __device__ unsigned int optixGetExceptionDetail_3()
2164{
2165 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 3 );
2166}
2167
2168static __forceinline__ __device__ unsigned int optixGetExceptionDetail_4()
2169{
2170 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 4 );
2171}
2172
2173static __forceinline__ __device__ unsigned int optixGetExceptionDetail_5()
2174{
2175 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 5 );
2176}
2177
2178static __forceinline__ __device__ unsigned int optixGetExceptionDetail_6()
2179{
2180 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 6 );
2181}
2182
2183static __forceinline__ __device__ unsigned int optixGetExceptionDetail_7()
2184{
2185 OPTIX_DEFINE_optixGetExceptionDetail_BODY( 7 );
2186}
2187
2188#undef OPTIX_DEFINE_optixGetExceptionDetail_BODY
2189
2190
2191static __forceinline__ __device__ char* optixGetExceptionLineInfo()
2192{
2193 unsigned long long ptr;
2194 asm( "call (%0), _optix_get_exception_line_info, ();" : "=l"(ptr) : );
2195 return (char*)ptr;
2196}
2197
2198template <typename ReturnT, typename... ArgTypes>
2199static __forceinline__ __device__ ReturnT optixDirectCall( unsigned int sbtIndex, ArgTypes... args )
2200{
2201 unsigned long long func;
2202 asm( "call (%0), _optix_call_direct_callable,(%1);" : "=l"( func ) : "r"( sbtIndex ) : );
2203 using funcT = ReturnT ( * )( ArgTypes... );
2204 funcT call = ( funcT )( func );
2205 return call( args... );
2206}
2207
2208template <typename ReturnT, typename... ArgTypes>
2209static __forceinline__ __device__ ReturnT optixContinuationCall( unsigned int sbtIndex, ArgTypes... args )
2210{
2211 unsigned long long func;
2212 asm( "call (%0), _optix_call_continuation_callable,(%1);" : "=l"( func ) : "r"( sbtIndex ) : );
2213 using funcT = ReturnT ( * )( ArgTypes... );
2214 funcT call = ( funcT )( func );
2215 return call( args... );
2216}
2217
2218static __forceinline__ __device__ uint4 optixTexFootprint2D( unsigned long long tex, unsigned int texInfo, float x, float y, unsigned int* singleMipLevel )
2219{
2220 uint4 result;
2221 unsigned long long resultPtr = reinterpret_cast<unsigned long long>( &result );
2222 unsigned long long singleMipLevelPtr = reinterpret_cast<unsigned long long>( singleMipLevel );
2223 // Cast float args to integers, because the intrinics take .b32 arguments when compiled to PTX.
2224 asm volatile(
2225 "call _optix_tex_footprint_2d_v2"
2226 ", (%0, %1, %2, %3, %4, %5);"
2227 :
2228 : "l"( tex ), "r"( texInfo ), "r"( __float_as_uint( x ) ), "r"( __float_as_uint( y ) ),
2229 "l"( singleMipLevelPtr ), "l"( resultPtr )
2230 : );
2231 return result;
2232}
2233
2234static __forceinline__ __device__ uint4 optixTexFootprint2DGrad( unsigned long long tex,
2235 unsigned int texInfo,
2236 float x,
2237 float y,
2238 float dPdx_x,
2239 float dPdx_y,
2240 float dPdy_x,
2241 float dPdy_y,
2242 bool coarse,
2243 unsigned int* singleMipLevel )
2244{
2245 uint4 result;
2246 unsigned long long resultPtr = reinterpret_cast<unsigned long long>( &result );
2247 unsigned long long singleMipLevelPtr = reinterpret_cast<unsigned long long>( singleMipLevel );
2248 // Cast float args to integers, because the intrinics take .b32 arguments when compiled to PTX.
2249 asm volatile(
2250 "call _optix_tex_footprint_2d_grad_v2"
2251 ", (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10);"
2252 :
2253 : "l"( tex ), "r"( texInfo ), "r"( __float_as_uint( x ) ), "r"( __float_as_uint( y ) ),
2254 "r"( __float_as_uint( dPdx_x ) ), "r"( __float_as_uint( dPdx_y ) ), "r"( __float_as_uint( dPdy_x ) ),
2255 "r"( __float_as_uint( dPdy_y ) ), "r"( static_cast<unsigned int>( coarse ) ), "l"( singleMipLevelPtr ), "l"( resultPtr )
2256 : );
2257
2258 return result;
2259}
2260
2261static __forceinline__ __device__ uint4
2262optixTexFootprint2DLod( unsigned long long tex, unsigned int texInfo, float x, float y, float level, bool coarse, unsigned int* singleMipLevel )
2263{
2264 uint4 result;
2265 unsigned long long resultPtr = reinterpret_cast<unsigned long long>( &result );
2266 unsigned long long singleMipLevelPtr = reinterpret_cast<unsigned long long>( singleMipLevel );
2267 // Cast float args to integers, because the intrinics take .b32 arguments when compiled to PTX.
2268 asm volatile(
2269 "call _optix_tex_footprint_2d_lod_v2"
2270 ", (%0, %1, %2, %3, %4, %5, %6, %7);"
2271 :
2272 : "l"( tex ), "r"( texInfo ), "r"( __float_as_uint( x ) ), "r"( __float_as_uint( y ) ),
2273 "r"( __float_as_uint( level ) ), "r"( static_cast<unsigned int>( coarse ) ), "l"( singleMipLevelPtr ), "l"( resultPtr )
2274 : );
2275 return result;
2276}
2277
2278#endif // OPTIX_OPTIX_DEVICE_IMPL_H
OptixTransformType
OptixTransformType
Transform.
Definition: optix_types.h:1855
CUdeviceptr
unsigned long long CUdeviceptr
CUDA device pointer.
Definition: optix_types.h:52
OptixVisibilityMask
unsigned int OptixVisibilityMask
Visibility mask.
Definition: optix_types.h:80
OptixTraversableHandle
unsigned long long OptixTraversableHandle
Traversable handle.
Definition: optix_types.h:77
OptixPrimitiveType
OptixPrimitiveType
Builtin primitive types.
Definition: optix_types.h:727
OptixPayloadTypeID
OptixPayloadTypeID
Payload type identifiers.
Definition: optix_types.h:1982
OPTIX_HIT_KIND_TRIANGLE_BACK_FACE
@ OPTIX_HIT_KIND_TRIANGLE_BACK_FACE
Ray hit the triangle on the back face.
Definition: optix_types.h:342
OPTIX_HIT_KIND_TRIANGLE_FRONT_FACE
@ OPTIX_HIT_KIND_TRIANGLE_FRONT_FACE
Ray hit the triangle on the front face.
Definition: optix_types.h:340
OPTIX_PRIMITIVE_TYPE_DISPLACED_MICROMESH_TRIANGLE
@ OPTIX_PRIMITIVE_TYPE_DISPLACED_MICROMESH_TRIANGLE
Triangle with an applied displacement micromap.
Definition: optix_types.h:747
OPTIX_PAYLOAD_TYPE_DEFAULT
@ OPTIX_PAYLOAD_TYPE_DEFAULT
Definition: optix_types.h:1983
optix_impl::optixGetObjectToWorldTransformMatrix
static __forceinline__ __device__ void optixGetObjectToWorldTransformMatrix(float4 &m0, float4 &m1, float4 &m2)
Definition: optix_device_impl_transformations.h:369
optix_impl::optixTransformPoint
static __forceinline__ __device__ float3 optixTransformPoint(const float4 &m0, const float4 &m1, const float4 &m2, const float3 &p)
Definition: optix_device_impl_transformations.h:400
optix_impl::optixTransformVector
static __forceinline__ __device__ float3 optixTransformVector(const float4 &m0, const float4 &m1, const float4 &m2, const float3 &v)
Definition: optix_device_impl_transformations.h:410
optix_impl::optixTransformNormal
static __forceinline__ __device__ float3 optixTransformNormal(const float4 &m0, const float4 &m1, const float4 &m2, const float3 &n)
Definition: optix_device_impl_transformations.h:421
optix_impl::optixGetWorldToObjectTransformMatrix
static __forceinline__ __device__ void optixGetWorldToObjectTransformMatrix(float4 &m0, float4 &m1, float4 &m2)
Definition: optix_device_impl_transformations.h:338
optix_internal
Definition: optix_device_impl.h:43
optixGetExceptionDetail_7
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_7()
Definition: optix_device_impl.h:2183
optixContinuationCall
static __forceinline__ __device__ ReturnT optixContinuationCall(unsigned int sbtIndex, ArgTypes... args)
Definition: optix_device_impl.h:2209
optixThrowException
static __forceinline__ __device__ void optixThrowException(int exceptionCode)
Definition: optix_device_impl.h:2055
optixGetPayload_2
static __forceinline__ __device__ unsigned int optixGetPayload_2()
Definition: optix_device_impl.h:1026
optixSetPayload_16
static __forceinline__ __device__ void optixSetPayload_16(unsigned int p)
Definition: optix_device_impl.h:932
optixTransformNormalFromWorldToObjectSpace
static __forceinline__ __device__ float3 optixTransformNormalFromWorldToObjectSpace(float3 normal)
Definition: optix_device_impl.h:1599
optixGetPayload_25
static __forceinline__ __device__ unsigned int optixGetPayload_25()
Definition: optix_device_impl.h:1187
optixGetPayload_1
static __forceinline__ __device__ unsigned int optixGetPayload_1()
Definition: optix_device_impl.h:1019
optixGetPayload_15
static __forceinline__ __device__ unsigned int optixGetPayload_15()
Definition: optix_device_impl.h:1117
optixGetGASTraversableHandle
static __forceinline__ __device__ OptixTraversableHandle optixGetGASTraversableHandle()
Definition: optix_device_impl.h:1481
optixHitObjectGetSbtGASIndex
static __forceinline__ __device__ unsigned int optixHitObjectGetSbtGASIndex()
Definition: optix_device_impl.h:624
optixGetLaunchDimensions
static __forceinline__ __device__ uint3 optixGetLaunchDimensions()
Definition: optix_device_impl.h:2039
optixSetPayload_11
static __forceinline__ __device__ void optixSetPayload_11(unsigned int p)
Definition: optix_device_impl.h:907
optixSetPayloadTypes
static __forceinline__ __device__ void optixSetPayloadTypes(unsigned int types)
Definition: optix_device_impl.h:1236
optixSetPayload_8
static __forceinline__ __device__ void optixSetPayload_8(unsigned int p)
Definition: optix_device_impl.h:892
optixSetPayload_5
static __forceinline__ __device__ void optixSetPayload_5(unsigned int p)
Definition: optix_device_impl.h:877
optixGetPayload_9
static __forceinline__ __device__ unsigned int optixGetPayload_9()
Definition: optix_device_impl.h:1075
optixSetPayload_29
static __forceinline__ __device__ void optixSetPayload_29(unsigned int p)
Definition: optix_device_impl.h:997
optixSetPayload_14
static __forceinline__ __device__ void optixSetPayload_14(unsigned int p)
Definition: optix_device_impl.h:922
optixGetExceptionDetail_2
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_2()
Definition: optix_device_impl.h:2158
optixGetGASMotionStepCount
static __forceinline__ __device__ unsigned int optixGetGASMotionStepCount(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1502
optixGetStaticTransformFromHandle
static __forceinline__ __device__ const OptixStaticTransform * optixGetStaticTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1660
optixSetPayload_19
static __forceinline__ __device__ void optixSetPayload_19(unsigned int p)
Definition: optix_device_impl.h:947
optixGetWorldRayDirection
static __forceinline__ __device__ float3 optixGetWorldRayDirection()
Definition: optix_device_impl.h:1257
optixGetObjectRayOrigin
static __forceinline__ __device__ float3 optixGetObjectRayOrigin()
Definition: optix_device_impl.h:1266
optixIsFrontFaceHit
static __forceinline__ __device__ bool optixIsFrontFaceHit(unsigned int hitKind)
Definition: optix_device_impl.h:1958
optixGetAttribute_4
static __forceinline__ __device__ unsigned int optixGetAttribute_4()
Definition: optix_device_impl.h:1877
optixGetPayload_4
static __forceinline__ __device__ unsigned int optixGetPayload_4()
Definition: optix_device_impl.h:1040
optixSetPayload_10
static __forceinline__ __device__ void optixSetPayload_10(unsigned int p)
Definition: optix_device_impl.h:902
optixGetPayload_18
static __forceinline__ __device__ unsigned int optixGetPayload_18()
Definition: optix_device_impl.h:1138
optixHitObjectIsHit
static __forceinline__ __device__ bool optixHitObjectIsHit()
Definition: optix_device_impl.h:528
optixGetAttribute_6
static __forceinline__ __device__ unsigned int optixGetAttribute_6()
Definition: optix_device_impl.h:1887
optixMakeNopHitObject
static __forceinline__ __device__ void optixMakeNopHitObject()
Definition: optix_device_impl.h:516
optixGetAttribute_3
static __forceinline__ __device__ unsigned int optixGetAttribute_3()
Definition: optix_device_impl.h:1872
optixHitObjectGetAttribute_4
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_4()
Definition: optix_device_impl.h:780
optixSetPayload_20
static __forceinline__ __device__ void optixSetPayload_20(unsigned int p)
Definition: optix_device_impl.h:952
optixSetPayload_4
static __forceinline__ __device__ void optixSetPayload_4(unsigned int p)
Definition: optix_device_impl.h:872
optixGetPayload_5
static __forceinline__ __device__ unsigned int optixGetPayload_5()
Definition: optix_device_impl.h:1047
optixGetPayload_22
static __forceinline__ __device__ unsigned int optixGetPayload_22()
Definition: optix_device_impl.h:1166
optixIsTriangleHit
static __forceinline__ __device__ bool optixIsTriangleHit()
Definition: optix_device_impl.h:1979
optixGetCubicBSplineVertexData
static __forceinline__ __device__ void optixGetCubicBSplineVertexData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float4 data[4])
Definition: optix_device_impl.h:1388
optixGetPayload_27
static __forceinline__ __device__ unsigned int optixGetPayload_27()
Definition: optix_device_impl.h:1201
optixTexFootprint2DGrad
static __forceinline__ __device__ uint4 optixTexFootprint2DGrad(unsigned long long tex, unsigned int texInfo, float x, float y, float dPdx_x, float dPdx_y, float dPdy_x, float dPdy_y, bool coarse, unsigned int *singleMipLevel)
Definition: optix_device_impl.h:2234
optixSetPayload_6
static __forceinline__ __device__ void optixSetPayload_6(unsigned int p)
Definition: optix_device_impl.h:882
optixMakeHitObjectWithRecord
static __forceinline__ __device__ void optixMakeHitObjectWithRecord(OptixTraversableHandle handle, float3 rayOrigin, float3 rayDirection, float tmin, float tmax, float rayTime, unsigned int sbtRecordIndex, unsigned int instIdx, const OptixTraversableHandle *transforms, unsigned int numTransforms, unsigned int sbtGASIdx, unsigned int primIdx, unsigned int hitKind, RegAttributes... regAttributes)
Definition: optix_device_impl.h:452
optixHitObjectGetWorldRayDirection
static __forceinline__ __device__ float3 optixHitObjectGetWorldRayDirection()
Definition: optix_device_impl.h:672
optixGetAttribute_1
static __forceinline__ __device__ unsigned int optixGetAttribute_1()
Definition: optix_device_impl.h:1862
optixIsTriangleBackFaceHit
static __forceinline__ __device__ bool optixIsTriangleBackFaceHit()
Definition: optix_device_impl.h:1989
optixGetAttribute_5
static __forceinline__ __device__ unsigned int optixGetAttribute_5()
Definition: optix_device_impl.h:1882
optixHitObjectGetTransformListHandle
static __forceinline__ __device__ OptixTraversableHandle optixHitObjectGetTransformListHandle(unsigned int index)
Definition: optix_device_impl.h:612
optixGetRibbonVertexData
static __forceinline__ __device__ void optixGetRibbonVertexData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float4 data[3])
Definition: optix_device_impl.h:1437
optixSetPayload_25
static __forceinline__ __device__ void optixSetPayload_25(unsigned int p)
Definition: optix_device_impl.h:977
optixGetRayFlags
static __forceinline__ __device__ unsigned int optixGetRayFlags()
Definition: optix_device_impl.h:1305
optixHitObjectGetInstanceId
static __forceinline__ __device__ unsigned int optixHitObjectGetInstanceId()
Definition: optix_device_impl.h:564
optixHitObjectGetSbtRecordIndex
static __forceinline__ __device__ unsigned int optixHitObjectGetSbtRecordIndex()
Definition: optix_device_impl.h:828
optixTransformPointFromObjectToWorldSpace
static __forceinline__ __device__ float3 optixTransformPointFromObjectToWorldSpace(float3 point)
Definition: optix_device_impl.h:1609
optixGetPayload_26
static __forceinline__ __device__ unsigned int optixGetPayload_26()
Definition: optix_device_impl.h:1194
optixMakeMissHitObject
static __forceinline__ __device__ void optixMakeMissHitObject(unsigned int missSBTIndex, float3 rayOrigin, float3 rayDirection, float tmin, float tmax, float rayTime)
Definition: optix_device_impl.h:495
optixSetPayload_17
static __forceinline__ __device__ void optixSetPayload_17(unsigned int p)
Definition: optix_device_impl.h:937
optixSetPayload_7
static __forceinline__ __device__ void optixSetPayload_7(unsigned int p)
Definition: optix_device_impl.h:887
optixDirectCall
static __forceinline__ __device__ ReturnT optixDirectCall(unsigned int sbtIndex, ArgTypes... args)
Definition: optix_device_impl.h:2199
optixGetPayload_30
static __forceinline__ __device__ unsigned int optixGetPayload_30()
Definition: optix_device_impl.h:1222
optixSetPayload_28
static __forceinline__ __device__ void optixSetPayload_28(unsigned int p)
Definition: optix_device_impl.h:992
optixGetPrimitiveType
static __forceinline__ __device__ OptixPrimitiveType optixGetPrimitiveType(unsigned int hitKind)
Definition: optix_device_impl.h:1944
optixTransformVectorFromObjectToWorldSpace
static __forceinline__ __device__ float3 optixTransformVectorFromObjectToWorldSpace(float3 vec)
Definition: optix_device_impl.h:1619
optixGetInstanceId
static __forceinline__ __device__ unsigned int optixGetInstanceId()
Definition: optix_device_impl.h:1923
optixGetExceptionLineInfo
static __forceinline__ __device__ char * optixGetExceptionLineInfo()
Definition: optix_device_impl.h:2191
optixSetPayload_27
static __forceinline__ __device__ void optixSetPayload_27(unsigned int p)
Definition: optix_device_impl.h:987
optixSetPayload_2
static __forceinline__ __device__ void optixSetPayload_2(unsigned int p)
Definition: optix_device_impl.h:862
optixGetCatmullRomVertexData
static __forceinline__ __device__ void optixGetCatmullRomVertexData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float4 data[4])
Definition: optix_device_impl.h:1405
optixGetExceptionCode
static __forceinline__ __device__ int optixGetExceptionCode()
Definition: optix_device_impl.h:2136
optixSetPayload_12
static __forceinline__ __device__ void optixSetPayload_12(unsigned int p)
Definition: optix_device_impl.h:912
optixGetObjectToWorldTransformMatrix
static __forceinline__ __device__ void optixGetObjectToWorldTransformMatrix(float m[12])
Definition: optix_device_impl.h:1544
optixGetPayload_24
static __forceinline__ __device__ unsigned int optixGetPayload_24()
Definition: optix_device_impl.h:1180
optixGetPayload_17
static __forceinline__ __device__ unsigned int optixGetPayload_17()
Definition: optix_device_impl.h:1131
optixGetGASMotionTimeBegin
static __forceinline__ __device__ float optixGetGASMotionTimeBegin(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1488
optixHitObjectGetRayTmax
static __forceinline__ __device__ float optixHitObjectGetRayTmax()
Definition: optix_device_impl.h:708
optixGetMatrixMotionTransformFromHandle
static __forceinline__ __device__ const OptixMatrixMotionTransform * optixGetMatrixMotionTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1674
optixGetSphereData
static __forceinline__ __device__ void optixGetSphereData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float4 data[1])
Definition: optix_device_impl.h:1467
optixGetPayload_13
static __forceinline__ __device__ unsigned int optixGetPayload_13()
Definition: optix_device_impl.h:1103
optixTexFootprint2D
static __forceinline__ __device__ uint4 optixTexFootprint2D(unsigned long long tex, unsigned int texInfo, float x, float y, unsigned int *singleMipLevel)
Definition: optix_device_impl.h:2218
optixSetPayload_23
static __forceinline__ __device__ void optixSetPayload_23(unsigned int p)
Definition: optix_device_impl.h:967
optixGetAttribute_0
static __forceinline__ __device__ unsigned int optixGetAttribute_0()
Definition: optix_device_impl.h:1857
optixHitObjectGetAttribute_6
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_6()
Definition: optix_device_impl.h:804
optixHitObjectGetRayTmin
static __forceinline__ __device__ float optixHitObjectGetRayTmin()
Definition: optix_device_impl.h:696
optixGetPayload_10
static __forceinline__ __device__ unsigned int optixGetPayload_10()
Definition: optix_device_impl.h:1082
optixHitObjectGetAttribute_1
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_1()
Definition: optix_device_impl.h:744
optixGetQuadraticBSplineVertexData
static __forceinline__ __device__ void optixGetQuadraticBSplineVertexData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float4 data[3])
Definition: optix_device_impl.h:1373
optixGetLaunchIndex
static __forceinline__ __device__ uint3 optixGetLaunchIndex()
Definition: optix_device_impl.h:2030
optixGetPayload_31
static __forceinline__ __device__ unsigned int optixGetPayload_31()
Definition: optix_device_impl.h:1229
optixSetPayload_24
static __forceinline__ __device__ void optixSetPayload_24(unsigned int p)
Definition: optix_device_impl.h:972
optixTerminateRay
static __forceinline__ __device__ void optixTerminateRay()
Definition: optix_device_impl.h:1899
optixGetInstanceTraversableFromIAS
static __forceinline__ __device__ OptixTraversableHandle optixGetInstanceTraversableFromIAS(OptixTraversableHandle ias, unsigned int instIdx)
Definition: optix_device_impl.h:1319
optixHitObjectGetAttribute_0
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_0()
Definition: optix_device_impl.h:732
optixGetExceptionDetail_6
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_6()
Definition: optix_device_impl.h:2178
optixGetAttribute_7
static __forceinline__ __device__ unsigned int optixGetAttribute_7()
Definition: optix_device_impl.h:1892
optixHitObjectIsNop
static __forceinline__ __device__ bool optixHitObjectIsNop()
Definition: optix_device_impl.h:552
optixGetRayVisibilityMask
static __forceinline__ __device__ unsigned int optixGetRayVisibilityMask()
Definition: optix_device_impl.h:1312
optixHitObjectGetAttribute_2
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_2()
Definition: optix_device_impl.h:756
optixSetPayload_30
static __forceinline__ __device__ void optixSetPayload_30(unsigned int p)
Definition: optix_device_impl.h:1002
optixGetInstanceIdFromHandle
static __forceinline__ __device__ unsigned int optixGetInstanceIdFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1681
optixHitObjectGetInstanceIndex
static __forceinline__ __device__ unsigned int optixHitObjectGetInstanceIndex()
Definition: optix_device_impl.h:576
optixGetLinearCurveVertexData
static __forceinline__ __device__ void optixGetLinearCurveVertexData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float4 data[2])
Definition: optix_device_impl.h:1359
optixGetPayload_3
static __forceinline__ __device__ unsigned int optixGetPayload_3()
Definition: optix_device_impl.h:1033
optixHitObjectGetRayTime
static __forceinline__ __device__ float optixHitObjectGetRayTime()
Definition: optix_device_impl.h:720
optixReportIntersection
static __forceinline__ __device__ bool optixReportIntersection(float hitT, unsigned int hitKind)
Definition: optix_device_impl.h:1709
optixGetPayload_23
static __forceinline__ __device__ unsigned int optixGetPayload_23()
Definition: optix_device_impl.h:1173
optixGetPayload_20
static __forceinline__ __device__ unsigned int optixGetPayload_20()
Definition: optix_device_impl.h:1152
optixHitObjectGetAttribute_3
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_3()
Definition: optix_device_impl.h:768
optixGetTransformListHandle
static __forceinline__ __device__ OptixTraversableHandle optixGetTransformListHandle(unsigned int index)
Definition: optix_device_impl.h:1646
optixGetExceptionDetail_0
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_0()
Definition: optix_device_impl.h:2148
optixUndefinedValue
static __forceinline__ __device__ unsigned int optixUndefinedValue()
Definition: optix_device_impl.h:1241
OPTIX_DEFINE_optixGetAttribute_BODY
#define OPTIX_DEFINE_optixGetAttribute_BODY(which)
Definition: optix_device_impl.h:1852
optixGetRibbonParameters
static __forceinline__ __device__ float2 optixGetRibbonParameters()
Definition: optix_device_impl.h:2016
optixHitObjectGetTransformListSize
static __forceinline__ __device__ unsigned int optixHitObjectGetTransformListSize()
Definition: optix_device_impl.h:600
optixIgnoreIntersection
static __forceinline__ __device__ void optixIgnoreIntersection()
Definition: optix_device_impl.h:1904
optixGetRayTmax
static __forceinline__ __device__ float optixGetRayTmax()
Definition: optix_device_impl.h:1291
optixHitObjectGetSbtDataPointer
static __forceinline__ __device__ CUdeviceptr optixHitObjectGetSbtDataPointer()
Definition: optix_device_impl.h:840
optixTransformNormalFromObjectToWorldSpace
static __forceinline__ __device__ float3 optixTransformNormalFromObjectToWorldSpace(float3 normal)
Definition: optix_device_impl.h:1629
optixGetCubicBezierVertexData
static __forceinline__ __device__ void optixGetCubicBezierVertexData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float4 data[4])
Definition: optix_device_impl.h:1421
optixGetExceptionDetail_1
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_1()
Definition: optix_device_impl.h:2153
optixGetTriangleBarycentrics
static __forceinline__ __device__ float2 optixGetTriangleBarycentrics()
Definition: optix_device_impl.h:2023
optixHitObjectGetWorldRayOrigin
static __forceinline__ __device__ float3 optixHitObjectGetWorldRayOrigin()
Definition: optix_device_impl.h:648
optixGetObjectRayDirection
static __forceinline__ __device__ float3 optixGetObjectRayDirection()
Definition: optix_device_impl.h:1275
optixGetPayload_6
static __forceinline__ __device__ unsigned int optixGetPayload_6()
Definition: optix_device_impl.h:1054
optixSetPayload_9
static __forceinline__ __device__ void optixSetPayload_9(unsigned int p)
Definition: optix_device_impl.h:897
optixSetPayload_26
static __forceinline__ __device__ void optixSetPayload_26(unsigned int p)
Definition: optix_device_impl.h:982
optixGetPayload_11
static __forceinline__ __device__ unsigned int optixGetPayload_11()
Definition: optix_device_impl.h:1089
optixGetMicroTriangleVertexData
static __forceinline__ __device__ void optixGetMicroTriangleVertexData(float3 data[3])
Definition: optix_device_impl.h:1343
optixHitObjectGetHitKind
static __forceinline__ __device__ unsigned int optixHitObjectGetHitKind()
Definition: optix_device_impl.h:636
optixHitObjectIsMiss
static __forceinline__ __device__ bool optixHitObjectIsMiss()
Definition: optix_device_impl.h:540
optixGetWorldToObjectTransformMatrix
static __forceinline__ __device__ void optixGetWorldToObjectTransformMatrix(float m[12])
Definition: optix_device_impl.h:1509
optixGetInstanceIndex
static __forceinline__ __device__ unsigned int optixGetInstanceIndex()
Definition: optix_device_impl.h:1930
optixIsDisplacedMicromeshTriangleFrontFaceHit
static __forceinline__ __device__ bool optixIsDisplacedMicromeshTriangleFrontFaceHit()
Definition: optix_device_impl.h:1999
optixGetWorldRayOrigin
static __forceinline__ __device__ float3 optixGetWorldRayOrigin()
Definition: optix_device_impl.h:1248
optixSetPayload_22
static __forceinline__ __device__ void optixSetPayload_22(unsigned int p)
Definition: optix_device_impl.h:962
optixHitObjectGetAttribute_5
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_5()
Definition: optix_device_impl.h:792
optixTrace
static __forceinline__ __device__ void optixTrace(OptixTraversableHandle handle, float3 rayOrigin, float3 rayDirection, float tmin, float tmax, float rayTime, OptixVisibilityMask visibilityMask, unsigned int rayFlags, unsigned int SBToffset, unsigned int SBTstride, unsigned int missSBTIndex, Payload &... payload)
Definition: optix_device_impl.h:49
optixTexFootprint2DLod
static __forceinline__ __device__ uint4 optixTexFootprint2DLod(unsigned long long tex, unsigned int texInfo, float x, float y, float level, bool coarse, unsigned int *singleMipLevel)
Definition: optix_device_impl.h:2262
optixTraverse
static __forceinline__ __device__ void optixTraverse(OptixTraversableHandle handle, float3 rayOrigin, float3 rayDirection, float tmin, float tmax, float rayTime, OptixVisibilityMask visibilityMask, unsigned int rayFlags, unsigned int SBToffset, unsigned int SBTstride, unsigned int missSBTIndex, Payload &... payload)
Definition: optix_device_impl.h:101
optixGetPrimitiveIndex
static __forceinline__ __device__ unsigned int optixGetPrimitiveIndex()
Definition: optix_device_impl.h:1909
optixGetInstanceTransformFromHandle
static __forceinline__ __device__ const float4 * optixGetInstanceTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1695
optixSetPayload_3
static __forceinline__ __device__ void optixSetPayload_3(unsigned int p)
Definition: optix_device_impl.h:867
optixIsBackFaceHit
static __forceinline__ __device__ bool optixIsBackFaceHit(unsigned int hitKind)
Definition: optix_device_impl.h:1951
optixGetTransformListSize
static __forceinline__ __device__ unsigned int optixGetTransformListSize()
Definition: optix_device_impl.h:1639
optixSetPayload_18
static __forceinline__ __device__ void optixSetPayload_18(unsigned int p)
Definition: optix_device_impl.h:942
optixInvoke
static __forceinline__ __device__ void optixInvoke(OptixPayloadTypeID type, Payload &... payload)
Definition: optix_device_impl.h:284
optixIsDisplacedMicromeshTriangleHit
static __forceinline__ __device__ bool optixIsDisplacedMicromeshTriangleHit()
Definition: optix_device_impl.h:1994
optixGetSbtGASIndex
static __forceinline__ __device__ unsigned int optixGetSbtGASIndex()
Definition: optix_device_impl.h:1916
optixHitObjectGetPrimitiveIndex
static __forceinline__ __device__ unsigned int optixHitObjectGetPrimitiveIndex()
Definition: optix_device_impl.h:588
optixSetPayload_21
static __forceinline__ __device__ void optixSetPayload_21(unsigned int p)
Definition: optix_device_impl.h:957
optixIsDisplacedMicromeshTriangleBackFaceHit
static __forceinline__ __device__ bool optixIsDisplacedMicromeshTriangleBackFaceHit()
Definition: optix_device_impl.h:2004
optixSetPayload_0
static __forceinline__ __device__ void optixSetPayload_0(unsigned int p)
Definition: optix_device_impl.h:852
optixGetSRTMotionTransformFromHandle
static __forceinline__ __device__ const OptixSRTMotionTransform * optixGetSRTMotionTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1667
optixGetPayload_14
static __forceinline__ __device__ unsigned int optixGetPayload_14()
Definition: optix_device_impl.h:1110
optixGetSbtDataPointer
static __forceinline__ __device__ CUdeviceptr optixGetSbtDataPointer()
Definition: optix_device_impl.h:2048
optixGetRibbonNormal
static __forceinline__ __device__ float3 optixGetRibbonNormal(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float2 ribbonParameters)
Definition: optix_device_impl.h:1451
optixGetMicroTriangleBarycentricsData
static __forceinline__ __device__ void optixGetMicroTriangleBarycentricsData(float2 data[3])
Definition: optix_device_impl.h:1351
optixSetPayload_1
static __forceinline__ __device__ void optixSetPayload_1(unsigned int p)
Definition: optix_device_impl.h:857
optixSetPayload_13
static __forceinline__ __device__ void optixSetPayload_13(unsigned int p)
Definition: optix_device_impl.h:917
OPTIX_DEFINE_optixGetExceptionDetail_BODY
#define OPTIX_DEFINE_optixGetExceptionDetail_BODY(which)
Definition: optix_device_impl.h:2143
optixGetPayload_21
static __forceinline__ __device__ unsigned int optixGetPayload_21()
Definition: optix_device_impl.h:1159
optixGetPayload_8
static __forceinline__ __device__ unsigned int optixGetPayload_8()
Definition: optix_device_impl.h:1068
optixGetCurveParameter
static __forceinline__ __device__ float optixGetCurveParameter()
Definition: optix_device_impl.h:2009
optixGetPayload_19
static __forceinline__ __device__ unsigned int optixGetPayload_19()
Definition: optix_device_impl.h:1145
optixGetExceptionDetail_4
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_4()
Definition: optix_device_impl.h:2168
optixTransformVectorFromWorldToObjectSpace
static __forceinline__ __device__ float3 optixTransformVectorFromWorldToObjectSpace(float3 vec)
Definition: optix_device_impl.h:1589
optixSetPayload_15
static __forceinline__ __device__ void optixSetPayload_15(unsigned int p)
Definition: optix_device_impl.h:927
optixIsTriangleFrontFaceHit
static __forceinline__ __device__ bool optixIsTriangleFrontFaceHit()
Definition: optix_device_impl.h:1984
optixGetPayload_0
static __forceinline__ __device__ unsigned int optixGetPayload_0()
Definition: optix_device_impl.h:1012
optixTransformPointFromWorldToObjectSpace
static __forceinline__ __device__ float3 optixTransformPointFromWorldToObjectSpace(float3 point)
Definition: optix_device_impl.h:1579
optixGetAttribute_2
static __forceinline__ __device__ unsigned int optixGetAttribute_2()
Definition: optix_device_impl.h:1867
optixGetHitKind
static __forceinline__ __device__ unsigned int optixGetHitKind()
Definition: optix_device_impl.h:1937
optixGetPayload_28
static __forceinline__ __device__ unsigned int optixGetPayload_28()
Definition: optix_device_impl.h:1208
optixGetPayload_7
static __forceinline__ __device__ unsigned int optixGetPayload_7()
Definition: optix_device_impl.h:1061
optixGetTriangleVertexData
static __forceinline__ __device__ void optixGetTriangleVertexData(OptixTraversableHandle gas, unsigned int primIdx, unsigned int sbtGASIndex, float time, float3 data[3])
Definition: optix_device_impl.h:1329
optixGetGASMotionTimeEnd
static __forceinline__ __device__ float optixGetGASMotionTimeEnd(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1495
optixGetRayTmin
static __forceinline__ __device__ float optixGetRayTmin()
Definition: optix_device_impl.h:1284
optixReorder
static __forceinline__ __device__ void optixReorder(unsigned int coherenceHint, unsigned int numCoherenceHintBits)
Definition: optix_device_impl.h:257
optixGetRayTime
static __forceinline__ __device__ float optixGetRayTime()
Definition: optix_device_impl.h:1298
optixSetPayload_31
static __forceinline__ __device__ void optixSetPayload_31(unsigned int p)
Definition: optix_device_impl.h:1007
optixGetPayload_16
static __forceinline__ __device__ unsigned int optixGetPayload_16()
Definition: optix_device_impl.h:1124
optixGetTransformTypeFromHandle
static __forceinline__ __device__ OptixTransformType optixGetTransformTypeFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1653
optixMakeHitObject
static __forceinline__ __device__ void optixMakeHitObject(OptixTraversableHandle handle, float3 rayOrigin, float3 rayDirection, float tmin, float tmax, float rayTime, unsigned int sbtOffset, unsigned int sbtStride, unsigned int instIdx, unsigned int sbtGASIdx, unsigned int primIdx, unsigned int hitKind, RegAttributes... regAttributes)
Definition: optix_device_impl.h:361
optixGetInstanceInverseTransformFromHandle
static __forceinline__ __device__ const float4 * optixGetInstanceInverseTransformFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1702
optixGetPayload_12
static __forceinline__ __device__ unsigned int optixGetPayload_12()
Definition: optix_device_impl.h:1096
optixGetExceptionDetail_5
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_5()
Definition: optix_device_impl.h:2173
optixHitObjectGetAttribute_7
static __forceinline__ __device__ unsigned int optixHitObjectGetAttribute_7()
Definition: optix_device_impl.h:816
optixGetPayload_29
static __forceinline__ __device__ unsigned int optixGetPayload_29()
Definition: optix_device_impl.h:1215
optixGetInstanceChildFromHandle
static __forceinline__ __device__ OptixTraversableHandle optixGetInstanceChildFromHandle(OptixTraversableHandle handle)
Definition: optix_device_impl.h:1688
optixGetExceptionDetail_3
static __forceinline__ __device__ unsigned int optixGetExceptionDetail_3()
Definition: optix_device_impl.h:2163
optix_device_impl_transformations.h
OptiX public API.
OptixMatrixMotionTransform
Represents a matrix motion transformation.
Definition: optix_types.h:1454
OptixSRTMotionTransform
Represents an SRT motion transformation.
Definition: optix_types.h:1537
OptixStaticTransform
Static transform.
Definition: optix_types.h:1414
optix_internal::TypePack
Definition: optix_device_impl.h:45