OptiX  3.9
NVIDIA OptiX Acceleration Engine
 All Classes Files Functions Variables Typedefs Enumerations Enumerator Macros Groups
optix_device.h
Go to the documentation of this file.
1 
2 /*
3  * Copyright (c) 2008 - 2009 NVIDIA Corporation. All rights reserved.
4  *
5  * NVIDIA Corporation and its licensors retain all intellectual property and proprietary
6  * rights in and to this software, related documentation and any modifications thereto.
7  * Any use, reproduction, disclosure or distribution of this software and related
8  * documentation without an express license agreement from NVIDIA Corporation is strictly
9  * prohibited.
10  *
11  * TO THE MAXIMUM EXTENT PERMITTED BY APPLICABLE LAW, THIS SOFTWARE IS PROVIDED *AS IS*
12  * AND NVIDIA AND ITS SUPPLIERS DISCLAIM ALL WARRANTIES, EITHER EXPRESS OR IMPLIED,
13  * INCLUDING, BUT NOT LIMITED TO, IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
14  * PARTICULAR PURPOSE. IN NO EVENT SHALL NVIDIA OR ITS SUPPLIERS BE LIABLE FOR ANY
15  * SPECIAL, INCIDENTAL, INDIRECT, OR CONSEQUENTIAL DAMAGES WHATSOEVER (INCLUDING, WITHOUT
16  * LIMITATION, DAMAGES FOR LOSS OF BUSINESS PROFITS, BUSINESS INTERRUPTION, LOSS OF
17  * BUSINESS INFORMATION, OR ANY OTHER PECUNIARY LOSS) ARISING OUT OF THE USE OF OR
18  * INABILITY TO USE THIS SOFTWARE, EVEN IF NVIDIA HAS BEEN ADVISED OF THE POSSIBILITY OF
19  * SUCH DAMAGES
20  */
21 
30 /******************************************************************************\
31  * optix_cuda.h
32  *
33  * This file provides the nvcc interface for generating PTX that the OptiX is
34  * capable of parsing and weaving into the final kernel. This is included by
35  * optix.h automatically if compiling device code. It can be included explicitly
36  * in host code if desired.
37  *
38 \******************************************************************************/
39 
40 #ifndef __optix_optix_cuda__internal_h__
41 #define __optix_optix_cuda__internal_h__
42 
43 #include "internal/optix_datatypes.h"
44 #include "internal/optix_declarations.h"
45 #include "internal/optix_internal.h"
46 
47 /*
48  Augment vector types
49 */
50 
51 namespace optix {
52 
53  template<typename T, int Dim> struct VectorTypes {};
54  template<> struct VectorTypes<int, 1> {
55  typedef int Type;
56  template<class S> static __device__ __forceinline__
57  Type make(S s) { return make_int(s); }
58  };
59  template<> struct VectorTypes<int, 2> {
60  typedef int2 Type;
61  template<class S> static __device__ __forceinline__
62  Type make(S s) { return make_int2(s); }
63  };
64  template<> struct VectorTypes<int, 3> {
65  typedef int3 Type;
66  template<class S> static __device__ __forceinline__
67  Type make(S s) { return make_int3(s); }
68  };
69  template<> struct VectorTypes<int, 4> {
70  typedef int4 Type;
71  template<class S> static __device__ __forceinline__
72  Type make(S s) { return make_int4(s); }
73  };
74  template<> struct VectorTypes<unsigned int, 1> {
75  typedef unsigned int Type;
76  static __device__ __forceinline__
77  Type make(unsigned int s) { return s; }
78  template<class S> static __device__ __forceinline__
79  Type make(S s) { return (unsigned int)s.x; }
80  };
81  template<> struct VectorTypes<unsigned int, 2> {
82  typedef uint2 Type;
83  template<class S> static __device__ __forceinline__
84  Type make(S s) { return make_uint2(s); }
85  };
86  template<> struct VectorTypes<unsigned int, 3> {
87  typedef uint3 Type;
88  template<class S> static __device__ __forceinline__
89  Type make(S s) { return make_uint3(s); }
90  };
91  template<> struct VectorTypes<unsigned int, 4> {
92  typedef uint4 Type;
93  template<class S> static __device__ __forceinline__
94  Type make(S s) { return make_uint4(s); }
95  };
96  template<> struct VectorTypes<float, 1> {
97  typedef float Type;
98  template<class S> static __device__ __forceinline__
99  Type make(S s) { return make_float(s); }
100  };
101  template<> struct VectorTypes<float, 2> {
102  typedef float2 Type;
103  template<class S> static __device__ __forceinline__
104  Type make(S s) { return make_float2(s); }
105  };
106  template<> struct VectorTypes<float, 3> {
107  typedef float3 Type;
108  template<class S> static __device__ __forceinline__
109  Type make(S s) { return make_float3(s); }
110  };
111  template<> struct VectorTypes<float, 4> {
112  typedef float4 Type;
113  template<class S> static __device__ __forceinline__
114  Type make(S s) { return make_float4(s); }
115  };
116 
117 #if defined(__APPLE__) || defined(__x86_64) || defined(AMD64) || defined(_M_AMD64)
118  template<> struct VectorTypes<size_t, 1> {
119  typedef size_t Type;
120  static __device__ __forceinline__
121  Type make(unsigned int s) { return s; }
122  template<class S> static __device__ __forceinline__
123  Type make(S s) { return (unsigned int)s.x; }
124  };
125  template<> struct VectorTypes<size_t, 2> {
126  typedef size_t2 Type;
127  template<class S> static __device__ __forceinline__
128  Type make(S s) { return make_size_t2(s); }
129  };
130  template<> struct VectorTypes<size_t, 3> {
131  typedef size_t3 Type;
132  template<class S> static __device__ __forceinline__
133  Type make(S s) { return make_size_t3(s); }
134  };
135  template<> struct VectorTypes<size_t, 4> {
136  typedef size_t4 Type;
137  template<class S> static __device__ __forceinline__
138  Type make(S s) { return make_size_t4(s); }
139  };
140 #endif
141 }
142 
143 /*
144  Variables
145 */
146 
179 struct rtObject {
180 protected:
181  unsigned int handle;
182  /* Bogus use of handle to quiet warnings from compilers that warn about unused private
183  * data members. */
184  void never_call() { handle = 0; }
185 };
186 
241 #define rtDeclareVariable(type, name, semantic, annotation) \
242  namespace rti_internal_typeinfo { \
243  __device__ ::rti_internal_typeinfo::rti_typeinfo name = { ::rti_internal_typeinfo::_OPTIX_VARIABLE, sizeof(type)}; \
244  } \
245  namespace rti_internal_typename { \
246  __device__ char name[] = #type; \
247  } \
248  namespace rti_internal_typeenum { \
249  __device__ int name = ::rti_internal_typeinfo::rti_typeenum<type>::m_typeenum; \
250  } \
251  namespace rti_internal_semantic { \
252  __device__ char name[] = #semantic; \
253  } \
254  namespace rti_internal_annotation { \
255  __device__ char name[] = #annotation; \
256  } \
257  __device__ type name
258 
259 
319 #define rtDeclareAnnotation(variable, annotation) \
320  namespace rti_internal_annotation { \
321  __device__ char variable[] = #annotation; \
322  }
323 
324 /* Declares a function that can be set via rtVariableSetObject and called from CUDA C.
325  Once declared the function variable can be used as if it were a regular function. Note
326  that the parameters argument to the macro need to have parentheses even if they will be
327  empty.
328 
329  Example: rtCallableProgram(float, times2, (float));
330  Example: rtCallableProgram(float, doStuff, ());
331  */
332 
333 template<typename T> struct rtCallableProgramSizeofWrapper { static const size_t value = sizeof(T); };
334 template<> struct rtCallableProgramSizeofWrapper<void> { static const size_t value = 0; };
335 
336 
374 #ifndef RT_USE_TEMPLATED_RTCALLABLEPROGRAM
375 #if (__CUDA_ARCH__ < 200) || (CUDART_VERSION < 4010)
376 #define rtCallableProgram(return_type, function_name, parameter_list) \
377  namespace rti_internal_typeinfo { \
378  __device__ ::rti_internal_typeinfo::rti_typeinfo function_name = { ::rti_internal_typeinfo::_OPTIX_VARIABLE, rtCallableProgramSizeofWrapper<return_type>::value }; \
379  } \
380  namespace rti_internal_typename { \
381  __device__ char function_name[] = #return_type; \
382  } \
383  namespace rti_internal_semantic { \
384  __device__ char function_name[] = ""; /* used to be rt_call, but not needed anymore */ \
385  } \
386  namespace rti_internal_annotation { \
387  __device__ char function_name[] = #parameter_list; \
388  } \
389  __noinline__ __device__ return_type function_name parameter_list { typedef return_type localtype; return localtype(); }
390 #else
391 #define rtCallableProgram(return_type, function_name, parameter_list) \
392  rtDeclareVariable(optix::boundCallableProgramId<return_type parameter_list>, function_name,,);
393 #endif
394 #endif
395 
396 /* Helper functions for converting pointers to/from integers so we can support
397  call by reference on SM1.x */
398 namespace optix {
399 
400  typedef unsigned int rtPickledLocalPointer;
401  static inline __device__ rtPickledLocalPointer rtPickleLocalPointer( void *p ) {
402  return optix::rt_pickle_pointer( p );
403  }
404 
405  static inline __device__ void * rtUnpickleLocalPointer( rtPickledLocalPointer p ) {
406  return optix::rt_unpickle_pointer( p );
407  }
408 }
409 
410 /*
411  Buffer
412 */
413 
414 namespace optix {
415  template<typename T, int Dim> struct bufferId;
416 
417  template<typename T, int Dim = 1> struct buffer {
418  typedef VectorTypes<size_t, Dim> WrapperType;
419  typedef typename VectorTypes<size_t, Dim>::Type IndexType;
420 
421  __device__ __forceinline__ IndexType size() const {
422  return WrapperType::make(rt_buffer_get_size(this, Dim, sizeof(T)));
423  }
424  __device__ __forceinline__ T& operator[](IndexType i) {
425  size_t4 c = make_index(i);
426  return *(T*)create(type<T>(), rt_buffer_get(this, Dim, sizeof(T), c.x, c.y, c.z, c.w));
427  }
428  protected:
429  __inline__ __device__ static size_t4 make_index(size_t v0) { return make_size_t4(v0, 0, 0, 0); }
430  __inline__ __device__ static size_t4 make_index(size_t2 v0) { return make_size_t4(v0.x, v0.y, 0, 0); }
431  __inline__ __device__ static size_t4 make_index(size_t3 v0) { return make_size_t4(v0.x, v0.y, v0.z, 0); }
432  __inline__ __device__ static size_t4 make_index(size_t4 v0) { return make_size_t4(v0.x, v0.y, v0.z, v0.w); }
433 
434  // This struct is used to create overloaded methods based on the type of buffer
435  // element. Note that we use a different name for the typemplate typename to avoid
436  // confusing it with the template type of buffer.
437  template<typename T2> struct type { };
438 
439  // Regular type: just return the pointer
440  template<typename T2> __device__ __forceinline__ static void* create(type<T2>, void* v) { return v; }
441  // bufferId type. Read the ID from the buffer than assign it to a new bufferId to be
442  // used later.
443  template<typename T2, int Dim2>
444  __device__ __forceinline__ static void* create(type<bufferId<T2,Dim2> >, void* v)
445  {
446  // Returning a pointer to a locally created thing is generally a bad idea,
447  // however since this function and its caller are always inlined the
448  // object is created on the same stack that the buffer::operator[] was
449  // called from.
450  bufferId<T,Dim> b(*reinterpret_cast<int*>(v));
451  void* result = &b;
452  return result;
453  }
454  };
455 
456  // Helper class for encapsulating a buffer ID with methods to allow it to behave as a buffer.
457  template<typename T, int Dim = 1> struct bufferId : public buffer<T,Dim> {
458  typedef typename buffer<T,Dim>::WrapperType WrapperType;
459  typedef typename buffer<T,Dim>::IndexType IndexType;
460 
461  // Default constructor
462  __device__ __forceinline__ bufferId() {}
463  // Constructor that initializes the id with null.
464  __device__ __forceinline__ bufferId(RTbufferidnull nullid) { m_id = (int)nullid; }
465  // Constructor that initializes the id.
466  __device__ __forceinline__ explicit bufferId(int id) : m_id(id) {}
467 
468  // assigment that initializes the id with null.
469  __device__ __forceinline__ bufferId& operator= (RTbufferidnull nullid) { m_id = nullid; return *this; }
470 
471  // Buffer access methods that use m_id as the argument to identify which buffer is
472  // being accessed.
473  __device__ __forceinline__ IndexType size() const {
474  return WrapperType::make(rt_buffer_get_size_id(m_id, Dim, sizeof(T)));
475  }
476  __device__ __forceinline__ T& operator[](IndexType i) const {
477  size_t4 c = make_index(i);
478  return *(T*)create(typename buffer<T,Dim>::template type<T>(),
479  rt_buffer_get_id(m_id, Dim, sizeof(T), c.x, c.y, c.z, c.w));
480  }
481 
482  __device__ __forceinline__ int getId() const { return m_id; }
483 
484  __device__ __forceinline__ operator bool() const { return m_id; }
485 
486  private:
487  // Member variable
488  int m_id;
489  };
490 }
491 
537 #define rtBuffer __device__ optix::buffer
538 
565 #define rtBufferId optix::bufferId
566 
567 /*
568  Texture - they are defined in CUDA
569 */
570 
600 #define rtTextureSampler texture
601 
602 namespace optix {
603 
604  typedef int rtTextureId;
605 
606  #define _OPTIX_TEX_FUNC_DECLARE_(FUNC, SIGNATURE, PARAMS ) \
607  template<> inline __device__ unsigned char FUNC SIGNATURE \
608  { \
609  uint4 tmp = FUNC <uint4> PARAMS; \
610  return (unsigned char)(tmp.x); \
611  } \
612  template<> inline __device__ char FUNC SIGNATURE \
613  { \
614  int4 tmp = FUNC <int4> PARAMS; \
615  return (char)(tmp.x); \
616  } \
617  template<> inline __device__ unsigned short FUNC SIGNATURE \
618  { \
619  uint4 tmp = FUNC <uint4> PARAMS; \
620  return (unsigned short)(tmp.x); \
621  } \
622  template<> inline __device__ short FUNC SIGNATURE \
623  { \
624  int4 tmp = FUNC <int4> PARAMS; \
625  return (short)(tmp.x); \
626  } \
627  template<> inline __device__ int FUNC SIGNATURE \
628  { \
629  int4 tmp = FUNC <int4> PARAMS; \
630  return tmp.x; \
631  } \
632  template<> inline __device__ unsigned int FUNC SIGNATURE \
633  { \
634  uint4 tmp = FUNC <uint4> PARAMS; \
635  return tmp.x; \
636  } \
637  template<> inline __device__ uchar1 FUNC SIGNATURE \
638  { \
639  uint4 tmp = FUNC <uint4> PARAMS; \
640  return make_uchar1(tmp.x); \
641  } \
642  template<> inline __device__ char1 FUNC SIGNATURE \
643  { \
644  int4 tmp = FUNC <int4> PARAMS; \
645  return make_char1(tmp.x); \
646  } \
647  template<> inline __device__ ushort1 FUNC SIGNATURE \
648  { \
649  uint4 tmp = FUNC <uint4> PARAMS; \
650  return make_ushort1(tmp.x); \
651  } \
652  template<> inline __device__ short1 FUNC SIGNATURE \
653  { \
654  int4 tmp = FUNC <int4> PARAMS; \
655  return make_short1(tmp.x); \
656  } \
657  template<> inline __device__ uint1 FUNC SIGNATURE \
658  { \
659  uint4 tmp = FUNC <uint4> PARAMS; \
660  return make_uint1(tmp.x); \
661  } \
662  template<> inline __device__ int1 FUNC SIGNATURE \
663  { \
664  int4 tmp = FUNC <int4> PARAMS; \
665  return make_int1(tmp.x); \
666  } \
667  template<> inline __device__ float FUNC SIGNATURE \
668  { \
669  float4 tmp = FUNC <float4> PARAMS; \
670  return tmp.x; \
671  } \
672  template<> inline __device__ uchar2 FUNC SIGNATURE \
673  { \
674  uint4 tmp = FUNC <uint4> PARAMS; \
675  return make_uchar2(tmp.x, tmp.y); \
676  } \
677  template<> inline __device__ char2 FUNC SIGNATURE \
678  { \
679  int4 tmp = FUNC <int4> PARAMS; \
680  return make_char2(tmp.x, tmp.y); \
681  } \
682  template<> inline __device__ ushort2 FUNC SIGNATURE \
683  { \
684  uint4 tmp = FUNC <uint4> PARAMS; \
685  return make_ushort2(tmp.x, tmp.y); \
686  } \
687  template<> inline __device__ short2 FUNC SIGNATURE \
688  { \
689  int4 tmp = FUNC <int4> PARAMS; \
690  return make_short2(tmp.x, tmp.y); \
691  } \
692  template<> inline __device__ uint2 FUNC SIGNATURE \
693  { \
694  uint4 tmp = FUNC <uint4> PARAMS; \
695  return make_uint2(tmp.x, tmp.y); \
696  } \
697  template<> inline __device__ int2 FUNC SIGNATURE \
698  { \
699  int4 tmp = FUNC <int4> PARAMS; \
700  return make_int2(tmp.x, tmp.y); \
701  } \
702  template<> inline __device__ float2 FUNC SIGNATURE \
703  { \
704  float4 tmp = FUNC <float4> PARAMS; \
705  return ::make_float2(tmp.x, tmp.y); \
706  } \
707  template<> inline __device__ uchar4 FUNC SIGNATURE \
708  { \
709  uint4 tmp = FUNC <uint4> PARAMS; \
710  return make_uchar4(tmp.x, tmp.y, tmp.z, tmp.w); \
711  } \
712  template<> inline __device__ char4 FUNC SIGNATURE \
713  { \
714  int4 tmp = FUNC <int4> PARAMS; \
715  return make_char4(tmp.x, tmp.y, tmp.z, tmp.w); \
716  } \
717  template<> inline __device__ ushort4 FUNC SIGNATURE \
718  { \
719  uint4 tmp = FUNC <uint4> PARAMS; \
720  return make_ushort4(tmp.x, tmp.y, tmp.z, tmp.w); \
721  } \
722  template<> inline __device__ short4 FUNC SIGNATURE \
723  { \
724  int4 tmp = FUNC <int4> PARAMS; \
725  return make_short4(tmp.x, tmp.y, tmp.z, tmp.w); \
726  }
727 
728  inline __device__ int4 float4AsInt4( float4 f4 ) {
729  return make_int4(__float_as_int(f4.x), __float_as_int(f4.y), __float_as_int(f4.z), __float_as_int(f4.w));
730  }
731 
732  inline __device__ uint4 float4AsUInt4( float4 f4 ) {
733  return make_uint4(__float_as_int(f4.x), __float_as_int(f4.y), __float_as_int(f4.z), __float_as_int(f4.w));
734  }
735 
784  inline __device__ uint3 rtTexSize(rtTextureId id)
785  {
786  return optix::rt_texture_get_size_id(id);
787  }
788 
789  template<typename T>
790  inline __device__ T rtTex1D(rtTextureId id, float x);
791  template<> inline __device__ float4 rtTex1D(rtTextureId id, float x)
792  {
793  return optix::rt_texture_get_f_id(id, 1, x, 0, 0, 0);
794  }
795  template<> inline __device__ int4 rtTex1D(rtTextureId id, float x)
796  {
797  return optix::rt_texture_get_i_id(id, 1, x, 0, 0, 0);
798  }
799  template<> inline __device__ uint4 rtTex1D(rtTextureId id, float x)
800  {
801  return optix::rt_texture_get_u_id(id, 1, x, 0, 0, 0);
802  }
803  _OPTIX_TEX_FUNC_DECLARE_(rtTex1D, (rtTextureId id, float x), (id, x) )
804  template<typename T>
805  inline __device__ void rtTex1D(T* retVal, rtTextureId id, float x)
806  {
807  T tmp = rtTex1D<T>(id, x);
808  *retVal = tmp;
809  }
810 
811  template<typename T>
812  inline __device__ T rtTex1DFetch(rtTextureId id, int x);
813  template<> inline __device__ float4 rtTex1DFetch(rtTextureId id, int x)
814  {
815  return optix::rt_texture_get_fetch_id(id, 1, x, 0, 0, 0);
816  }
817  template<> inline __device__ int4 rtTex1DFetch(rtTextureId id, int x)
818  {
819  return float4AsInt4(optix::rt_texture_get_fetch_id(id, 1, x, 0, 0, 0));
820  }
821  template<> inline __device__ uint4 rtTex1DFetch(rtTextureId id, int x)
822  {
823  return float4AsUInt4(optix::rt_texture_get_fetch_id(id, 1, x, 0, 0, 0));
824  }
825  _OPTIX_TEX_FUNC_DECLARE_(rtTex1DFetch, (rtTextureId id, int x), (id, x) )
826  template<typename T>
827  inline __device__ void rtTex1DFetch(T* retVal, rtTextureId id, int x)
828  {
829  T tmp = rtTex1DFetch<T>(id, x);
830  *retVal = tmp;
831  }
832 
833  template<typename T>
834  inline __device__ T rtTex2D(rtTextureId id, float x, float y);
835  template<>
836  inline __device__ float4 rtTex2D(rtTextureId id, float x, float y)
837  {
838  return optix::rt_texture_get_f_id(id, 2, x, y, 0, 0);
839  }
840  template<>
841  inline __device__ int4 rtTex2D(rtTextureId id, float x, float y)
842  {
843  return optix::rt_texture_get_i_id(id, 2, x, y, 0, 0);
844  }
845  template<>
846  inline __device__ uint4 rtTex2D(rtTextureId id, float x, float y)
847  {
848  return optix::rt_texture_get_u_id(id, 2, x, y, 0, 0);
849  }
850  _OPTIX_TEX_FUNC_DECLARE_(rtTex2D, (rtTextureId id, float x, float y), (id, x, y) )
851  template<typename T>
852  inline __device__ void rtTex2D(T* retVal, rtTextureId id, float x, float y)
853  {
854  T tmp = rtTex2D<T>(id, x, y);
855  *retVal = tmp;
856  }
857 
858  template<typename T>
859  inline __device__ T rtTex2DFetch(rtTextureId id, int x, int y);
860  template<> inline __device__ float4 rtTex2DFetch(rtTextureId id, int x, int y)
861  {
862  return optix::rt_texture_get_fetch_id(id, 2, x, y, 0, 0);
863  }
864  template<> inline __device__ int4 rtTex2DFetch(rtTextureId id, int x, int y)
865  {
866  return float4AsInt4(optix::rt_texture_get_fetch_id(id, 2, x, y, 0, 0));
867  }
868  template<> inline __device__ uint4 rtTex2DFetch(rtTextureId id, int x, int y)
869  {
870  return float4AsUInt4(optix::rt_texture_get_fetch_id(id, 2, x, y, 0, 0));
871  }
872  _OPTIX_TEX_FUNC_DECLARE_(rtTex2DFetch, (rtTextureId id, int x, int y), (id, x, y) )
873  template<typename T>
874  inline __device__ void rtTex2DFetch(T* retVal, rtTextureId id, int x, int y)
875  {
876  T tmp = rtTex2DFetch<T>(id, x, y);
877  *retVal = tmp;
878  }
879 
880  template<typename T>
881  inline __device__ T rtTex3D(rtTextureId id, float x, float y, float z);
882  template<> inline __device__ float4 rtTex3D(rtTextureId id, float x, float y, float z)
883  {
884  return optix::rt_texture_get_f_id(id, 3, x, y, z, 0);
885  }
886  template<> inline __device__ int4 rtTex3D(rtTextureId id, float x, float y, float z)
887  {
888  return optix::rt_texture_get_i_id(id, 3, x, y, z, 0);
889  }
890  template<> inline __device__ uint4 rtTex3D(rtTextureId id, float x, float y, float z)
891  {
892  return optix::rt_texture_get_u_id(id, 3, x, y, z, 0);
893  }
894  _OPTIX_TEX_FUNC_DECLARE_(rtTex3D, (rtTextureId id, float x, float y, float z), (id, x, y, z) )
895  template<typename T>
896  inline __device__ void rtTex3D(T* retVal, rtTextureId id, float x, float y, float z)
897  {
898  T tmp = rtTex3D<T>(id, x, y, z);
899  *retVal = tmp;
900  }
901 
902  template<typename T>
903  inline __device__ T rtTex3DFetch(rtTextureId id, int x, int y, int z);
904  template<> inline __device__ float4 rtTex3DFetch(rtTextureId id, int x, int y, int z)
905  {
906  return optix::rt_texture_get_fetch_id(id, 3, x, y, z, 0);
907  }
908  template<> inline __device__ int4 rtTex3DFetch(rtTextureId id, int x, int y, int z)
909  {
910  return float4AsInt4(optix::rt_texture_get_fetch_id(id, 3, x, y, z, 0));
911  }
912  template<> inline __device__ uint4 rtTex3DFetch(rtTextureId id, int x, int y, int z)
913  {
914  return float4AsUInt4(optix::rt_texture_get_fetch_id(id, 3, x, y, z, 0));
915  }
916  _OPTIX_TEX_FUNC_DECLARE_(rtTex3DFetch, (rtTextureId id, int x, int y, int z), (id, x, y, z) )
917  template<typename T>
918  inline __device__ void rtTex3DFetch(T* retVal, rtTextureId id, int x, int y, int z)
919  {
920  T tmp = rtTex3DFetch<T>(id, x, y, z);
921  *retVal = tmp;
922  }
923 
924  template<typename T>
925  inline __device__ T rtTex2DGather(rtTextureId id, float x, float y, int comp = 0);
926  template<> inline __device__ float4 rtTex2DGather(rtTextureId id, float x, float y, int comp)
927  {
928  return optix::rt_texture_get_gather_id(id, x, y, comp);
929  }
930  template<> inline __device__ int4 rtTex2DGather(rtTextureId id, float x, float y, int comp)
931  {
932  return float4AsInt4(optix::rt_texture_get_gather_id(id, x, y, comp));
933  }
934  template<> inline __device__ uint4 rtTex2DGather(rtTextureId id, float x, float y, int comp)
935  {
936  return float4AsUInt4(optix::rt_texture_get_gather_id(id, x, y, comp));
937  }
938  _OPTIX_TEX_FUNC_DECLARE_(rtTex2DGather, (rtTextureId id, float x, float y, int comp), (id, x, y, comp) )
939  template<typename T>
940  inline __device__ void rtTex2DGather(T* retVal, rtTextureId id, float x, float y, int comp = 0)
941  {
942  T tmp = rtTex2DGather<T>(id, x, y, comp);
943  *retVal = tmp;
944  }
945 
946  template<typename T>
947  inline __device__ T rtTex1DGrad(rtTextureId id, float x, float dPdx, float dPdy);
948  template<> inline __device__ float4 rtTex1DGrad(rtTextureId id, float x, float dPdx, float dPdy)
949  {
950  return optix::rt_texture_get_grad_id(id, TEX_LOOKUP_1D, x, 0, 0, 0, dPdx, 0, 0, dPdy, 0, 0);
951  }
952  template<> inline __device__ int4 rtTex1DGrad(rtTextureId id, float x, float dPdx, float dPdy)
953  {
954  return float4AsInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_1D, x, 0, 0, 0, dPdx, 0, 0, dPdy, 0, 0));
955  }
956  template<> inline __device__ uint4 rtTex1DGrad(rtTextureId id, float x, float dPdx, float dPdy)
957  {
958  return float4AsUInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_1D, x, 0, 0, 0, dPdx, 0, 0, dPdy, 0, 0));
959  }
960  _OPTIX_TEX_FUNC_DECLARE_(rtTex1DGrad, (rtTextureId id, float x, float dPdx, float dPdy), (id, x, dPdx, dPdy) )
961  template<typename T>
962  inline __device__ void rtTex1DGrad(T* retVal, rtTextureId id, float x, float dPdx, float dPdy)
963  {
964  T tmp = rtTex1DGrad<T>(id, x, dPdx, dPdy);
965  *retVal = tmp;
966  }
967 
968  template<typename T>
969  inline __device__ T rtTex2DGrad(rtTextureId id, float x, float y, float2 dPdx, float2 dPdy);
970  template<> inline __device__ float4 rtTex2DGrad(rtTextureId id, float x, float y, float2 dPdx, float2 dPdy)
971  {
972  return optix::rt_texture_get_grad_id(id, TEX_LOOKUP_2D, x, y, 0, 0, dPdx.x, dPdx.y, 0, dPdy.x, dPdy.y, 0);
973  }
974  template<> inline __device__ int4 rtTex2DGrad(rtTextureId id, float x, float y, float2 dPdx, float2 dPdy)
975  {
976  return float4AsInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_2D, x, y, 0, 0, dPdx.x, dPdx.y, 0, dPdy.x, dPdy.y, 0));
977  }
978  template<> inline __device__ uint4 rtTex2DGrad(rtTextureId id, float x, float y, float2 dPdx, float2 dPdy)
979  {
980  return float4AsUInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_2D, x, y, 0, 0, dPdx.x, dPdx.y, 0, dPdy.x, dPdy.y, 0));
981  }
982  _OPTIX_TEX_FUNC_DECLARE_(rtTex2DGrad, (rtTextureId id, float x, float y, float2 dPdx, float2 dPdy), (id, x, y, dPdx, dPdy) )
983  template<typename T>
984  inline __device__ void rtTex2DGrad(T* retVal, rtTextureId id, float x, float y, float2 dPdx, float2 dPdy)
985  {
986  T tmp = rtTex2DGrad<T>(id, x, y, dPdx, dPdy);
987  *retVal = tmp;
988  }
989 
990  template<typename T>
991  inline __device__ T rtTex3DGrad(rtTextureId id, float x, float y, float z, float4 dPdx, float4 dPdy);
992  template<> inline __device__ float4 rtTex3DGrad(rtTextureId id, float x, float y, float z, float4 dPdx, float4 dPdy)
993  {
994  return optix::rt_texture_get_grad_id(id, TEX_LOOKUP_3D, x, y, z, 0, dPdx.x, dPdx.y, dPdx.z, dPdy.x, dPdy.y, dPdy.z);
995  }
996  template<> inline __device__ int4 rtTex3DGrad(rtTextureId id, float x, float y, float z, float4 dPdx, float4 dPdy)
997  {
998  return float4AsInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_3D, x, y, z, 0, dPdx.x, dPdx.y, dPdx.z, dPdy.x, dPdy.y, dPdy.z));
999  }
1000  template<> inline __device__ uint4 rtTex3DGrad(rtTextureId id, float x, float y, float z, float4 dPdx, float4 dPdy)
1001  {
1002  return float4AsUInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_3D, x, y, z, 0, dPdx.x, dPdx.y, dPdx.z, dPdy.x, dPdy.y, dPdy.z));
1003  }
1004  _OPTIX_TEX_FUNC_DECLARE_(rtTex3DGrad, (rtTextureId id, float x, float y, float z, float4 dPdx, float4 dPdy), (id, x, y, z, dPdx, dPdy) )
1005  template<typename T>
1006  inline __device__ void rtTex3DGrad(T* retVal, rtTextureId id, float x, float y, float z, float4 dPdx, float4 dPdy)
1007  {
1008  T tmp = rtTex3DGrad<T>(id, x, y, z, dPdx, dPdy);
1009  *retVal = tmp;
1010  }
1011 
1012  template<typename T>
1013  inline __device__ T rtTex1DLayeredGrad(rtTextureId id, float x, int layer, float dPdx, float dPdy);
1014  template<> inline __device__ float4 rtTex1DLayeredGrad(rtTextureId id, float x, int layer, float dPdx, float dPdy)
1015  {
1016  return optix::rt_texture_get_grad_id(id, TEX_LOOKUP_A1, x, 0, 0, layer, dPdx, 0, 0, dPdy, 0, 0);
1017  }
1018  template<> inline __device__ int4 rtTex1DLayeredGrad(rtTextureId id, float x, int layer, float dPdx, float dPdy)
1019  {
1020  return float4AsInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_A1, x, 0, 0, layer, dPdx, 0, 0, dPdy, 0, 0));
1021  }
1022  template<> inline __device__ uint4 rtTex1DLayeredGrad(rtTextureId id, float x, int layer, float dPdx, float dPdy)
1023  {
1024  return float4AsUInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_A1, x, 0, 0, layer, dPdx, 0, 0, dPdy, 0, 0));
1025  }
1026  _OPTIX_TEX_FUNC_DECLARE_(rtTex1DLayeredGrad, (rtTextureId id, float x, int layer, float dPdx, float dPdy), (id, x, layer, dPdx, dPdy) )
1027  template<typename T>
1028  inline __device__ void rtTex1DLayeredGrad(T* retVal, rtTextureId id, float x, int layer, float dPdx, float dPdy)
1029  {
1030  T tmp = rtTex1DLayeredGrad<T>(id, x, layer, dPdx, dPdy);
1031  *retVal = tmp;
1032  }
1033 
1034  template<typename T>
1035  inline __device__ T rtTex2DLayeredGrad(rtTextureId id, float x, float y, int layer, float2 dPdx, float2 dPdy);
1036  template<> inline __device__ float4 rtTex2DLayeredGrad(rtTextureId id, float x, float y, int layer, float2 dPdx, float2 dPdy)
1037  {
1038  return optix::rt_texture_get_grad_id(id, TEX_LOOKUP_A2, x, y, 0, layer, dPdx.x, dPdx.y, 0, dPdy.x, dPdy.y, 0);
1039  }
1040  template<> inline __device__ int4 rtTex2DLayeredGrad(rtTextureId id, float x, float y, int layer, float2 dPdx, float2 dPdy)
1041  {
1042  return float4AsInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_A2, x, y, 0, layer, dPdx.x, dPdx.y, 0, dPdy.x, dPdy.y, 0));
1043  }
1044  template<> inline __device__ uint4 rtTex2DLayeredGrad(rtTextureId id, float x, float y, int layer, float2 dPdx, float2 dPdy)
1045  {
1046  return float4AsUInt4(optix::rt_texture_get_grad_id(id, TEX_LOOKUP_A2, x, y, 0, layer, dPdx.x, dPdx.y, 0, dPdy.x, dPdy.y, 0));
1047  }
1048  _OPTIX_TEX_FUNC_DECLARE_(rtTex2DLayeredGrad, (rtTextureId id, float x, float y, int layer, float2 dPdx, float2 dPdy), (id, x, y, layer, dPdx, dPdy) )
1049  template<typename T>
1050  inline __device__ void rtTex2DLayeredGrad(T* retVal, rtTextureId id, float x, float y, int layer, float2 dPdx, float2 dPdy)
1051  {
1052  T tmp = rtTex2DLayeredGrad<T>(id, x, y, layer, dPdx, dPdy);
1053  *retVal = tmp;
1054  }
1055 
1056  template<typename T>
1057  inline __device__ T rtTex1DLod(rtTextureId id, float x, float level);
1058  template<> inline __device__ float4 rtTex1DLod(rtTextureId id, float x, float level)
1059  {
1060  return optix::rt_texture_get_level_id(id, TEX_LOOKUP_1D, x, 0, 0, 0, level );
1061  }
1062  template<> inline __device__ int4 rtTex1DLod(rtTextureId id, float x, float level)
1063  {
1064  return float4AsInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_1D, x, 0, 0, 0, level ));
1065  }
1066  template<> inline __device__ uint4 rtTex1DLod(rtTextureId id, float x, float level)
1067  {
1068  return float4AsUInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_1D, x, 0, 0, 0, level ));
1069  }
1070  _OPTIX_TEX_FUNC_DECLARE_(rtTex1DLod, (rtTextureId id, float x, float level), (id, x, level) )
1071  template<typename T>
1072  inline __device__ void rtTex1DLod(T* retVal, rtTextureId id, float x, float level)
1073  {
1074  T tmp = rtTex1DLod<T>(id, x, level);
1075  *retVal = tmp;
1076  }
1077 
1078  template<typename T>
1079  inline __device__ T rtTex2DLod(rtTextureId id, float x, float y, float level);
1080  template<> inline __device__ float4 rtTex2DLod(rtTextureId id, float x, float y, float level)
1081  {
1082  return optix::rt_texture_get_level_id(id, TEX_LOOKUP_2D, x, y, 0, 0, level );
1083  }
1084  template<> inline __device__ int4 rtTex2DLod(rtTextureId id, float x, float y, float level)
1085  {
1086  return float4AsInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_2D, x, y, 0, 0, level ));
1087  }
1088  template<> inline __device__ uint4 rtTex2DLod(rtTextureId id, float x, float y, float level)
1089  {
1090  return float4AsUInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_2D, x, y, 0, 0, level ));
1091  }
1092  _OPTIX_TEX_FUNC_DECLARE_(rtTex2DLod, (rtTextureId id, float x, float y, float level), (id, x, y, level) )
1093  template<typename T>
1094  inline __device__ void rtTex2DLod(T* retVal, rtTextureId id, float x, float y, float level)
1095  {
1096  T tmp = rtTex2DLod<T>(id, x, y, level);
1097  *retVal = tmp;
1098  }
1099 
1100  template<typename T>
1101  inline __device__ T rtTex3DLod(rtTextureId id, float x, float y, float z, float level);
1102  template<> inline __device__ float4 rtTex3DLod(rtTextureId id, float x, float y, float z, float level)
1103  {
1104  return optix::rt_texture_get_level_id(id, TEX_LOOKUP_3D, x, y, z, 0, level );
1105  }
1106  template<> inline __device__ int4 rtTex3DLod(rtTextureId id, float x, float y, float z, float level)
1107  {
1108  return float4AsInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_3D, x, y, z, 0, level ));
1109  }
1110  template<> inline __device__ uint4 rtTex3DLod(rtTextureId id, float x, float y, float z, float level)
1111  {
1112  return float4AsUInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_3D, x, y, z, 0, level ));
1113  }
1114  _OPTIX_TEX_FUNC_DECLARE_(rtTex3DLod, (rtTextureId id, float x, float y, float z, float level), (id, x, y, z, level) )
1115  template<typename T>
1116  inline __device__ void rtTex3DLod(T* retVal, rtTextureId id, float x, float y, float z, float level)
1117  {
1118  T tmp = rtTex3DLod<T>(id, x, y, z, level);
1119  *retVal = tmp;
1120  }
1121 
1122  template<typename T>
1123  inline __device__ T rtTex1DLayeredLod(rtTextureId id, float x, int layer, float level);
1124  template<> inline __device__ float4 rtTex1DLayeredLod(rtTextureId id, float x, int layer, float level)
1125  {
1126  return optix::rt_texture_get_level_id(id, TEX_LOOKUP_A1, x, 0, 0, layer, level );
1127  }
1128  template<> inline __device__ int4 rtTex1DLayeredLod(rtTextureId id, float x, int layer, float level)
1129  {
1130  return float4AsInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_A1, x, 0, 0, layer, level ));
1131  }
1132  template<> inline __device__ uint4 rtTex1DLayeredLod(rtTextureId id, float x, int layer, float level)
1133  {
1134  return float4AsUInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_A1, x, 0, 0, layer, level ));
1135  }
1136  _OPTIX_TEX_FUNC_DECLARE_(rtTex1DLayeredLod, (rtTextureId id, float x, int layer, float level), (id, x, layer, level) )
1137  template<typename T>
1138  inline __device__ void rtTex1DLayeredLod(T* retVal, rtTextureId id, float x, int layer, float level)
1139  {
1140  T tmp = rtTex1DLayeredLod<T>(id, x, layer, level);
1141  *retVal = tmp;
1142  }
1143 
1144  template<typename T>
1145  inline __device__ T rtTex2DLayeredLod(rtTextureId id, float x, float y, int layer, float level);
1146  template<> inline __device__ float4 rtTex2DLayeredLod(rtTextureId id, float x, float y, int layer, float level)
1147  {
1148  return optix::rt_texture_get_level_id(id, TEX_LOOKUP_A2, x, y, 0, layer, level );
1149  }
1150  template<> inline __device__ int4 rtTex2DLayeredLod(rtTextureId id, float x, float y, int layer, float level)
1151  {
1152  return float4AsInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_A2, x, y, 0, layer, level ));
1153  }
1154  template<> inline __device__ uint4 rtTex2DLayeredLod(rtTextureId id, float x, float y, int layer, float level)
1155  {
1156  return float4AsUInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_A2, x, y, 0, layer, level ));
1157  }
1158  _OPTIX_TEX_FUNC_DECLARE_(rtTex2DLayeredLod, (rtTextureId id, float x, float y, int layer, float level), (id, x, y, layer, level) )
1159  template<typename T>
1160  inline __device__ void rtTex2DLayeredLod(T* retVal, rtTextureId id, float x, float y, int layer, float level)
1161  {
1162  T tmp = rtTex2DLayeredLod<T>(id, x, y, layer, level);
1163  *retVal = tmp;
1164  }
1165 
1166  template<typename T>
1167  inline __device__ T rtTex1DLayered(rtTextureId id, float x, int layer);
1168  template<> inline __device__ float4 rtTex1DLayered(rtTextureId id, float x, int layer)
1169  {
1170  return optix::rt_texture_get_base_id(id, TEX_LOOKUP_A1, x, 0, 0, layer );
1171  }
1172  template<> inline __device__ int4 rtTex1DLayered(rtTextureId id, float x, int layer)
1173  {
1174  return float4AsInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_A1, x, 0, 0, layer ));
1175  }
1176  template<> inline __device__ uint4 rtTex1DLayered(rtTextureId id, float x, int layer)
1177  {
1178  return float4AsUInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_A1, x, 0, 0, layer ));
1179  }
1180  _OPTIX_TEX_FUNC_DECLARE_(rtTex1DLayered, (rtTextureId id, float x, int layer), (id, x, layer) )
1181  template<typename T>
1182  inline __device__ void rtTex1DLayered(T* retVal, rtTextureId id, float x, int layer)
1183  {
1184  T tmp = rtTex1DLayered<T>(id, x, layer);
1185  *retVal = tmp;
1186  }
1187 
1188  template<typename T>
1189  inline __device__ T rtTex2DLayered(rtTextureId id, float x, float y, int layer);
1190  template<> inline __device__ float4 rtTex2DLayered(rtTextureId id, float x, float y, int layer)
1191  {
1192  return optix::rt_texture_get_base_id(id, TEX_LOOKUP_A2, x, y, 0, layer );
1193  }
1194  template<> inline __device__ int4 rtTex2DLayered(rtTextureId id, float x, float y, int layer)
1195  {
1196  return float4AsInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_A2, x, y, 0, layer ));
1197  }
1198  template<> inline __device__ uint4 rtTex2DLayered(rtTextureId id, float x, float y, int layer)
1199  {
1200  return float4AsUInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_A2, x, y, 0, layer ));
1201  }
1202  _OPTIX_TEX_FUNC_DECLARE_(rtTex2DLayered, (rtTextureId id, float x, float y, int layer), (id, x, y, layer) )
1203  template<typename T>
1204  inline __device__ void rtTex2DLayered(T* retVal, rtTextureId id, float x, float y, int layer)
1205  {
1206  T tmp = rtTex2DLayered<T>(id, x, y, layer);
1207  *retVal = tmp;
1208  }
1209 
1210  template<typename T>
1211  inline __device__ T rtTexCubemap(rtTextureId id, float x, float y, float z);
1212  template<> inline __device__ float4 rtTexCubemap(rtTextureId id, float x, float y, float z)
1213  {
1214  return optix::rt_texture_get_base_id(id, TEX_LOOKUP_CUBE, x, y, z, 0 );
1215  }
1216  template<> inline __device__ int4 rtTexCubemap(rtTextureId id, float x, float y, float z)
1217  {
1218  return float4AsInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_CUBE, x, y, z, 0 ));
1219  }
1220  template<> inline __device__ uint4 rtTexCubemap(rtTextureId id, float x, float y, float z)
1221  {
1222  return float4AsUInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_CUBE, x, y, z, 0 ));
1223  }
1224  _OPTIX_TEX_FUNC_DECLARE_(rtTexCubemap, (rtTextureId id, float x, float y, float z), (id, x, y, z) )
1225  template<typename T>
1226  inline __device__ void rtTexCubemap(T* retVal, rtTextureId id, float x, float y, float z)
1227  {
1228  T tmp = rtTexCubemap<T>(id, x, y, z);
1229  *retVal = tmp;
1230  }
1231 
1232  template<typename T>
1233  inline __device__ T rtTexCubemapLayered(rtTextureId id, float x, float y, float z, int layer);
1234  template<> inline __device__ float4 rtTexCubemapLayered(rtTextureId id, float x, float y, float z, int layer)
1235  {
1236  return optix::rt_texture_get_base_id(id, TEX_LOOKUP_ACUBE, x, y, z, layer );
1237  }
1238  template<> inline __device__ int4 rtTexCubemapLayered(rtTextureId id, float x, float y, float z, int layer)
1239  {
1240  return float4AsInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_ACUBE, x, y, z, layer ));
1241  }
1242  template<> inline __device__ uint4 rtTexCubemapLayered(rtTextureId id, float x, float y, float z, int layer)
1243  {
1244  return float4AsUInt4(optix::rt_texture_get_base_id(id, TEX_LOOKUP_ACUBE, x, y, z, layer ));
1245  }
1246  _OPTIX_TEX_FUNC_DECLARE_(rtTexCubemapLayered, (rtTextureId id, float x, float y, float z, int layer), (id, x, y, z, layer) )
1247  template<typename T>
1248  inline __device__ void rtTexCubemapLayered(T* retVal, rtTextureId id, float x, float y, float z, int layer)
1249  {
1250  T tmp = rtTexCubemapLayered<T>(id, x, y, z, layer);
1251  *retVal = tmp;
1252  }
1253 
1254  template<typename T>
1255  inline __device__ T rtTexCubemapLod(rtTextureId id, float x, float y, float z, float level);
1256  template<> inline __device__ float4 rtTexCubemapLod(rtTextureId id, float x, float y, float z, float level)
1257  {
1258  return optix::rt_texture_get_level_id(id, TEX_LOOKUP_CUBE, x, y, z, 0, level );
1259  }
1260  template<> inline __device__ int4 rtTexCubemapLod(rtTextureId id, float x, float y, float z, float level)
1261  {
1262  return float4AsInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_CUBE, x, y, z, 0, level ));
1263  }
1264  template<> inline __device__ uint4 rtTexCubemapLod(rtTextureId id, float x, float y, float z, float level)
1265  {
1266  return float4AsUInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_CUBE, x, y, z, 0, level ));
1267  }
1268  _OPTIX_TEX_FUNC_DECLARE_(rtTexCubemapLod, (rtTextureId id, float x, float y, float z, float level), (id, x, y, z, level) )
1269  template<typename T>
1270  inline __device__ void rtTexCubemapLod(T* retVal, rtTextureId id, float x, float y, float z, float level)
1271  {
1272  T tmp = rtTexCubemapLod<T>(id, x, y, z, level);
1273  *retVal = tmp;
1274  }
1275 
1276  template<typename T>
1277  inline __device__ T rtTexCubemapLayeredLod(rtTextureId id, float x, float y, float z, int layer, float level);
1278  template<> inline __device__ float4 rtTexCubemapLayeredLod(rtTextureId id, float x, float y, float z, int layer, float level)
1279  {
1280  return optix::rt_texture_get_level_id(id, TEX_LOOKUP_ACUBE, x, y, z, layer, level );
1281  }
1282  template<> inline __device__ int4 rtTexCubemapLayeredLod(rtTextureId id, float x, float y, float z, int layer, float level)
1283  {
1284  return float4AsInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_ACUBE, x, y, z, layer, level ));
1285  }
1286  template<> inline __device__ uint4 rtTexCubemapLayeredLod(rtTextureId id, float x, float y, float z, int layer, float level)
1287  {
1288  return float4AsUInt4(optix::rt_texture_get_level_id(id, TEX_LOOKUP_ACUBE, x, y, z, layer, level ));
1289  }
1290  _OPTIX_TEX_FUNC_DECLARE_(rtTexCubemapLayeredLod, (rtTextureId id, float x, float y, float z, int layer, float level), (id, x, y, z, layer, level) )
1291  template<typename T>
1292  inline __device__ void rtTexCubemapLayeredLod(T* retVal, rtTextureId id, float x, float y, float z, int layer, float level)
1293  {
1294  T tmp = rtTexCubemapLayeredLod<T>(id, x, y, z, layer, level);
1295  *retVal = tmp;
1296  }
1297 
1300  #undef _OPTIX_TEX_FUNC_DECLARE_
1301 };
1302 
1303 /*
1304  Program
1305 */
1306 
1336 #define RT_PROGRAM __global__
1337 
1338 /* This is used to declare programs that can be attached to variables and called from
1339  * within other RT_PROGRAMS.
1340  *
1341  * There are some limitations with PTX that is targetted at sm_1x devices.
1342  *
1343  * 1. Functions declared with RT_CALLABLE_PROGRAM will not be emitted in the PTX unless
1344  * another function calls it. This can be fixed by declaring a __global__ helper
1345  * function that calls the desired function.
1346  *
1347  * RT_CALLABLE_PROGRAM
1348  * float3 simple_shade(float multiplier, float3 input_color)
1349  * {
1350  * return multiplier * input_color;
1351  * }
1352  *
1353  * #if __CUDA_ARCH__ < 200
1354  * __global__ void stub() {
1355  * (void) simple_shade( 0, make_float3(0,0,0) );
1356  * }
1357  * #endif
1358  *
1359  * 2. You can't pass pointers to functions or use integers for pointers. In the first
1360  * case CUDA will force the inline of the proxy function removing the call altogether,
1361  * and in the case of passing pointers as integers, CUDA will assume that any pointer
1362  * that was cast from an integer will point to global memory and could cause errors
1363  * when loading from that pointer. If you need to pass pointers, you should target
1364  * sm_20.
1365  */
1366 
1367 #define RT_CALLABLE_PROGRAM __device__ __noinline__
1368 
1369 
1370 namespace rti_internal_callableprogram {
1371  /* Any classes or types in the rti_internal_callableprogram namespace are used to help
1372  * implement callable program features and should not be used directly.
1373  */
1374 
1375 
1376  /* CPArgVoid a special class to act as an unspecified argument type that we can
1377  * statically query to determine if we have called our function with the wrong number of
1378  * arguments.
1379  */
1380 
1381  class CPArgVoid {};
1382  template< typename T1>
1383  struct is_CPArgVoid { static const bool result = false; };
1384 
1385  template<>
1386  struct is_CPArgVoid<CPArgVoid> { static const bool result = true; };
1387 
1388  template< bool Condition, typename Dummy = void >
1389  struct check_is_CPArgVoid {
1390  typedef bool result;
1391  };
1392 
1393  template<typename IntentionalError>
1394  struct check_is_CPArgVoid<false,IntentionalError> {
1395  typedef typename IntentionalError::does_not_exist result;
1396  };
1397 
1398  /* callableProgramIdBase is the underlying class for handling both bound and bindless
1399  * callable program calls. It should not be used directly, but instead the derived
1400  * classes of rtCallableProgramId and rtCallableProgramX should be used.
1401  */
1402  template <typename ReturnT
1403  ,typename Arg0T=rti_internal_callableprogram::CPArgVoid
1404  ,typename Arg1T=rti_internal_callableprogram::CPArgVoid
1405  ,typename Arg2T=rti_internal_callableprogram::CPArgVoid
1406  ,typename Arg3T=rti_internal_callableprogram::CPArgVoid
1407  ,typename Arg4T=rti_internal_callableprogram::CPArgVoid
1408  ,typename Arg5T=rti_internal_callableprogram::CPArgVoid
1409  ,typename Arg6T=rti_internal_callableprogram::CPArgVoid
1410  ,typename Arg7T=rti_internal_callableprogram::CPArgVoid
1411  ,typename Arg8T=rti_internal_callableprogram::CPArgVoid
1412  ,typename Arg9T=rti_internal_callableprogram::CPArgVoid
1413  >
1414  class callableProgramIdBase {
1415  public:
1416 #if (CUDART_VERSION >= 4010)
1417  // Default constructor
1418  __device__ __forceinline__ callableProgramIdBase() {}
1419  // Constructor that initializes the id with null.
1420  __device__ __forceinline__ callableProgramIdBase(RTprogramidnull nullid) { m_id = (int)nullid; }
1421  // Constructor that initializes the id.
1422  __device__ __forceinline__ explicit callableProgramIdBase(int id) : m_id(id) {}
1423 #endif
1424 
1426  // Call operators
1427  //
1428  // If you call the function with the wrong number of argument, you will get a
1429  // compilation error. If you have too many, then you will warned that an argument
1430  // doesn't match the CPArgVoid type. If you have too few, then the check_is_CPArgVoid
1431  // typedef will error out complaining that check_is_CPArgVoid::result isn't a type.
1432  __device__ __forceinline__ ReturnT operator()()
1433  {
1434  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg0T>::result>::result Arg0_test;
1435  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg1T>::result>::result Arg1_test;
1436  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg2T>::result>::result Arg2_test;
1437  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg3T>::result>::result Arg3_test;
1438  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg4T>::result>::result Arg4_test;
1439  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg5T>::result>::result Arg5_test;
1440  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg6T>::result>::result Arg6_test;
1441  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1442  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1443  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1444  typedef ReturnT (*funcT)();
1445  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1446  return call();
1447  }
1448  __device__ __forceinline__ ReturnT operator()(Arg0T arg0)
1449  {
1450  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg1T>::result>::result Arg1_test;
1451  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg2T>::result>::result Arg2_test;
1452  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg3T>::result>::result Arg3_test;
1453  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg4T>::result>::result Arg4_test;
1454  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg5T>::result>::result Arg5_test;
1455  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg6T>::result>::result Arg6_test;
1456  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1457  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1458  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1459  typedef ReturnT (*funcT)(Arg0T);
1460  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1461  return call(arg0);
1462  }
1463  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1)
1464  {
1465  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg2T>::result>::result Arg2_test;
1466  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg3T>::result>::result Arg3_test;
1467  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg4T>::result>::result Arg4_test;
1468  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg5T>::result>::result Arg5_test;
1469  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg6T>::result>::result Arg6_test;
1470  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1471  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1472  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1473  typedef ReturnT (*funcT)(Arg0T,Arg1T);
1474  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1475  return call(arg0,arg1);
1476  }
1477  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2)
1478  {
1479  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg3T>::result>::result Arg3_test;
1480  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg4T>::result>::result Arg4_test;
1481  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg5T>::result>::result Arg5_test;
1482  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg6T>::result>::result Arg6_test;
1483  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1484  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1485  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1486  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T);
1487  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1488  return call(arg0,arg1,arg2);
1489  }
1490  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2, Arg3T arg3)
1491  {
1492  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg4T>::result>::result Arg4_test;
1493  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg5T>::result>::result Arg5_test;
1494  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg6T>::result>::result Arg6_test;
1495  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1496  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1497  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1498  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T,Arg3T);
1499  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1500  return call(arg0,arg1,arg2,arg3);
1501  }
1502  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2, Arg3T arg3,
1503  Arg4T arg4)
1504  {
1505  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg5T>::result>::result Arg5_test;
1506  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg6T>::result>::result Arg6_test;
1507  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1508  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1509  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1510  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T);
1511  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1512  return call(arg0,arg1,arg2,arg3,arg4);
1513  }
1514  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2, Arg3T arg3,
1515  Arg4T arg4, Arg5T arg5)
1516  {
1517  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg6T>::result>::result Arg6_test;
1518  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1519  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1520  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1521  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T);
1522  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1523  return call(arg0,arg1,arg2,arg3,arg4,arg5);
1524  }
1525  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2, Arg3T arg3,
1526  Arg4T arg4, Arg5T arg5, Arg6T arg6)
1527  {
1528  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg7T>::result>::result Arg7_test;
1529  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1530  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1531  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T);
1532  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1533  return call(arg0,arg1,arg2,arg3,arg4,arg5,arg6);
1534  }
1535  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2, Arg3T arg3,
1536  Arg4T arg4, Arg5T arg5, Arg6T arg6, Arg7T arg7)
1537  {
1538  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg8T>::result>::result Arg8_test;
1539  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1540  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T);
1541  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1542  return call(arg0,arg1,arg2,arg3,arg4,arg5,arg6,arg7);
1543  }
1544  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2, Arg3T arg3,
1545  Arg4T arg4, Arg5T arg5, Arg6T arg6, Arg7T arg7,
1546  Arg8T arg8)
1547  {
1548  typedef typename check_is_CPArgVoid<is_CPArgVoid<Arg9T>::result>::result Arg9_test;
1549  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T);
1550  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1551  return call(arg0,arg1,arg2,arg3,arg4,arg5,arg6,arg7,arg8);
1552  }
1553  __device__ __forceinline__ ReturnT operator()(Arg0T arg0, Arg1T arg1, Arg2T arg2, Arg3T arg3,
1554  Arg4T arg4, Arg5T arg5, Arg6T arg6, Arg7T arg7,
1555  Arg8T arg8, Arg9T arg9)
1556  {
1557  typedef ReturnT (*funcT)(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T,Arg9T);
1558  funcT call = (funcT)optix::rt_callable_program_from_id(m_id);
1559  return call(arg0,arg1,arg2,arg3,arg4,arg5,arg6,arg7,arg8,arg9);
1560  }
1561  protected:
1562  int m_id;
1563  };
1564 } // end namespace rti_internal_callableprogram
1565 
1566 namespace optix {
1567 
1568  /* RT_INTERNAL_CALLABLE_PROGRAM_DEFS is a helper macro to define the body of each
1569  * version of the callableProgramId class. Variadic macro arguments are our friend here,
1570  * so we can use arguments such as (ReturnT,Arg0T) and (ReturnT,Arg0T,Arg1T) and get the
1571  * correct template types defined.
1572  */
1573 #define RT_INTERNAL_CALLABLE_PROGRAM_DEFS(...) public rti_internal_callableprogram::callableProgramIdBase<__VA_ARGS__> \
1574  { \
1575  public: \
1576  /* Default constructor */ \
1577  __device__ __forceinline__ callableProgramId() {} \
1578  /* Constructor that initializes the id with null.*/ \
1579  __device__ __forceinline__ callableProgramId(RTprogramidnull nullid) \
1580  : rti_internal_callableprogram::callableProgramIdBase<__VA_ARGS__>(nullid) {} \
1581  /* Constructor that initializes the id.*/ \
1582  __device__ __forceinline__ explicit callableProgramId(int id) \
1583  : rti_internal_callableprogram::callableProgramIdBase<__VA_ARGS__>(id) {} \
1584  /* assigment that initializes the id with null. */ \
1585  __device__ __forceinline__ callableProgramId& operator= (RTprogramidnull nullid) \
1586  { this->m_id = nullid; return *this; } \
1587  /* Return the id */ \
1588  __device__ __forceinline__ int getId() const { return this->m_id; } \
1589  /* Return whether the id is valid */ \
1590  __device__ __forceinline__ operator bool() const \
1591  { return this->m_id != RT_PROGRAM_ID_NULL; } \
1592  }
1593 
1594  /* callableProgramId should not be used directly. Use rtCallableProgramId instead to
1595  * make sure compatibility with future versions of OptiX is maintained.
1596  */
1597 
1598  /* The default template version is left undefined on purpose. Only the specialized
1599  * versions should be used. */
1600  template<typename T>
1601  class callableProgramId;
1602 
1603  /* These are specializations designed to be used like: <ReturnT(argument types)> */
1604  template<typename ReturnT>
1605  class callableProgramId<ReturnT()>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT);
1606  template<typename ReturnT, typename Arg0T>
1607  class callableProgramId<ReturnT(Arg0T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T);
1608  template<typename ReturnT, typename Arg0T, typename Arg1T>
1609  class callableProgramId<ReturnT(Arg0T,Arg1T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T);
1610  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T>
1611  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T);
1612  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T>
1613  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T);
1614  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1615  typename Arg4T>
1616  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T);
1617  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1618  typename Arg4T, typename Arg5T>
1619  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T);
1620  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1621  typename Arg4T, typename Arg5T, typename Arg6T>
1622  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T);
1623  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1624  typename Arg4T, typename Arg5T, typename Arg6T, typename Arg7T>
1625  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T);
1626  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1627  typename Arg4T, typename Arg5T, typename Arg6T, typename Arg7T, typename Arg8T>
1628  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T);
1629  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1630  typename Arg4T, typename Arg5T, typename Arg6T, typename Arg7T, typename Arg8T, typename Arg9T>
1631  class callableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T,Arg9T)>: RT_INTERNAL_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T,Arg9T);
1632 
1633  /* RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS is a helper macro to define the body of each
1634  * version of the boundCallableProgramId class. Variadic macro arguments are our friend
1635  * here, so we can use arguments such as (ReturnT,Arg0T) and (ReturnT,Arg0T,Arg1T) and
1636  * get the correct template types defined.
1637  *
1638  * Also, the constructors (except the default) and operators are made private, because
1639  * the objects should not be set, copied or otherwise changed from the value set by
1640  * OptiX from the host.
1641  *
1642  * The getId and bool operator (from the parent) are redefined and made private, because
1643  * the internal ID (m_id) should never be accessible. Using this ID is likely to cause
1644  * problems, because OptiX is free to compile the called program in a method that would
1645  * be incompatible with bindless callable programs.
1646  */
1647 #if (CUDART_VERSION >= 4010)
1648 #define RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(...) public rti_internal_callableprogram::callableProgramIdBase<__VA_ARGS__> \
1649  { \
1650  public: \
1651  /* Default constructor */ \
1652  __device__ __forceinline__ boundCallableProgramId() {} \
1653  private: \
1654  /* No copying of this class*/ \
1655  __device__ __forceinline__ boundCallableProgramId(const boundCallableProgramId& ); \
1656  __device__ __forceinline__ boundCallableProgramId& operator= (const boundCallableProgramId& ); \
1657  }
1658 #else
1659 #define RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(...) public rti_internal_callableprogram::callableProgramIdBase<__VA_ARGS__> \
1660  { \
1661  }
1662 #endif
1663 
1664  /* boundCallableProgramId should not be used directly. Use rtCallableProgramX
1665  * instead to make sure compatibility with future versions of OptiX is maintained.
1666  */
1667 
1668  /* The default template version is left undefined on purpose. Only the specialized
1669  * versions should be used. */
1670  template<typename T>
1671  class boundCallableProgramId;
1672 
1673  /* These are specializations designed to be used like: <ReturnT(argument types)> */
1674  template<typename ReturnT>
1675  class boundCallableProgramId<ReturnT()>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT);
1676  template<typename ReturnT, typename Arg0T>
1677  class boundCallableProgramId<ReturnT(Arg0T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T);
1678  template<typename ReturnT, typename Arg0T, typename Arg1T>
1679  class boundCallableProgramId<ReturnT(Arg0T,Arg1T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T);
1680  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T>
1681  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T);
1682  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T>
1683  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T);
1684  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1685  typename Arg4T>
1686  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T);
1687  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1688  typename Arg4T, typename Arg5T>
1689  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T);
1690  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1691  typename Arg4T, typename Arg5T, typename Arg6T>
1692  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T);
1693  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1694  typename Arg4T, typename Arg5T, typename Arg6T, typename Arg7T>
1695  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T);
1696  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1697  typename Arg4T, typename Arg5T, typename Arg6T, typename Arg7T, typename Arg8T>
1698  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T);
1699  template<typename ReturnT, typename Arg0T, typename Arg1T, typename Arg2T, typename Arg3T,
1700  typename Arg4T, typename Arg5T, typename Arg6T, typename Arg7T, typename Arg8T, typename Arg9T>
1701  class boundCallableProgramId<ReturnT(Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T,Arg9T)>: RT_INTERNAL_BOUND_CALLABLE_PROGRAM_DEFS(ReturnT,Arg0T,Arg1T,Arg2T,Arg3T,Arg4T,Arg5T,Arg6T,Arg7T,Arg8T,Arg9T);
1702 
1703 } // end namespace optix
1704 
1705 namespace rti_internal_typeinfo {
1706  // Specialization for callableProgramId types
1707  template <typename T>
1708  struct rti_typeenum<optix::callableProgramId<T> >
1709  {
1710  static const int m_typeenum = _OPTIX_TYPE_ENUM_PROGRAM_ID;
1711  };
1712 
1713  // Specialization for boundCallableProgramId types
1714  template <typename T>
1715  struct rti_typeenum<optix::boundCallableProgramId<T> >
1716  {
1717  static const int m_typeenum = _OPTIX_TYPE_ENUM_PROGRAM_AS_ID;
1718  };
1719 
1720 }
1721 
1751 #define rtCallableProgramId optix::callableProgramId
1752 
1787 #define rtCallableProgramX optix::boundCallableProgramId
1788 #ifdef RT_USE_TEMPLATED_RTCALLABLEPROGRAM
1789 # undef rtCallableProgram
1790 # define rtCallableProgram optix::boundCallableProgramId
1791 #endif
1792 
1793 /*
1794  Functions
1795 */
1796 
1824 template<class T>
1825 static inline __device__ void rtTrace( rtObject topNode, optix::Ray ray, T& prd )
1826 {
1827  optix::rt_trace(*(unsigned int*)&topNode, ray.origin, ray.direction, ray.ray_type, ray.tmin, ray.tmax, &prd, sizeof(T));
1828 }
1829 
1875 static inline __device__ bool rtPotentialIntersection( float tmin )
1876 {
1877  return optix::rt_potential_intersection( tmin );
1878 }
1879 
1904 static inline __device__ bool rtReportIntersection( unsigned int material )
1905 {
1906  return optix::rt_report_intersection( material );
1907 }
1908 
1940 static inline __device__ void rtIgnoreIntersection()
1941 {
1942  optix::rt_ignore_intersection();
1943 }
1944 
1969 static inline __device__ void rtTerminateRay()
1970 {
1971  optix::rt_terminate_ray();
1972 }
1973 
2009 static inline __device__ void rtIntersectChild( unsigned int index )
2010 {
2011  optix::rt_intersect_child( index );
2012 }
2013 
2049 static inline __device__ float3 rtTransformPoint( RTtransformkind kind, const float3& p )
2050 {
2051  return optix::rt_transform_point( kind, p );
2052 }
2053 
2090 static inline __device__ float3 rtTransformVector( RTtransformkind kind, const float3& v )
2091 {
2092  return optix::rt_transform_vector( kind, v );
2093 }
2094 
2131 static inline __device__ float3 rtTransformNormal( RTtransformkind kind, const float3& n )
2132 {
2133  return optix::rt_transform_normal( kind, n );
2134 }
2135 
2171 static inline __device__ void rtGetTransform( RTtransformkind kind, float matrix[16] )
2172 {
2173  return optix::rt_get_transform( kind, matrix );
2174 }
2175 
2176 
2177 /*
2178  Printing
2179 */
2180 
2213 static inline __device__ void rtPrintf( const char* fmt )
2214 {
2215  _RT_PRINTF_1();
2216  optix::rt_print_start(fmt,sz);
2217 }
2218 template<typename T1>
2219 static inline __device__ void rtPrintf( const char* fmt, T1 arg1 )
2220 {
2221  _RT_PRINTF_1();
2222  _RT_PRINTF_ARG_1( arg1 );
2223  _RT_PRINTF_2();
2224  _RT_PRINTF_ARG_2( arg1 );
2225 }
2226 template<typename T1, typename T2>
2227 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2 )
2228 {
2229  _RT_PRINTF_1();
2230  _RT_PRINTF_ARG_1( arg1 );
2231  _RT_PRINTF_ARG_1( arg2 );
2232  _RT_PRINTF_2();
2233  _RT_PRINTF_ARG_2( arg1 );
2234  _RT_PRINTF_ARG_2( arg2 );
2235 }
2236 template<typename T1, typename T2, typename T3>
2237 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3 )
2238 {
2239  _RT_PRINTF_1();
2240  _RT_PRINTF_ARG_1( arg1 );
2241  _RT_PRINTF_ARG_1( arg2 );
2242  _RT_PRINTF_ARG_1( arg3 );
2243  _RT_PRINTF_2();
2244  _RT_PRINTF_ARG_2( arg1 );
2245  _RT_PRINTF_ARG_2( arg2 );
2246  _RT_PRINTF_ARG_2( arg3 );
2247 }
2248 template<typename T1, typename T2, typename T3, typename T4>
2249 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4 )
2250 {
2251  _RT_PRINTF_1();
2252  _RT_PRINTF_ARG_1( arg1 );
2253  _RT_PRINTF_ARG_1( arg2 );
2254  _RT_PRINTF_ARG_1( arg3 );
2255  _RT_PRINTF_ARG_1( arg4 );
2256  _RT_PRINTF_2();
2257  _RT_PRINTF_ARG_2( arg1 );
2258  _RT_PRINTF_ARG_2( arg2 );
2259  _RT_PRINTF_ARG_2( arg3 );
2260  _RT_PRINTF_ARG_2( arg4 );
2261 }
2262 template<typename T1, typename T2, typename T3, typename T4, typename T5>
2263 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5 )
2264 {
2265  _RT_PRINTF_1();
2266  _RT_PRINTF_ARG_1( arg1 );
2267  _RT_PRINTF_ARG_1( arg2 );
2268  _RT_PRINTF_ARG_1( arg3 );
2269  _RT_PRINTF_ARG_1( arg4 );
2270  _RT_PRINTF_ARG_1( arg5 );
2271  _RT_PRINTF_2();
2272  _RT_PRINTF_ARG_2( arg1 );
2273  _RT_PRINTF_ARG_2( arg2 );
2274  _RT_PRINTF_ARG_2( arg3 );
2275  _RT_PRINTF_ARG_2( arg4 );
2276  _RT_PRINTF_ARG_2( arg5 );
2277 }
2278 template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6>
2279 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6 )
2280 {
2281  _RT_PRINTF_1();
2282  _RT_PRINTF_ARG_1( arg1 );
2283  _RT_PRINTF_ARG_1( arg2 );
2284  _RT_PRINTF_ARG_1( arg3 );
2285  _RT_PRINTF_ARG_1( arg4 );
2286  _RT_PRINTF_ARG_1( arg5 );
2287  _RT_PRINTF_ARG_1( arg6 );
2288  _RT_PRINTF_2();
2289  _RT_PRINTF_ARG_2( arg1 );
2290  _RT_PRINTF_ARG_2( arg2 );
2291  _RT_PRINTF_ARG_2( arg3 );
2292  _RT_PRINTF_ARG_2( arg4 );
2293  _RT_PRINTF_ARG_2( arg5 );
2294  _RT_PRINTF_ARG_2( arg6 );
2295 }
2296 template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7>
2297 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7 )
2298 {
2299  _RT_PRINTF_1();
2300  _RT_PRINTF_ARG_1( arg1 );
2301  _RT_PRINTF_ARG_1( arg2 );
2302  _RT_PRINTF_ARG_1( arg3 );
2303  _RT_PRINTF_ARG_1( arg4 );
2304  _RT_PRINTF_ARG_1( arg5 );
2305  _RT_PRINTF_ARG_1( arg6 );
2306  _RT_PRINTF_ARG_1( arg7 );
2307  _RT_PRINTF_2();
2308  _RT_PRINTF_ARG_2( arg1 );
2309  _RT_PRINTF_ARG_2( arg2 );
2310  _RT_PRINTF_ARG_2( arg3 );
2311  _RT_PRINTF_ARG_2( arg4 );
2312  _RT_PRINTF_ARG_2( arg5 );
2313  _RT_PRINTF_ARG_2( arg6 );
2314  _RT_PRINTF_ARG_2( arg7 );
2315 }
2316 template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8>
2317 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8 )
2318 {
2319  _RT_PRINTF_1();
2320  _RT_PRINTF_ARG_1( arg1 );
2321  _RT_PRINTF_ARG_1( arg2 );
2322  _RT_PRINTF_ARG_1( arg3 );
2323  _RT_PRINTF_ARG_1( arg4 );
2324  _RT_PRINTF_ARG_1( arg5 );
2325  _RT_PRINTF_ARG_1( arg6 );
2326  _RT_PRINTF_ARG_1( arg7 );
2327  _RT_PRINTF_ARG_1( arg8 );
2328  _RT_PRINTF_2();
2329  _RT_PRINTF_ARG_2( arg1 );
2330  _RT_PRINTF_ARG_2( arg2 );
2331  _RT_PRINTF_ARG_2( arg3 );
2332  _RT_PRINTF_ARG_2( arg4 );
2333  _RT_PRINTF_ARG_2( arg5 );
2334  _RT_PRINTF_ARG_2( arg6 );
2335  _RT_PRINTF_ARG_2( arg7 );
2336  _RT_PRINTF_ARG_2( arg8 );
2337 }
2338 template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
2339 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9 )
2340 {
2341  _RT_PRINTF_1();
2342  _RT_PRINTF_ARG_1( arg1 );
2343  _RT_PRINTF_ARG_1( arg2 );
2344  _RT_PRINTF_ARG_1( arg3 );
2345  _RT_PRINTF_ARG_1( arg4 );
2346  _RT_PRINTF_ARG_1( arg5 );
2347  _RT_PRINTF_ARG_1( arg6 );
2348  _RT_PRINTF_ARG_1( arg7 );
2349  _RT_PRINTF_ARG_1( arg8 );
2350  _RT_PRINTF_ARG_1( arg9 );
2351  _RT_PRINTF_2();
2352  _RT_PRINTF_ARG_2( arg1 );
2353  _RT_PRINTF_ARG_2( arg2 );
2354  _RT_PRINTF_ARG_2( arg3 );
2355  _RT_PRINTF_ARG_2( arg4 );
2356  _RT_PRINTF_ARG_2( arg5 );
2357  _RT_PRINTF_ARG_2( arg6 );
2358  _RT_PRINTF_ARG_2( arg7 );
2359  _RT_PRINTF_ARG_2( arg8 );
2360  _RT_PRINTF_ARG_2( arg9 );
2361 }
2362 template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9, typename T10>
2363 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9, T10 arg10 )
2364 {
2365  _RT_PRINTF_1();
2366  _RT_PRINTF_ARG_1( arg1 );
2367  _RT_PRINTF_ARG_1( arg2 );
2368  _RT_PRINTF_ARG_1( arg3 );
2369  _RT_PRINTF_ARG_1( arg4 );
2370  _RT_PRINTF_ARG_1( arg5 );
2371  _RT_PRINTF_ARG_1( arg6 );
2372  _RT_PRINTF_ARG_1( arg7 );
2373  _RT_PRINTF_ARG_1( arg8 );
2374  _RT_PRINTF_ARG_1( arg9 );
2375  _RT_PRINTF_ARG_1( arg10 );
2376  _RT_PRINTF_2();
2377  _RT_PRINTF_ARG_2( arg1 );
2378  _RT_PRINTF_ARG_2( arg2 );
2379  _RT_PRINTF_ARG_2( arg3 );
2380  _RT_PRINTF_ARG_2( arg4 );
2381  _RT_PRINTF_ARG_2( arg5 );
2382  _RT_PRINTF_ARG_2( arg6 );
2383  _RT_PRINTF_ARG_2( arg7 );
2384  _RT_PRINTF_ARG_2( arg8 );
2385  _RT_PRINTF_ARG_2( arg9 );
2386  _RT_PRINTF_ARG_2( arg10 );
2387 }
2388 template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9, typename T10, typename T11>
2389 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9, T10 arg10, T11 arg11 )
2390 {
2391  _RT_PRINTF_1();
2392  _RT_PRINTF_ARG_1( arg1 );
2393  _RT_PRINTF_ARG_1( arg2 );
2394  _RT_PRINTF_ARG_1( arg3 );
2395  _RT_PRINTF_ARG_1( arg4 );
2396  _RT_PRINTF_ARG_1( arg5 );
2397  _RT_PRINTF_ARG_1( arg6 );
2398  _RT_PRINTF_ARG_1( arg7 );
2399  _RT_PRINTF_ARG_1( arg8 );
2400  _RT_PRINTF_ARG_1( arg9 );
2401  _RT_PRINTF_ARG_1( arg10 );
2402  _RT_PRINTF_ARG_1( arg11 );
2403  _RT_PRINTF_2();
2404  _RT_PRINTF_ARG_2( arg1 );
2405  _RT_PRINTF_ARG_2( arg2 );
2406  _RT_PRINTF_ARG_2( arg3 );
2407  _RT_PRINTF_ARG_2( arg4 );
2408  _RT_PRINTF_ARG_2( arg5 );
2409  _RT_PRINTF_ARG_2( arg6 );
2410  _RT_PRINTF_ARG_2( arg7 );
2411  _RT_PRINTF_ARG_2( arg8 );
2412  _RT_PRINTF_ARG_2( arg9 );
2413  _RT_PRINTF_ARG_2( arg10 );
2414  _RT_PRINTF_ARG_2( arg11 );
2415 }
2416 template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9, typename T10, typename T11, typename T12>
2417 static inline __device__ void rtPrintf( const char* fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9, T10 arg10, T11 arg11, T12 arg12 )
2418 {
2419  _RT_PRINTF_1();
2420  _RT_PRINTF_ARG_1( arg1 );
2421  _RT_PRINTF_ARG_1( arg2 );
2422  _RT_PRINTF_ARG_1( arg3 );
2423  _RT_PRINTF_ARG_1( arg4 );
2424  _RT_PRINTF_ARG_1( arg5 );
2425  _RT_PRINTF_ARG_1( arg6 );
2426  _RT_PRINTF_ARG_1( arg7 );
2427  _RT_PRINTF_ARG_1( arg8 );
2428  _RT_PRINTF_ARG_1( arg9 );
2429  _RT_PRINTF_ARG_1( arg10 );
2430  _RT_PRINTF_ARG_1( arg11 );
2431  _RT_PRINTF_ARG_1( arg12 );
2432  _RT_PRINTF_2();
2433  _RT_PRINTF_ARG_2( arg1 );
2434  _RT_PRINTF_ARG_2( arg2 );
2435  _RT_PRINTF_ARG_2( arg3 );
2436  _RT_PRINTF_ARG_2( arg4 );
2437  _RT_PRINTF_ARG_2( arg5 );
2438  _RT_PRINTF_ARG_2( arg6 );
2439  _RT_PRINTF_ARG_2( arg7 );
2440  _RT_PRINTF_ARG_2( arg8 );
2441  _RT_PRINTF_ARG_2( arg9 );
2442  _RT_PRINTF_ARG_2( arg10 );
2443  _RT_PRINTF_ARG_2( arg11 );
2444  _RT_PRINTF_ARG_2( arg12 );
2445 }
2448 #undef _RT_PRINTF_1
2449 #undef _RT_PRINTF_2
2450 #undef _RT_PRINTF_ARG_1
2451 #undef _RT_PRINTF_ARG_2
2452 
2454 namespace rti_internal_register {
2455  extern __device__ void* reg_bitness_detector;
2456  extern __device__ volatile unsigned long long reg_exception_64_detail0;
2457  extern __device__ volatile unsigned long long reg_exception_64_detail1;
2458  extern __device__ volatile unsigned long long reg_exception_64_detail2;
2459  extern __device__ volatile unsigned long long reg_exception_64_detail3;
2460  extern __device__ volatile unsigned long long reg_exception_64_detail4;
2461  extern __device__ volatile unsigned long long reg_exception_64_detail5;
2462  extern __device__ volatile unsigned long long reg_exception_64_detail6;
2463  extern __device__ volatile unsigned long long reg_exception_64_detail7;
2464  extern __device__ volatile unsigned long long reg_exception_64_detail8;
2465  extern __device__ volatile unsigned long long reg_exception_64_detail9;
2466  extern __device__ volatile unsigned int reg_exception_detail0;
2467  extern __device__ volatile unsigned int reg_exception_detail1;
2468  extern __device__ volatile unsigned int reg_exception_detail2;
2469  extern __device__ volatile unsigned int reg_exception_detail3;
2470  extern __device__ volatile unsigned int reg_exception_detail4;
2471  extern __device__ volatile unsigned int reg_exception_detail5;
2472  extern __device__ volatile unsigned int reg_exception_detail6;
2473  extern __device__ volatile unsigned int reg_exception_detail7;
2474  extern __device__ volatile unsigned int reg_exception_detail8;
2475  extern __device__ volatile unsigned int reg_exception_detail9;
2476  extern __device__ volatile unsigned int reg_rayIndex_x;
2477  extern __device__ volatile unsigned int reg_rayIndex_y;
2478  extern __device__ volatile unsigned int reg_rayIndex_z;
2479 }
2516 static inline __device__ void rtThrow( unsigned int code )
2517 {
2518  optix::rt_throw( code );
2519 }
2520 
2548 static inline __device__ unsigned int rtGetExceptionCode()
2549 {
2550  return optix::rt_get_exception_code();
2551 }
2552 
2582 static inline __device__ void rtPrintExceptionDetails()
2583 {
2584  const unsigned int code = rtGetExceptionCode();
2585 
2586  if( code == RT_EXCEPTION_STACK_OVERFLOW )
2587  {
2588  rtPrintf( "Caught RT_EXCEPTION_STACK_OVERFLOW\n"
2589  " launch index : %d, %d, %d\n",
2590  rti_internal_register::reg_rayIndex_x,
2591  rti_internal_register::reg_rayIndex_y,
2592  rti_internal_register::reg_rayIndex_z
2593  );
2594  }
2595  else if( code == RT_EXCEPTION_BUFFER_INDEX_OUT_OF_BOUNDS )
2596  {
2597 #if defined(__x86_64) || defined(AMD64) || defined(_M_AMD64)
2598  const unsigned int dim = rti_internal_register::reg_exception_detail0;
2599 
2600  rtPrintf( "Caught RT_EXCEPTION_BUFFER_INDEX_OUT_OF_BOUNDS\n"
2601  " launch index : %d, %d, %d\n"
2602  " buffer address : 0x%llX\n"
2603  " dimensionality : %d\n"
2604  " size : %lldx%lldx%lld\n"
2605  " element size : %d\n"
2606  " accessed index : %lld, %lld, %lld\n",
2607  rti_internal_register::reg_rayIndex_x,
2608  rti_internal_register::reg_rayIndex_y,
2609  rti_internal_register::reg_rayIndex_z,
2610  rti_internal_register::reg_exception_64_detail0,
2611  rti_internal_register::reg_exception_detail0,
2612  rti_internal_register::reg_exception_64_detail1,
2613  dim > 1 ? rti_internal_register::reg_exception_64_detail2 : 1,
2614  dim > 2 ? rti_internal_register::reg_exception_64_detail3 : 1,
2615  rti_internal_register::reg_exception_detail1,
2616  rti_internal_register::reg_exception_64_detail4,
2617  rti_internal_register::reg_exception_64_detail5,
2618  rti_internal_register::reg_exception_64_detail6
2619  );
2620 #else
2621  const unsigned int dim = rti_internal_register::reg_exception_detail1;
2622 
2623  rtPrintf( "Caught RT_EXCEPTION_BUFFER_INDEX_OUT_OF_BOUNDS\n"
2624  " launch index : %d, %d, %d\n"
2625  " buffer address : 0x%X\n"
2626  " dimensionality : %d\n"
2627  " size : %dx%dx%d\n"
2628  " element size : %d\n"
2629  " accessed index : %d, %d, %d\n",
2630  rti_internal_register::reg_rayIndex_x,
2631  rti_internal_register::reg_rayIndex_y,
2632  rti_internal_register::reg_rayIndex_z,
2633  rti_internal_register::reg_exception_detail0,
2634  rti_internal_register::reg_exception_detail1,
2635  rti_internal_register::reg_exception_detail2,
2636  dim > 1 ? rti_internal_register::reg_exception_detail3 : 1,
2637  dim > 2 ? rti_internal_register::reg_exception_detail4 : 1,
2638  rti_internal_register::reg_exception_detail5,
2639  rti_internal_register::reg_exception_detail6,
2640  rti_internal_register::reg_exception_detail7,
2641  rti_internal_register::reg_exception_detail8
2642  );
2643 #endif
2644  }
2645  else if( code == RT_EXCEPTION_PROGRAM_ID_INVALID )
2646  {
2647  rtPrintf( "Caught RT_EXCEPTION_PROGRAM_ID_INVALID\n");
2648  switch(rti_internal_register::reg_exception_detail1)
2649  {
2650  case 0:
2651  rtPrintf( "\tprogram ID equal to RT_PROGRAM_ID_NULL used\n");
2652  break;
2653  case 1:
2654  rtPrintf( "\tprogram ID (%d) is not in the valid range of [1,size)\n", rti_internal_register::reg_exception_detail0);
2655  break;
2656  case 2:
2657  rtPrintf( "\tprogram ID of a deleted program used\n");
2658  break;
2659  }
2660  }
2661  else if( code == RT_EXCEPTION_TEXTURE_ID_INVALID )
2662  {
2663  rtPrintf( "Caught RT_EXCEPTION_TEXTURE_ID_INVALID\n");
2664  switch(rti_internal_register::reg_exception_detail1)
2665  {
2666  case 0:
2667  rtPrintf( "\ttexture ID (%d) is invalid (0)\n", rti_internal_register::reg_exception_detail0);
2668  break;
2669  case 1:
2670  rtPrintf( "\ttexture ID (%d) is not in the valid range of [1,size)\n", rti_internal_register::reg_exception_detail0);
2671  break;
2672  case 2:
2673  rtPrintf( "\ttexture ID (%d) is invalid (-1)\n", rti_internal_register::reg_exception_detail0);
2674  break;
2675  }
2676  }
2677  else if( code == RT_EXCEPTION_BUFFER_ID_INVALID )
2678  {
2679  rtPrintf( "Caught RT_EXCEPTION_BUFFER_ID_INVALID\n");
2680  switch(rti_internal_register::reg_exception_detail1)
2681  {
2682  case 0:
2683  rtPrintf( "\tbuffer ID equal to RT_BUFFER_ID_NULL used\n");
2684  break;
2685  case 1:
2686  rtPrintf( "\tbuffer ID (%d) is not in the valid range of [1,size)\n", rti_internal_register::reg_exception_detail0);
2687  break;
2688  case 2:
2689  rtPrintf( "\tBuffer ID of a deleted buffer used\n");
2690  break;
2691  }
2692  }
2693  else if( code == RT_EXCEPTION_INDEX_OUT_OF_BOUNDS )
2694  {
2695 #if defined(__x86_64) || defined(AMD64) || defined(_M_AMD64)
2696  const unsigned int dim = rti_internal_register::reg_exception_detail0;
2697 
2698  rtPrintf( "Caught RT_EXCEPTION_INDEX_OUT_OF_BOUNDS\n"
2699  " launch index : %d, %d, %d\n"
2700  " buffer address : 0x%llX\n"
2701  " size : %lld\n"
2702  " accessed index : %lld\n",
2703  rti_internal_register::reg_rayIndex_x,
2704  rti_internal_register::reg_rayIndex_y,
2705  rti_internal_register::reg_rayIndex_z,
2706  rti_internal_register::reg_exception_64_detail0,
2707  rti_internal_register::reg_exception_64_detail1,
2708  rti_internal_register::reg_exception_64_detail2
2709  );
2710 #else
2711  const unsigned int dim = rti_internal_register::reg_exception_detail1;
2712 
2713  rtPrintf( "Caught RT_EXCEPTION_INDEX_OUT_OF_BOUNDS\n"
2714  " launch index : %d, %d, %d\n"
2715  " buffer address : 0x%X\n"
2716  " size : %d\n"
2717  " accessed index : %d\n",
2718  rti_internal_register::reg_rayIndex_x,
2719  rti_internal_register::reg_rayIndex_y,
2720  rti_internal_register::reg_rayIndex_z,
2721  rti_internal_register::reg_exception_detail0,
2722  rti_internal_register::reg_exception_detail1,
2723  rti_internal_register::reg_exception_detail2
2724  );
2725 #endif
2726  }
2727  else if( code == RT_EXCEPTION_INVALID_RAY )
2728  {
2729  rtPrintf( "Caught RT_EXCEPTION_INVALID_RAY\n"
2730  " launch index : %d, %d, %d\n"
2731  " ray origin : %f %f %f\n"
2732  " ray direction : %f %f %f\n"
2733  " ray type : %d\n"
2734  " ray tmin : %f\n"
2735  " ray tmax : %f\n",
2736  rti_internal_register::reg_rayIndex_x,
2737  rti_internal_register::reg_rayIndex_y,
2738  rti_internal_register::reg_rayIndex_z,
2739  __int_as_float(rti_internal_register::reg_exception_detail0),
2740  __int_as_float(rti_internal_register::reg_exception_detail1),
2741  __int_as_float(rti_internal_register::reg_exception_detail2),
2742  __int_as_float(rti_internal_register::reg_exception_detail3),
2743  __int_as_float(rti_internal_register::reg_exception_detail4),
2744  __int_as_float(rti_internal_register::reg_exception_detail5),
2745  rti_internal_register::reg_exception_detail6,
2746  __int_as_float(rti_internal_register::reg_exception_detail7),
2747  __int_as_float(rti_internal_register::reg_exception_detail8)
2748  );
2749  }
2750  else if( code == RT_EXCEPTION_INTERNAL_ERROR )
2751  {
2752  // Should never happen.
2753  rtPrintf( "Caught RT_EXCEPTION_INTERNAL_ERROR\n"
2754  " launch index : %d, %d, %d\n"
2755  " error id : %d\n",
2756  rti_internal_register::reg_rayIndex_x,
2757  rti_internal_register::reg_rayIndex_y,
2758  rti_internal_register::reg_rayIndex_z,
2759  rti_internal_register::reg_exception_detail0
2760  );
2761  }
2762  else if( code >= RT_EXCEPTION_USER && code <= 0xFFFF )
2763  {
2764  rtPrintf( "Caught RT_EXCEPTION_USER+%d\n"
2765  " launch index : %d, %d, %d\n",
2766  code-RT_EXCEPTION_USER,
2767  rti_internal_register::reg_rayIndex_x,
2768  rti_internal_register::reg_rayIndex_y,
2769  rti_internal_register::reg_rayIndex_z
2770  );
2771  }
2772  else
2773  {
2774  // Should never happen.
2775  rtPrintf( "Caught unknown exception\n"
2776  " launch index : %d, %d, %d\n",
2777  rti_internal_register::reg_rayIndex_x,
2778  rti_internal_register::reg_rayIndex_y,
2779  rti_internal_register::reg_rayIndex_z
2780  );
2781  }
2782 }
2783 
2784 #endif /* __optix_optix_cuda__internal_h__ */
rtObject
Opaque handle to a OptiX object.
Definition: optix_device.h:179
RT_EXCEPTION_PROGRAM_ID_INVALID
Definition: optix_declarations.h:222
RT_EXCEPTION_STACK_OVERFLOW
Definition: optix_declarations.h:226
optix::VectorTypes
Definition: optix_device.h:53
optix::buffer
Definition: optix_device.h:417
rtPotentialIntersection
static __device__ bool rtPotentialIntersection(float tmin)
Determine whether a computed intersection is potentially valid.
Definition: optix_device.h:1875
RT_EXCEPTION_BUFFER_ID_INVALID
Definition: optix_declarations.h:224
rti_internal_callableprogram::check_is_CPArgVoid
Definition: optix_device.h:1389
rtIntersectChild
static __device__ void rtIntersectChild(unsigned int index)
Visit child of selector.
Definition: optix_device.h:2009
RT_EXCEPTION_INVALID_RAY
Definition: optix_declarations.h:228
RTprogramidnull
RTprogramidnull
Definition: optix_declarations.h:344
optix::bufferId
bufferId is a host version of the device side bufferId.
Definition: optix_device.h:415
rtTransformVector
static __device__ float3 rtTransformVector(RTtransformkind kind, const float3 &v)
Apply the current transformation to a vector.
Definition: optix_device.h:2090
rtTrace
static __device__ void rtTrace(rtObject topNode, optix::Ray ray, T &prd)
Traces a ray.
Definition: optix_device.h:1825
rtPrintExceptionDetails
static __device__ void rtPrintExceptionDetails()
Print information on a caught exception.
Definition: optix_device.h:2582
RT_EXCEPTION_INDEX_OUT_OF_BOUNDS
Definition: optix_declarations.h:225
rtGetTransform
static __device__ void rtGetTransform(RTtransformkind kind, float matrix[16])
Get requested transform.
Definition: optix_device.h:2171
RT_EXCEPTION_INTERNAL_ERROR
Definition: optix_declarations.h:229
optix::buffer::type
Definition: optix_device.h:437
optix_datatypes.h
OptiX public API.
rtPrintf
static __device__ void rtPrintf(const char *fmt)
Prints text to the standard output.
Definition: optix_device.h:2213
rtTransformNormal
static __device__ float3 rtTransformNormal(RTtransformkind kind, const float3 &n)
Apply the current transformation to a normal.
Definition: optix_device.h:2131
RTbufferidnull
RTbufferidnull
Definition: optix_declarations.h:341
RT_INTERNAL_CALLABLE_PROGRAM_DEFS
#define RT_INTERNAL_CALLABLE_PROGRAM_DEFS()
callableProgramId is a host version of the device side callableProgramId.
Definition: optixpp_namespace.h:1770
rtReportIntersection
static __device__ bool rtReportIntersection(unsigned int material)
Report an intersection with the current object and the specified material.
Definition: optix_device.h:1904
rti_internal_callableprogram::callableProgramIdBase
Definition: optix_device.h:1414
rtThrow
static __device__ void rtThrow(unsigned int code)
Throw a user exception.
Definition: optix_device.h:2516
rti_internal_callableprogram::CPArgVoid
Definition: optix_device.h:1381
rti_internal_typeinfo::rti_typeenum
Definition: optix_defines.h:74
optix::callableProgramId
Definition: optix_device.h:1601
rtGetExceptionCode
static __device__ unsigned int rtGetExceptionCode()
Retrieves the type of a caught exception.
Definition: optix_device.h:2548
rtCallableProgramSizeofWrapper
Definition: optix_device.h:333
optix_declarations.h
OptiX public API declarations.
RT_EXCEPTION_USER
Definition: optix_declarations.h:230
RT_EXCEPTION_BUFFER_INDEX_OUT_OF_BOUNDS
Definition: optix_declarations.h:227
RT_EXCEPTION_TEXTURE_ID_INVALID
Definition: optix_declarations.h:223
rti_internal_callableprogram::is_CPArgVoid
Definition: optix_device.h:1383
optix::boundCallableProgramId
Definition: optix_device.h:1671
rtTransformPoint
static __device__ float3 rtTransformPoint(RTtransformkind kind, const float3 &p)
Apply the current transformation to a point.
Definition: optix_device.h:2049
rtTerminateRay
static __device__ void rtTerminateRay()
Terminate traversal associated with the current ray.
Definition: optix_device.h:1969
RTtransformkind
RTtransformkind
Definition: optix_defines.h:45
rtIgnoreIntersection
static __device__ void rtIgnoreIntersection()
Cancels the potential intersection with current ray.
Definition: optix_device.h:1940
Generated on Thu Mar 31 2016 09:25:40 for OptiX by   doxygen 1.8.6