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