eic-opticks/CSGOptiX/CSGOptiX6.cu

0001 /**
0002 CSGOptiX6.cu : incomplete backwards compat for new geometry model
0003 =====================================================================
0004
0005 NOTE THAT MANY OF THE BELOW HEADERS NOT USED AS SIMULATION NOT BROUGHT TO 6
0006 BUT THE UNUSED HEADERS ARE HELPFUL SO THAT COMPILATION ERRORS FROM 6
0007 ARE SIMILAR TO THOSE FROM 7 : FOR EARLY WARNING ON LAPTOP WHICH CANNOT RUN 7
0008
0009 TODO: COMBINE GEOM INTO HERE TO MAKE 6 MORE LIKE 7
0010
0011
0012 **/
0013
0014 #include "scuda.h"
0015 #include "squad.h"
0016
0017 #include "srec.h"
0018 #include "sqat4.h"
0019 #include "sphoton.h"
0020 #include "scerenkov.h"
0021 #include "sstate.h"
0022
0023
0024 #include "qbnd.h"
0025 #include "qcerenkov.h"
0026 #include "qsim.h"
0027
0028
0029
0030
0031 #include <optix_device.h>
0032
0033 rtDeclareVariable(float3,        eye, , );
0034 rtDeclareVariable(float3,        U, , );
0035 rtDeclareVariable(float3,        V, , );
0036 rtDeclareVariable(float3,        W, , );
0037 rtDeclareVariable(float,         tmin, , );
0038 rtDeclareVariable(unsigned,      radiance_ray_type, , );
0039 rtDeclareVariable(unsigned,      cameratype, , );
0040 rtDeclareVariable(unsigned,      raygenmode, , );
0041
0042 rtDeclareVariable(uint2, launch_index, rtLaunchIndex, );
0043 rtDeclareVariable(uint2, launch_dim,   rtLaunchDim, );
0044
0045 rtDeclareVariable(optix::Ray, ray, rtCurrentRay, );
0046 rtDeclareVariable(float,      t, rtIntersectionDistance, );
0047
0048
0049 // NB: cannot replace these "real" optix buffer as only input buffers can shims over CUDA buffers like geometry buffers
0050
0051 rtBuffer<uchar4, 2>   pixel_buffer;   // formerly pixels_buffer
0052 rtBuffer<float4, 2>   isect_buffer;   // formerly posi_buffer
0053 rtBuffer<quad4,  2>   photon_buffer;  // formerly isect_buffer
0054
0055
0056 static __device__ __inline__ uchar4 make_color(const float3& c)
0057 {
0058     return make_uchar4( static_cast<unsigned char>(__saturatef(c.x)*255.99f),
0059                         static_cast<unsigned char>(__saturatef(c.y)*255.99f),
0060                         static_cast<unsigned char>(__saturatef(c.z)*255.99f),
0061                         255u);
0062 }
0063
0064
0065 struct PerRayData
0066 {
0067     float3   result;
0068     int      mode ;
0069     float4   isect ;
0070 };
0071
0072 rtDeclareVariable(float3, position,         attribute position, );
0073 rtDeclareVariable(float3, shading_normal,   attribute shading_normal, );
0074 rtDeclareVariable(unsigned,  intersect_identity,   attribute intersect_identity, );
0075
0076 rtDeclareVariable(PerRayData, prd, rtPayload, );
0077
0078 rtDeclareVariable(rtObject,      top_object, , );
0079
0080
0081 RT_PROGRAM void raygen()
0082 {
0083     PerRayData prd;
0084     prd.result = make_float3( 1.f, 0.f, 0.f ) ;
0085     prd.isect  = make_float4( 0.f, 0.f, 0.f, 0.f );
0086     prd.mode = 0 ;
0087
0088     const float2 d = make_float2(launch_index) / make_float2(launch_dim) * 2.f - 1.f ;
0089     const float3 dxyUV = d.x*U + d.y*V ;
0090     //                       cameratype     0u perspective           1u orthographic
0091     const float3 origin    = cameratype == 0u ? eye                    : eye + dxyUV    ;
0092     const float3 direction = cameratype == 0u ? normalize( dxyUV + W ) : normalize( W ) ;
0093
0094     optix::Ray ray = optix::make_Ray( origin, direction, radiance_ray_type, tmin, RT_DEFAULT_MAX) ;
0095     rtTrace(top_object, ray, prd);
0096
0097     const bool yflip = true ;
0098     uint2 index = make_uint2( launch_index.x , yflip ? launch_dim.y - 1u - launch_index.y : launch_index.y );
0099     pixel_buffer[index] = make_color( prd.result ) ;
0100     isect_buffer[index] = prd.isect ;
0101
0102     quad4 photon ;
0103     photon.q0.f.x = prd.result.x ;
0104     photon.q0.f.y = prd.result.y ;
0105     photon.q0.f.z = prd.result.z ;
0106     photon.q0.f.w = 0.f ;
0107
0108     photon.q1.f.x = prd.isect.x ;
0109     photon.q1.f.y = prd.isect.y ;
0110     photon.q1.f.z = prd.isect.z ;
0111     photon.q1.f.w = prd.isect.w ;
0112
0113     photon.q2.f.x = origin.x ;
0114     photon.q2.f.y = origin.y ;
0115     photon.q2.f.z = origin.z ;
0116     photon.q2.f.w = tmin ;
0117
0118     photon.q3.f.x = direction.x ;
0119     photon.q3.f.y = direction.y ;
0120     photon.q3.f.z = direction.z ;
0121     photon.q3.i.w = prd.mode ;
0122
0123     photon_buffer[index] = photon ;
0124
0125 #ifdef DEBUG_SIX
0126     //rtPrintf("//DEBUG_SIX/OptiXTest.cu:raygen prd.mode %d \n", prd.mode );
0127 #endif
0128
0129 }
0130
0131 RT_PROGRAM void miss()
0132 {
0133     //prd.result = make_float3(0.5f, 1.f, 1.f) ;  // cyan
0134     prd.result = make_float3(1.f, 1.f, 1.f) ;
0135     prd.mode = 1 ;
0136 }
0137
0138 RT_PROGRAM void closest_hit()
0139 {
0140     prd.result = normalize(rtTransformNormal(RT_OBJECT_TO_WORLD, shading_normal))*0.5f + 0.5f;
0141     float3 isect = ray.origin + t*ray.direction ;
0142     prd.isect = make_float4( isect, __uint_as_float(intersect_identity) );
0143     prd.mode = 2 ;
0144 }
0145