NVIDIA OptiX 7.1 API nvidia_logo_transpbg.gif Up
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups
optix_7_device_impl.h
Go to the documentation of this file.
1 /*
2 * Copyright (c) 2020 NVIDIA Corporation. All rights reserved.
3 *
4 * NVIDIA Corporation and its licensors retain all intellectual property and proprietary
5 * rights in and to this software, related documentation and any modifications thereto.
6 * Any use, reproduction, disclosure or distribution of this software and related
7 * documentation without an express license agreement from NVIDIA Corporation is strictly
8 * prohibited.
9 *
10 * TO THE MAXIMUM EXTENT PERMITTED BY APPLICABLE LAW, THIS SOFTWARE IS PROVIDED *AS IS*
11 * AND NVIDIA AND ITS SUPPLIERS DISCLAIM ALL WARRANTIES, EITHER EXPRESS OR IMPLIED,
12 * INCLUDING, BUT NOT LIMITED TO, IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
13 * PARTICULAR PURPOSE. IN NO EVENT SHALL NVIDIA OR ITS SUPPLIERS BE LIABLE FOR ANY
14 * SPECIAL, INCIDENTAL, INDIRECT, OR CONSEQUENTIAL DAMAGES WHATSOEVER (INCLUDING, WITHOUT
15 * LIMITATION, DAMAGES FOR LOSS OF BUSINESS PROFITS, BUSINESS INTERRUPTION, LOSS OF
16 * BUSINESS INFORMATION, OR ANY OTHER PECUNIARY LOSS) ARISING OUT OF THE USE OF OR
17 * INABILITY TO USE THIS SOFTWARE, EVEN IF NVIDIA HAS BEEN ADVISED OF THE POSSIBILITY OF
18 * SUCH DAMAGES
19 */
20 
29 #if !defined( __OPTIX_INCLUDE_INTERNAL_HEADERS__ )
30 #error("optix_7_device_impl.h is an internal header file and must not be used directly. Please use optix_device.h or optix.h instead.")
31 #endif
32 
33 #ifndef __optix_optix_7_device_impl_h__
34 #define __optix_optix_7_device_impl_h__
35 
38 
39 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
40  float3 rayOrigin,
41  float3 rayDirection,
42  float tmin,
43  float tmax,
44  float rayTime,
45  OptixVisibilityMask visibilityMask,
46  unsigned int rayFlags,
47  unsigned int SBToffset,
48  unsigned int SBTstride,
49  unsigned int missSBTIndex )
50 {
51  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
52  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
53  asm volatile(
54  "call _optix_trace_0"
55  ", (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14"
56  ");"
57  :
58  /* no return value */
59  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
60  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ), "r"( missSBTIndex )
61  : );
62 }
63 
64 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
65  float3 rayOrigin,
66  float3 rayDirection,
67  float tmin,
68  float tmax,
69  float rayTime,
70  OptixVisibilityMask visibilityMask,
71  unsigned int rayFlags,
72  unsigned int SBToffset,
73  unsigned int SBTstride,
74  unsigned int missSBTIndex,
75  unsigned int& p0 )
76 {
77  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
78  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
79  unsigned int p0_out;
80  asm volatile(
81  "call (%0), _optix_trace_1"
82  ", (%1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15"
83  ", %16"
84  ");"
85  : "=r"( p0_out )
86  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ), "f"( rayTime ),
87  "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ), "r"( missSBTIndex ), "r"( p0 )
88  : );
89  p0 = p0_out;
90 }
91 
92 static __forceinline__ __device__ void optixTrace( 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  unsigned int& p0,
104  unsigned int& p1 )
105 {
106  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
107  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
108  unsigned int p0_out, p1_out;
109  asm volatile(
110  "call (%0, %1), _optix_trace_2"
111  ", (%2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15, %16"
112  ", %17, %18"
113  ");"
114  : "=r"( p0_out ), "=r"( p1_out )
115  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
116  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
117  "r"( missSBTIndex ), "r"( p0 ), "r"( p1 )
118  : );
119  p0 = p0_out;
120  p1 = p1_out;
121 }
122 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
123  float3 rayOrigin,
124  float3 rayDirection,
125  float tmin,
126  float tmax,
127  float rayTime,
128  OptixVisibilityMask visibilityMask,
129  unsigned int rayFlags,
130  unsigned int SBToffset,
131  unsigned int SBTstride,
132  unsigned int missSBTIndex,
133  unsigned int& p0,
134  unsigned int& p1,
135  unsigned int& p2 )
136 {
137  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
138  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
139  unsigned int p0_out, p1_out, p2_out;
140  asm volatile(
141  "call (%0, %1, %2), _optix_trace_3"
142  ", (%3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15, %16, %17"
143  ", %18, %19, %20"
144  ");"
145  : "=r"( p0_out ), "=r"( p1_out ), "=r"( p2_out )
146  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
147  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
148  "r"( missSBTIndex ), "r"( p0 ), "r"( p1 ), "r"( p2 )
149  : );
150  p0 = p0_out;
151  p1 = p1_out;
152  p2 = p2_out;
153 }
154 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
155  float3 rayOrigin,
156  float3 rayDirection,
157  float tmin,
158  float tmax,
159  float rayTime,
160  OptixVisibilityMask visibilityMask,
161  unsigned int rayFlags,
162  unsigned int SBToffset,
163  unsigned int SBTstride,
164  unsigned int missSBTIndex,
165  unsigned int& p0,
166  unsigned int& p1,
167  unsigned int& p2,
168  unsigned int& p3 )
169 {
170  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
171  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
172  unsigned int p0_out, p1_out, p2_out, p3_out;
173  asm volatile(
174  "call (%0, %1, %2, %3), _optix_trace_4"
175  ", (%4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15, %16, %17, %18"
176  ", %19, %20, %21, %22"
177  ");"
178  : "=r"( p0_out ), "=r"( p1_out ), "=r"( p2_out ), "=r"( p3_out )
179  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
180  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
181  "r"( missSBTIndex ), "r"( p0 ), "r"( p1 ), "r"( p2 ), "r"( p3 )
182  : );
183  p0 = p0_out;
184  p1 = p1_out;
185  p2 = p2_out;
186  p3 = p3_out;
187 }
188 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
189  float3 rayOrigin,
190  float3 rayDirection,
191  float tmin,
192  float tmax,
193  float rayTime,
194  OptixVisibilityMask visibilityMask,
195  unsigned int rayFlags,
196  unsigned int SBToffset,
197  unsigned int SBTstride,
198  unsigned int missSBTIndex,
199  unsigned int& p0,
200  unsigned int& p1,
201  unsigned int& p2,
202  unsigned int& p3,
203  unsigned int& p4 )
204 {
205  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
206  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
207  unsigned int p0_out, p1_out, p2_out, p3_out, p4_out;
208  asm volatile(
209  "call (%0, %1, %2, %3, %4), _optix_trace_5"
210  ", (%5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15, %16, %17, %18, %19"
211  ", %20, %21, %22, %23, %24"
212  ");"
213  : "=r"( p0_out ), "=r"( p1_out ), "=r"( p2_out ), "=r"( p3_out ), "=r"( p4_out )
214  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
215  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
216  "r"( missSBTIndex ), "r"( p0 ), "r"( p1 ), "r"( p2 ), "r"( p3 ), "r"( p4 )
217  : );
218  p0 = p0_out;
219  p1 = p1_out;
220  p2 = p2_out;
221  p3 = p3_out;
222  p4 = p4_out;
223 }
224 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
225  float3 rayOrigin,
226  float3 rayDirection,
227  float tmin,
228  float tmax,
229  float rayTime,
230  OptixVisibilityMask visibilityMask,
231  unsigned int rayFlags,
232  unsigned int SBToffset,
233  unsigned int SBTstride,
234  unsigned int missSBTIndex,
235  unsigned int& p0,
236  unsigned int& p1,
237  unsigned int& p2,
238  unsigned int& p3,
239  unsigned int& p4,
240  unsigned int& p5 )
241 {
242  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
243  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
244  unsigned int p0_out, p1_out, p2_out, p3_out, p4_out, p5_out;
245  asm volatile(
246  "call (%0, %1, %2, %3, %4, %5), _optix_trace_6"
247  ", (%6, %7, %8, %9, %10, %11, %12, %13, %14, %15, %16, %17, %18, %19, %20"
248  ", %21, %22, %23, %24, %25, %26"
249  ");"
250  : "=r"( p0_out ), "=r"( p1_out ), "=r"( p2_out ), "=r"( p3_out ), "=r"( p4_out ), "=r"( p5_out )
251  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
252  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
253  "r"( missSBTIndex ), "r"( p0 ), "r"( p1 ), "r"( p2 ), "r"( p3 ), "r"( p4 ), "r"( p5 )
254  : );
255  p0 = p0_out;
256  p1 = p1_out;
257  p2 = p2_out;
258  p3 = p3_out;
259  p4 = p4_out;
260  p5 = p5_out;
261 }
262 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
263  float3 rayOrigin,
264  float3 rayDirection,
265  float tmin,
266  float tmax,
267  float rayTime,
268  OptixVisibilityMask visibilityMask,
269  unsigned int rayFlags,
270  unsigned int SBToffset,
271  unsigned int SBTstride,
272  unsigned int missSBTIndex,
273  unsigned int& p0,
274  unsigned int& p1,
275  unsigned int& p2,
276  unsigned int& p3,
277  unsigned int& p4,
278  unsigned int& p5,
279  unsigned int& p6 )
280 {
281  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
282  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
283  unsigned int p0_out, p1_out, p2_out, p3_out, p4_out, p5_out, p6_out;
284  asm volatile(
285  "call (%0, %1, %2, %3, %4, %5, %6), _optix_trace_7"
286  ", (%7, %8, %9, %10, %11, %12, %13, %14, %15, %16, %17, %18, %19, %20, %21"
287  ", %22, %23, %24, %25, %26, %27, %28"
288  ");"
289  : "=r"( p0_out ), "=r"( p1_out ), "=r"( p2_out ), "=r"( p3_out ), "=r"( p4_out ), "=r"( p5_out ), "=r"( p6_out )
290  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
291  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
292  "r"( missSBTIndex ), "r"( p0 ), "r"( p1 ), "r"( p2 ), "r"( p3 ), "r"( p4 ), "r"( p5 ), "r"( p6 )
293  : );
294  p0 = p0_out;
295  p1 = p1_out;
296  p2 = p2_out;
297  p3 = p3_out;
298  p4 = p4_out;
299  p5 = p5_out;
300  p6 = p6_out;
301 }
302 static __forceinline__ __device__ void optixTrace( OptixTraversableHandle handle,
303  float3 rayOrigin,
304  float3 rayDirection,
305  float tmin,
306  float tmax,
307  float rayTime,
308  OptixVisibilityMask visibilityMask,
309  unsigned int rayFlags,
310  unsigned int SBToffset,
311  unsigned int SBTstride,
312  unsigned int missSBTIndex,
313  unsigned int& p0,
314  unsigned int& p1,
315  unsigned int& p2,
316  unsigned int& p3,
317  unsigned int& p4,
318  unsigned int& p5,
319  unsigned int& p6,
320  unsigned int& p7 )
321 {
322  float ox = rayOrigin.x, oy = rayOrigin.y, oz = rayOrigin.z;
323  float dx = rayDirection.x, dy = rayDirection.y, dz = rayDirection.z;
324  unsigned int p0_out, p1_out, p2_out, p3_out, p4_out, p5_out, p6_out, p7_out;
325  asm volatile(
326  "call (%0, %1, %2, %3, %4, %5, %6, %7), _optix_trace_8"
327  ", (%8, %9, %10, %11, %12, %13, %14, %15, %16, %17, %18, %19, %20, %21, %22"
328  ", %23, %24, %25, %26, %27, %28, %29, %30"
329  ");"
330  : "=r"( p0_out ), "=r"( p1_out ), "=r"( p2_out ), "=r"( p3_out ), "=r"( p4_out ), "=r"( p5_out ),
331  "=r"( p6_out ), "=r"( p7_out )
332  : "l"( handle ), "f"( ox ), "f"( oy ), "f"( oz ), "f"( dx ), "f"( dy ), "f"( dz ), "f"( tmin ), "f"( tmax ),
333  "f"( rayTime ), "r"( visibilityMask ), "r"( rayFlags ), "r"( SBToffset ), "r"( SBTstride ),
334  "r"( missSBTIndex ), "r"( p0 ), "r"( p1 ), "r"( p2 ), "r"( p3 ), "r"( p4 ), "r"( p5 ), "r"( p6 ), "r"( p7 )
335  : );
336  p0 = p0_out;
337  p1 = p1_out;
338  p2 = p2_out;
339  p3 = p3_out;
340  p4 = p4_out;
341  p5 = p5_out;
342  p6 = p6_out;
343  p7 = p7_out;
344 }
345 
346 #define OPTIX_DEFINE_optixSetPayload_BODY( which ) \
347  asm volatile( "call _optix_set_payload_" #which ", (%0);" : : "r"( p ) : );
348 
349 #define OPTIX_DEFINE_optixGetPayload_BODY( which ) \
350  unsigned int result; \
351  asm volatile( "call (%0), _optix_get_payload_" #which ", ();" : "=r"( result ) : ); \
352  return result;
353 
354 static __forceinline__ __device__ void optixSetPayload_0( unsigned int p )
355 {
357 }
358 
359 static __forceinline__ __device__ void optixSetPayload_1( unsigned int p )
360 {
362 }
363 
364 static __forceinline__ __device__ void optixSetPayload_2( unsigned int p )
365 {
367 }
368 
369 static __forceinline__ __device__ void optixSetPayload_3( unsigned int p )
370 {
372 }
373 
374 static __forceinline__ __device__ void optixSetPayload_4( unsigned int p )
375 {
377 }
378 
379 static __forceinline__ __device__ void optixSetPayload_5( unsigned int p )
380 {
382 }
383 
384 static __forceinline__ __device__ void optixSetPayload_6( unsigned int p )
385 {
387 }
388 
389 static __forceinline__ __device__ void optixSetPayload_7( unsigned int p )
390 {
392 }
393 
394 static __forceinline__ __device__ unsigned int optixGetPayload_0()
395 {
397 }
398 
399 static __forceinline__ __device__ unsigned int optixGetPayload_1()
400 {
402 }
403 
404 static __forceinline__ __device__ unsigned int optixGetPayload_2()
405 {
407 }
408 
409 static __forceinline__ __device__ unsigned int optixGetPayload_3()
410 {
412 }
413 
414 static __forceinline__ __device__ unsigned int optixGetPayload_4()
415 {
417 }
418 
419 static __forceinline__ __device__ unsigned int optixGetPayload_5()
420 {
422 }
423 
424 static __forceinline__ __device__ unsigned int optixGetPayload_6()
425 {
427 }
428 
429 static __forceinline__ __device__ unsigned int optixGetPayload_7()
430 {
432 }
433 
434 #undef OPTIX_DEFINE_optixSetPayload_BODY
435 #undef OPTIX_DEFINE_optixGetPayload_BODY
436 
437 static __forceinline__ __device__ unsigned int optixUndefinedValue()
438 {
439  unsigned int u0;
440  asm( "call (%0), _optix_undef_value, ();" : "=r"( u0 ) : );
441  return u0;
442 }
443 
444 static __forceinline__ __device__ float3 optixGetWorldRayOrigin()
445 {
446  float f0, f1, f2;
447  asm( "call (%0), _optix_get_world_ray_origin_x, ();" : "=f"( f0 ) : );
448  asm( "call (%0), _optix_get_world_ray_origin_y, ();" : "=f"( f1 ) : );
449  asm( "call (%0), _optix_get_world_ray_origin_z, ();" : "=f"( f2 ) : );
450  return make_float3( f0, f1, f2 );
451 }
452 
453 static __forceinline__ __device__ float3 optixGetWorldRayDirection()
454 {
455  float f0, f1, f2;
456  asm( "call (%0), _optix_get_world_ray_direction_x, ();" : "=f"( f0 ) : );
457  asm( "call (%0), _optix_get_world_ray_direction_y, ();" : "=f"( f1 ) : );
458  asm( "call (%0), _optix_get_world_ray_direction_z, ();" : "=f"( f2 ) : );
459  return make_float3( f0, f1, f2 );
460 }
461 
462 static __forceinline__ __device__ float3 optixGetObjectRayOrigin()
463 {
464  float f0, f1, f2;
465  asm( "call (%0), _optix_get_object_ray_origin_x, ();" : "=f"( f0 ) : );
466  asm( "call (%0), _optix_get_object_ray_origin_y, ();" : "=f"( f1 ) : );
467  asm( "call (%0), _optix_get_object_ray_origin_z, ();" : "=f"( f2 ) : );
468  return make_float3( f0, f1, f2 );
469 }
470 
471 static __forceinline__ __device__ float3 optixGetObjectRayDirection()
472 {
473  float f0, f1, f2;
474  asm( "call (%0), _optix_get_object_ray_direction_x, ();" : "=f"( f0 ) : );
475  asm( "call (%0), _optix_get_object_ray_direction_y, ();" : "=f"( f1 ) : );
476  asm( "call (%0), _optix_get_object_ray_direction_z, ();" : "=f"( f2 ) : );
477  return make_float3( f0, f1, f2 );
478 }
479 
480 static __forceinline__ __device__ float optixGetRayTmin()
481 {
482  float f0;
483  asm( "call (%0), _optix_get_ray_tmin, ();" : "=f"( f0 ) : );
484  return f0;
485 }
486 
487 static __forceinline__ __device__ float optixGetRayTmax()
488 {
489  float f0;
490  asm( "call (%0), _optix_get_ray_tmax, ();" : "=f"( f0 ) : );
491  return f0;
492 }
493 
494 static __forceinline__ __device__ float optixGetRayTime()
495 {
496  float f0;
497  asm( "call (%0), _optix_get_ray_time, ();" : "=f"( f0 ) : );
498  return f0;
499 }
500 
501 static __forceinline__ __device__ unsigned int optixGetRayFlags()
502 {
503  unsigned int u0;
504  asm( "call (%0), _optix_get_ray_flags, ();" : "=r"( u0 ) : );
505  return u0;
506 }
507 
508 static __forceinline__ __device__ unsigned int optixGetRayVisibilityMask()
509 {
510  unsigned int u0;
511  asm( "call (%0), _optix_get_ray_visibility_mask, ();" : "=r"( u0 ) : );
512  return u0;
513 }
514 
515 static __forceinline__ __device__ void optixGetTriangleVertexData( OptixTraversableHandle gas,
516  unsigned int primIdx,
517  unsigned int sbtGASIndex,
518  float time,
519  float3 data[3] )
520 {
521  asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8), _optix_get_triangle_vertex_data, "
522  "(%9, %10, %11, %12);"
523  : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[1].x ), "=f"( data[1].y ),
524  "=f"( data[1].z ), "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z )
525  : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
526  : );
527 }
528 
529 static __forceinline__ __device__ void optixGetLinearCurveVertexData( OptixTraversableHandle gas,
530  unsigned int primIdx,
531  unsigned int sbtGASIndex,
532  float time,
533  float4 data[2] )
534 {
535  asm( "call (%0, %1, %2, %3, %4, %5, %6, %7), _optix_get_linear_curve_vertex_data, "
536  "(%8, %9, %10, %11);"
537  : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ),
538  "=f"( data[1].x ), "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w )
539  : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
540  : );
541 }
542 
543 static __forceinline__ __device__ void optixGetQuadraticBSplineVertexData( OptixTraversableHandle gas,
544  unsigned int primIdx,
545  unsigned int sbtGASIndex,
546  float time,
547  float4 data[3] )
548 {
549  asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11), _optix_get_quadratic_bspline_vertex_data, "
550  "(%12, %13, %14, %15);"
551  : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ),
552  "=f"( data[1].x ), "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w ),
553  "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z ), "=f"( data[2].w )
554  : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
555  : );
556 }
557 
558 static __forceinline__ __device__ void optixGetCubicBSplineVertexData( OptixTraversableHandle gas,
559  unsigned int primIdx,
560  unsigned int sbtGASIndex,
561  float time,
562  float4 data[4] )
563 {
564  asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15), "
565  "_optix_get_cubic_bspline_vertex_data, "
566  "(%16, %17, %18, %19);"
567  : "=f"( data[0].x ), "=f"( data[0].y ), "=f"( data[0].z ), "=f"( data[0].w ),
568  "=f"( data[1].x ), "=f"( data[1].y ), "=f"( data[1].z ), "=f"( data[1].w ),
569  "=f"( data[2].x ), "=f"( data[2].y ), "=f"( data[2].z ), "=f"( data[2].w ),
570  "=f"( data[3].x ), "=f"( data[3].y ), "=f"( data[3].z ), "=f"( data[3].w )
571  : "l"( gas ), "r"( primIdx ), "r"( sbtGASIndex ), "f"( time )
572  : );
573 }
574 
575 static __forceinline__ __device__ OptixTraversableHandle optixGetGASTraversableHandle()
576 {
577  unsigned long long handle;
578  asm( "call (%0), _optix_get_gas_traversable_handle, ();" : "=l"( handle ) : );
579  return (OptixTraversableHandle)handle;
580 }
581 
582 static __forceinline__ __device__ float optixGetGASMotionTimeBegin( OptixTraversableHandle handle )
583 {
584  float f0;
585  asm( "call (%0), _optix_get_gas_motion_time_begin, (%1);" : "=f"( f0 ) : "l"( handle ) : );
586  return f0;
587 }
588 
589 static __forceinline__ __device__ float optixGetGASMotionTimeEnd( OptixTraversableHandle handle )
590 {
591  float f0;
592  asm( "call (%0), _optix_get_gas_motion_time_end, (%1);" : "=f"( f0 ) : "l"( handle ) : );
593  return f0;
594 }
595 
596 static __forceinline__ __device__ unsigned int optixGetGASMotionStepCount( OptixTraversableHandle handle )
597 {
598  unsigned int u0;
599  asm( "call (%0), _optix_get_gas_motion_step_count, (%1);" : "=r"( u0 ) : "l"( handle ) : );
600  return u0;
601 }
602 
603 static __forceinline__ __device__ void optixGetWorldToObjectTransformMatrix( float m[12] )
604 {
605  if( optixGetTransformListSize() == 0 )
606  {
607  m[0] = 1.0f;
608  m[1] = 0.0f;
609  m[2] = 0.0f;
610  m[3] = 0.0f;
611  m[4] = 0.0f;
612  m[5] = 1.0f;
613  m[6] = 0.0f;
614  m[7] = 0.0f;
615  m[8] = 0.0f;
616  m[9] = 0.0f;
617  m[10] = 1.0f;
618  m[11] = 0.0f;
619  return;
620  }
621 
622  float4 m0, m1, m2;
624  m[0] = m0.x;
625  m[1] = m0.y;
626  m[2] = m0.z;
627  m[3] = m0.w;
628  m[4] = m1.x;
629  m[5] = m1.y;
630  m[6] = m1.z;
631  m[7] = m1.w;
632  m[8] = m2.x;
633  m[9] = m2.y;
634  m[10] = m2.z;
635  m[11] = m2.w;
636 }
637 
638 static __forceinline__ __device__ void optixGetObjectToWorldTransformMatrix( float m[12] )
639 {
640  if( optixGetTransformListSize() == 0 )
641  {
642  m[0] = 1.0f;
643  m[1] = 0.0f;
644  m[2] = 0.0f;
645  m[3] = 0.0f;
646  m[4] = 0.0f;
647  m[5] = 1.0f;
648  m[6] = 0.0f;
649  m[7] = 0.0f;
650  m[8] = 0.0f;
651  m[9] = 0.0f;
652  m[10] = 1.0f;
653  m[11] = 0.0f;
654  return;
655  }
656 
657  float4 m0, m1, m2;
659  m[0] = m0.x;
660  m[1] = m0.y;
661  m[2] = m0.z;
662  m[3] = m0.w;
663  m[4] = m1.x;
664  m[5] = m1.y;
665  m[6] = m1.z;
666  m[7] = m1.w;
667  m[8] = m2.x;
668  m[9] = m2.y;
669  m[10] = m2.z;
670  m[11] = m2.w;
671 }
672 
673 static __forceinline__ __device__ float3 optixTransformPointFromWorldToObjectSpace( float3 point )
674 {
675  if( optixGetTransformListSize() == 0 )
676  return point;
677 
678  float4 m0, m1, m2;
680  return optix_impl::optixTransformPoint( m0, m1, m2, point );
681 }
682 
683 static __forceinline__ __device__ float3 optixTransformVectorFromWorldToObjectSpace( float3 vec )
684 {
685  if( optixGetTransformListSize() == 0 )
686  return vec;
687 
688  float4 m0, m1, m2;
690  return optix_impl::optixTransformVector( m0, m1, m2, vec );
691 }
692 
693 static __forceinline__ __device__ float3 optixTransformNormalFromWorldToObjectSpace( float3 normal )
694 {
695  if( optixGetTransformListSize() == 0 )
696  return normal;
697 
698  float4 m0, m1, m2;
699  optix_impl::optixGetObjectToWorldTransformMatrix( m0, m1, m2 ); // inverse of optixGetWorldToObjectTransformMatrix()
700  return optix_impl::optixTransformNormal( m0, m1, m2, normal );
701 }
702 
703 static __forceinline__ __device__ float3 optixTransformPointFromObjectToWorldSpace( float3 point )
704 {
705  if( optixGetTransformListSize() == 0 )
706  return point;
707 
708  float4 m0, m1, m2;
710  return optix_impl::optixTransformPoint( m0, m1, m2, point );
711 }
712 
713 static __forceinline__ __device__ float3 optixTransformVectorFromObjectToWorldSpace( float3 vec )
714 {
715  if( optixGetTransformListSize() == 0 )
716  return vec;
717 
718  float4 m0, m1, m2;
720  return optix_impl::optixTransformVector( m0, m1, m2, vec );
721 }
722 
723 static __forceinline__ __device__ float3 optixTransformNormalFromObjectToWorldSpace( float3 normal )
724 {
725  if( optixGetTransformListSize() == 0 )
726  return normal;
727 
728  float4 m0, m1, m2;
729  optix_impl::optixGetWorldToObjectTransformMatrix( m0, m1, m2 ); // inverse of optixGetObjectToWorldTransformMatrix()
730  return optix_impl::optixTransformNormal( m0, m1, m2, normal );
731 }
732 
733 static __forceinline__ __device__ unsigned int optixGetTransformListSize()
734 {
735  unsigned int u0;
736  asm( "call (%0), _optix_get_transform_list_size, ();" : "=r"( u0 ) : );
737  return u0;
738 }
739 
740 static __forceinline__ __device__ OptixTraversableHandle optixGetTransformListHandle( unsigned int index )
741 {
742  unsigned long long u0;
743  asm( "call (%0), _optix_get_transform_list_handle, (%1);" : "=l"( u0 ) : "r"( index ) : );
744  return u0;
745 }
746 
748 {
749  int i0;
750  asm( "call (%0), _optix_get_transform_type_from_handle, (%1);" : "=r"( i0 ) : "l"( handle ) : );
751  return (OptixTransformType)i0;
752 }
753 
754 static __forceinline__ __device__ const OptixStaticTransform* optixGetStaticTransformFromHandle( OptixTraversableHandle handle )
755 {
756  unsigned long long ptr;
757  asm( "call (%0), _optix_get_static_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
758  return (const OptixStaticTransform*)ptr;
759 }
760 
762 {
763  unsigned long long ptr;
764  asm( "call (%0), _optix_get_srt_motion_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
765  return (const OptixSRTMotionTransform*)ptr;
766 }
767 
769 {
770  unsigned long long ptr;
771  asm( "call (%0), _optix_get_matrix_motion_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
772  return (const OptixMatrixMotionTransform*)ptr;
773 }
774 
775 static __forceinline__ __device__ unsigned int optixGetInstanceIdFromHandle( OptixTraversableHandle handle )
776 {
777  int i0;
778  asm( "call (%0), _optix_get_instance_id_from_handle, (%1);" : "=r"( i0 ) : "l"( handle ) : );
779  return i0;
780 }
781 
782 static __forceinline__ __device__ const float4* optixGetInstanceTransformFromHandle( OptixTraversableHandle handle )
783 {
784  unsigned long long ptr;
785  asm( "call (%0), _optix_get_instance_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
786  return (const float4*)ptr;
787 }
788 
789 static __forceinline__ __device__ const float4* optixGetInstanceInverseTransformFromHandle( OptixTraversableHandle handle )
790 {
791  unsigned long long ptr;
792  asm( "call (%0), _optix_get_instance_inverse_transform_from_handle, (%1);" : "=l"( ptr ) : "l"( handle ) : );
793  return (const float4*)ptr;
794 }
795 
796 static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind )
797 {
798  int ret;
799  asm volatile(
800  "call (%0), _optix_report_intersection_0"
801  ", (%1, %2);"
802  : "=r"( ret )
803  : "f"( hitT ), "r"( hitKind )
804  : );
805  return ret;
806 }
807 
808 static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind, unsigned int a0 )
809 {
810  int ret;
811  asm volatile(
812  "call (%0), _optix_report_intersection_1"
813  ", (%1, %2, %3);"
814  : "=r"( ret )
815  : "f"( hitT ), "r"( hitKind ), "r"( a0 )
816  : );
817  return ret;
818 }
819 
820 static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind, unsigned int a0, unsigned int a1 )
821 {
822  int ret;
823  asm volatile(
824  "call (%0), _optix_report_intersection_2"
825  ", (%1, %2, %3, %4);"
826  : "=r"( ret )
827  : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 )
828  : );
829  return ret;
830 }
831 
832 static __forceinline__ __device__ bool optixReportIntersection( float hitT, unsigned int hitKind, unsigned int a0, unsigned int a1, unsigned int a2 )
833 {
834  int ret;
835  asm volatile(
836  "call (%0), _optix_report_intersection_3"
837  ", (%1, %2, %3, %4, %5);"
838  : "=r"( ret )
839  : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 )
840  : );
841  return ret;
842 }
843 
844 static __forceinline__ __device__ bool optixReportIntersection( float hitT,
845  unsigned int hitKind,
846  unsigned int a0,
847  unsigned int a1,
848  unsigned int a2,
849  unsigned int a3 )
850 {
851  int ret;
852  asm volatile(
853  "call (%0), _optix_report_intersection_4"
854  ", (%1, %2, %3, %4, %5, %6);"
855  : "=r"( ret )
856  : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 )
857  : );
858  return ret;
859 }
860 
861 static __forceinline__ __device__ bool optixReportIntersection( float hitT,
862  unsigned int hitKind,
863  unsigned int a0,
864  unsigned int a1,
865  unsigned int a2,
866  unsigned int a3,
867  unsigned int a4 )
868 {
869  int ret;
870  asm volatile(
871  "call (%0), _optix_report_intersection_5"
872  ", (%1, %2, %3, %4, %5, %6, %7);"
873  : "=r"( ret )
874  : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 )
875  : );
876  return ret;
877 }
878 
879 static __forceinline__ __device__ bool optixReportIntersection( float hitT,
880  unsigned int hitKind,
881  unsigned int a0,
882  unsigned int a1,
883  unsigned int a2,
884  unsigned int a3,
885  unsigned int a4,
886  unsigned int a5 )
887 {
888  int ret;
889  asm volatile(
890  "call (%0), _optix_report_intersection_6"
891  ", (%1, %2, %3, %4, %5, %6, %7, %8);"
892  : "=r"( ret )
893  : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 ), "r"( a5 )
894  : );
895  return ret;
896 }
897 
898 static __forceinline__ __device__ bool optixReportIntersection( float hitT,
899  unsigned int hitKind,
900  unsigned int a0,
901  unsigned int a1,
902  unsigned int a2,
903  unsigned int a3,
904  unsigned int a4,
905  unsigned int a5,
906  unsigned int a6 )
907 {
908  int ret;
909  asm volatile(
910  "call (%0), _optix_report_intersection_7"
911  ", (%1, %2, %3, %4, %5, %6, %7, %8, %9);"
912  : "=r"( ret )
913  : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 ), "r"( a5 ), "r"( a6 )
914  : );
915  return ret;
916 }
917 
918 static __forceinline__ __device__ bool optixReportIntersection( float hitT,
919  unsigned int hitKind,
920  unsigned int a0,
921  unsigned int a1,
922  unsigned int a2,
923  unsigned int a3,
924  unsigned int a4,
925  unsigned int a5,
926  unsigned int a6,
927  unsigned int a7 )
928 {
929  int ret;
930  asm volatile(
931  "call (%0), _optix_report_intersection_8"
932  ", (%1, %2, %3, %4, %5, %6, %7, %8, %9, %10);"
933  : "=r"( ret )
934  : "f"( hitT ), "r"( hitKind ), "r"( a0 ), "r"( a1 ), "r"( a2 ), "r"( a3 ), "r"( a4 ), "r"( a5 ), "r"( a6 ), "r"( a7 )
935  : );
936  return ret;
937 }
938 
939 #define OPTIX_DEFINE_optixGetAttribute_BODY( which ) \
940  unsigned int ret; \
941  asm( "call (%0), _optix_get_attribute_" #which ", ();" : "=r"( ret ) : ); \
942  return ret;
943 
944 static __forceinline__ __device__ unsigned int optixGetAttribute_0()
945 {
947 }
948 
949 static __forceinline__ __device__ unsigned int optixGetAttribute_1()
950 {
952 }
953 
954 static __forceinline__ __device__ unsigned int optixGetAttribute_2()
955 {
957 }
958 
959 static __forceinline__ __device__ unsigned int optixGetAttribute_3()
960 {
962 }
963 
964 static __forceinline__ __device__ unsigned int optixGetAttribute_4()
965 {
967 }
968 
969 static __forceinline__ __device__ unsigned int optixGetAttribute_5()
970 {
972 }
973 
974 static __forceinline__ __device__ unsigned int optixGetAttribute_6()
975 {
977 }
978 
979 static __forceinline__ __device__ unsigned int optixGetAttribute_7()
980 {
982 }
983 
984 #undef OPTIX_DEFINE_optixGetAttribute_BODY
985 
986 static __forceinline__ __device__ void optixTerminateRay()
987 {
988  asm volatile( "call _optix_terminate_ray, ();" );
989 }
990 
991 static __forceinline__ __device__ void optixIgnoreIntersection()
992 {
993  asm volatile( "call _optix_ignore_intersection, ();" );
994 }
995 
996 static __forceinline__ __device__ unsigned int optixGetPrimitiveIndex()
997 {
998  unsigned int u0;
999  asm( "call (%0), _optix_read_primitive_idx, ();" : "=r"( u0 ) : );
1000  return u0;
1001 }
1002 
1003 static __forceinline__ __device__ unsigned int optixGetSbtGASIndex()
1004 {
1005  unsigned int u0;
1006  asm( "call (%0), _optix_read_sbt_gas_idx, ();" : "=r"( u0 ) : );
1007  return u0;
1008 }
1009 
1010 static __forceinline__ __device__ unsigned int optixGetInstanceId()
1011 {
1012  unsigned int u0;
1013  asm( "call (%0), _optix_read_instance_id, ();" : "=r"( u0 ) : );
1014  return u0;
1015 }
1016 
1017 static __forceinline__ __device__ unsigned int optixGetInstanceIndex()
1018 {
1019  unsigned int u0;
1020  asm( "call (%0), _optix_read_instance_idx, ();" : "=r"( u0 ) : );
1021  return u0;
1022 }
1023 
1024 static __forceinline__ __device__ unsigned int optixGetHitKind()
1025 {
1026  unsigned int u0;
1027  asm( "call (%0), _optix_get_hit_kind, ();" : "=r"( u0 ) : );
1028  return u0;
1029 }
1030 
1031 static __forceinline__ __device__ OptixPrimitiveType optixGetPrimitiveType(unsigned int hitKind)
1032 {
1033  unsigned int u0;
1034  asm( "call (%0), _optix_get_primitive_type_from_hit_kind, (%1);" : "=r"( u0 ) : "r"( hitKind ) );
1035  return (OptixPrimitiveType)u0;
1036 }
1037 
1038 static __forceinline__ __device__ bool optixIsBackFaceHit( unsigned int hitKind )
1039 {
1040  unsigned int u0;
1041  asm( "call (%0), _optix_get_backface_from_hit_kind, (%1);" : "=r"( u0 ) : "r"( hitKind ) );
1042  return (u0 == 0x1);
1043 }
1044 
1045 static __forceinline__ __device__ bool optixIsFrontFaceHit( unsigned int hitKind )
1046 {
1047  return !optixIsBackFaceHit( hitKind );
1048 }
1049 
1050 
1051 static __forceinline__ __device__ OptixPrimitiveType optixGetPrimitiveType()
1052 {
1054 }
1055 
1056 static __forceinline__ __device__ bool optixIsBackFaceHit()
1057 {
1058  return optixIsBackFaceHit( optixGetHitKind() );
1059 }
1060 
1061 static __forceinline__ __device__ bool optixIsFrontFaceHit()
1062 {
1064 }
1065 
1066 static __forceinline__ __device__ bool optixIsTriangleHit()
1067 {
1069 }
1070 
1071 static __forceinline__ __device__ bool optixIsTriangleFrontFaceHit()
1072 {
1074 }
1075 
1076 static __forceinline__ __device__ bool optixIsTriangleBackFaceHit()
1077 {
1079 }
1080 
1081 static __forceinline__ __device__ float optixGetCurveParameter()
1082 {
1083  return __int_as_float( optixGetAttribute_0() );
1084 }
1085 
1086 static __forceinline__ __device__ float2 optixGetTriangleBarycentrics()
1087 {
1088  float f0, f1;
1089  asm( "call (%0, %1), _optix_get_triangle_barycentrics, ();" : "=f"( f0 ), "=f"( f1 ) : );
1090  return make_float2( f0, f1 );
1091 }
1092 
1093 static __forceinline__ __device__ uint3 optixGetLaunchIndex()
1094 {
1095  unsigned int u0, u1, u2;
1096  asm( "call (%0), _optix_get_launch_index_x, ();" : "=r"( u0 ) : );
1097  asm( "call (%0), _optix_get_launch_index_y, ();" : "=r"( u1 ) : );
1098  asm( "call (%0), _optix_get_launch_index_z, ();" : "=r"( u2 ) : );
1099  return make_uint3( u0, u1, u2 );
1100 }
1101 
1102 static __forceinline__ __device__ uint3 optixGetLaunchDimensions()
1103 {
1104  unsigned int u0, u1, u2;
1105  asm( "call (%0), _optix_get_launch_dimension_x, ();" : "=r"( u0 ) : );
1106  asm( "call (%0), _optix_get_launch_dimension_y, ();" : "=r"( u1 ) : );
1107  asm( "call (%0), _optix_get_launch_dimension_z, ();" : "=r"( u2 ) : );
1108  return make_uint3( u0, u1, u2 );
1109 }
1110 
1111 static __forceinline__ __device__ CUdeviceptr optixGetSbtDataPointer()
1112 {
1113  unsigned long long ptr;
1114  asm( "call (%0), _optix_get_sbt_data_ptr_64, ();" : "=l"( ptr ) : );
1115  return (CUdeviceptr)ptr;
1116 }
1117 
1118 static __forceinline__ __device__ void optixThrowException( int exceptionCode )
1119 {
1120  asm volatile(
1121  "call _optix_throw_exception_0, (%0);"
1122  : /* no return value */
1123  : "r"( exceptionCode )
1124  : );
1125 }
1126 
1127 static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0 )
1128 {
1129  asm volatile(
1130  "call _optix_throw_exception_1, (%0, %1);"
1131  : /* no return value */
1132  : "r"( exceptionCode ), "r"( exceptionDetail0 )
1133  : );
1134 }
1135 
1136 static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1 )
1137 {
1138  asm volatile(
1139  "call _optix_throw_exception_2, (%0, %1, %2);"
1140  : /* no return value */
1141  : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 )
1142  : );
1143 }
1144 
1145 static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2 )
1146 {
1147  asm volatile(
1148  "call _optix_throw_exception_3, (%0, %1, %2, %3);"
1149  : /* no return value */
1150  : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 )
1151  : );
1152 }
1153 
1154 static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3 )
1155 {
1156  asm volatile(
1157  "call _optix_throw_exception_4, (%0, %1, %2, %3, %4);"
1158  : /* no return value */
1159  : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 )
1160  : );
1161 }
1162 
1163 static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3, unsigned int exceptionDetail4 )
1164 {
1165  asm volatile(
1166  "call _optix_throw_exception_5, (%0, %1, %2, %3, %4, %5);"
1167  : /* no return value */
1168  : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 )
1169  : );
1170 }
1171 
1172 static __forceinline__ __device__ void optixThrowException( int exceptionCode, unsigned int exceptionDetail0, unsigned int exceptionDetail1, unsigned int exceptionDetail2, unsigned int exceptionDetail3, unsigned int exceptionDetail4, unsigned int exceptionDetail5 )
1173 {
1174  asm volatile(
1175  "call _optix_throw_exception_6, (%0, %1, %2, %3, %4, %5, %6);"
1176  : /* no return value */
1177  : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 ), "r"( exceptionDetail5 )
1178  : );
1179 }
1180 
1181 static __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 )
1182 {
1183  asm volatile(
1184  "call _optix_throw_exception_7, (%0, %1, %2, %3, %4, %5, %6, %7);"
1185  : /* no return value */
1186  : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 ), "r"( exceptionDetail5 ), "r"( exceptionDetail6 )
1187  : );
1188 }
1189 
1190 static __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 )
1191 {
1192  asm volatile(
1193  "call _optix_throw_exception_8, (%0, %1, %2, %3, %4, %5, %6, %7, %8);"
1194  : /* no return value */
1195  : "r"( exceptionCode ), "r"( exceptionDetail0 ), "r"( exceptionDetail1 ), "r"( exceptionDetail2 ), "r"( exceptionDetail3 ), "r"( exceptionDetail4 ), "r"( exceptionDetail5 ), "r"( exceptionDetail6 ), "r"( exceptionDetail7 )
1196  : );
1197 }
1198 
1199 static __forceinline__ __device__ int optixGetExceptionCode()
1200 {
1201  int s0;
1202  asm( "call (%0), _optix_get_exception_code, ();" : "=r"( s0 ) : );
1203  return s0;
1204 }
1205 
1206 #define OPTIX_DEFINE_optixGetExceptionDetail_BODY( which ) \
1207  unsigned int ret; \
1208  asm( "call (%0), _optix_get_exception_detail_" #which ", ();" : "=r"( ret ) : ); \
1209  return ret;
1210 
1211 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_0()
1212 {
1214 }
1215 
1216 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_1()
1217 {
1219 }
1220 
1221 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_2()
1222 {
1224 }
1225 
1226 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_3()
1227 {
1229 }
1230 
1231 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_4()
1232 {
1234 }
1235 
1236 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_5()
1237 {
1239 }
1240 
1241 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_6()
1242 {
1244 }
1245 
1246 static __forceinline__ __device__ unsigned int optixGetExceptionDetail_7()
1247 {
1249 }
1250 
1251 #undef OPTIX_DEFINE_optixGetExceptionDetail_BODY
1252 
1254 {
1255  unsigned long long handle;
1256  asm( "call (%0), _optix_get_exception_invalid_traversable, ();" : "=l"( handle ) : );
1257  return (OptixTraversableHandle)handle;
1258 }
1259 
1260 static __forceinline__ __device__ int optixGetExceptionInvalidSbtOffset()
1261 {
1262  int s0;
1263  asm( "call (%0), _optix_get_exception_invalid_sbt_offset, ();" : "=r"( s0 ) : );
1264  return s0;
1265 }
1266 
1267 static __forceinline__ __device__ OptixInvalidRayExceptionDetails optixGetExceptionInvalidRay()
1268 {
1269  float rayOriginX, rayOriginY, rayOriginZ, rayDirectionX, rayDirectionY, rayDirectionZ, tmin, tmax, rayTime;
1270  asm( "call (%0, %1, %2, %3, %4, %5, %6, %7, %8), _optix_get_exception_invalid_ray, ();"
1271  : "=f"( rayOriginX ), "=f"( rayOriginY ), "=f"( rayOriginZ ), "=f"( rayDirectionX ), "=f"( rayDirectionY ),
1272  "=f"( rayDirectionZ ), "=f"( tmin ), "=f"( tmax ), "=f"( rayTime )
1273  : );
1274  OptixInvalidRayExceptionDetails ray;
1275  ray.origin = make_float3( rayOriginX, rayOriginY, rayOriginZ );
1276  ray.direction = make_float3( rayDirectionX, rayDirectionY, rayDirectionZ );
1277  ray.tmin = tmin;
1278  ray.tmax = tmax;
1279  ray.time = rayTime;
1280  return ray;
1281 }
1282 
1283 static __forceinline__ __device__ OptixParameterMismatchExceptionDetails optixGetExceptionParameterMismatch()
1284 {
1285  unsigned int expected, actual, sbtIdx;
1286  unsigned long long calleeName;
1287  asm(
1288  "call (%0, %1, %2, %3), _optix_get_exception_parameter_mismatch, ();"
1289  : "=r"(expected), "=r"(actual), "=r"(sbtIdx), "=l"(calleeName) : );
1290  OptixParameterMismatchExceptionDetails details;
1291  details.expectedParameterCount = expected;
1292  details.passedArgumentCount = actual;
1293  details.sbtIndex = sbtIdx;
1294  details.callableName = (char*)calleeName;
1295  return details;
1296 }
1297 
1298 static __forceinline__ __device__ char* optixGetExceptionLineInfo()
1299 {
1300  unsigned long long ptr;
1301  asm( "call (%0), _optix_get_exception_line_info, ();" : "=l"(ptr) : );
1302  return (char*)ptr;
1303 }
1304 
1305 template <typename ReturnT, typename... ArgTypes>
1306 static __forceinline__ __device__ ReturnT optixDirectCall( unsigned int sbtIndex, ArgTypes... args )
1307 {
1308  unsigned long long func;
1309  asm( "call (%0), _optix_call_direct_callable,(%1);" : "=l"( func ) : "r"( sbtIndex ) : );
1310  using funcT = ReturnT ( * )( ArgTypes... );
1311  funcT call = ( funcT )( func );
1312  return call( args... );
1313 }
1314 
1315 template <typename ReturnT, typename... ArgTypes>
1316 static __forceinline__ __device__ ReturnT optixContinuationCall( unsigned int sbtIndex, ArgTypes... args )
1317 {
1318  unsigned long long func;
1319  asm( "call (%0), _optix_call_continuation_callable,(%1);" : "=l"( func ) : "r"( sbtIndex ) : );
1320  using funcT = ReturnT ( * )( ArgTypes... );
1321  funcT call = ( funcT )( func );
1322  return call( args... );
1323 }
1324 #endif
1325