eic-opticks/sysrap/SOPTIX.cu

0001 /**
0002 SOPTIX.cu
0003 ===========
0004
0005
0006 Functions
0007 -----------
0008
0009 trace
0010     populate quad2 prd by call to optixTrace
0011
0012 make_normal_pixel
0013     minimal normal "shader"
0014
0015 render
0016     raygen function : calling trace and "shading" pixels
0017
0018 __raygen__rg
0019     calls one of the above raygen functions depending on params.raygenmode
0020
0021 setPayload
0022     mechanics of communication when not using WITH_PRD
0023
0024 __miss_ms
0025     default quad2 prd OR payload for rays that miss
0026
0027 __closesthit__ch
0028     populate quad2 prd OR payload for rays that intersect
0029
0030 **/
0031
0032 #include <optix.h>
0033 #include <stdint.h>
0034
0035 #include "scuda.h"
0036 #include "squad.h"
0037
0038 #include "SOPTIX_Binding.h"
0039 #include "SOPTIX_Params.h"
0040
0041 #include "scuda_pointer.h"
0042 #include "SOPTIX_getPRD.h"
0043
0044 extern "C" { __constant__ SOPTIX_Params params ;  }
0045
0046 /**
0047 trace : pure function, with no use of params, everything via args
0048 -------------------------------------------------------------------
0049
0050 Outcome of trace is to populate *prd* by payload and attribute passing.
0051 When WITH_PRD macro is defined only 2 32-bit payload values are used to
0052 pass the 64-bit  pointer, otherwise more payload and attributes values
0053 are used to pass the contents IS->CH->RG.
0054
0055 See __closesthit__ch to see where the payload p0-p7 comes from.
0056 **/
0057
0058 static __forceinline__ __device__ void trace(
0059         OptixTraversableHandle handle,
0060         float3                 ray_origin,
0061         float3                 ray_direction,
0062         float                  tmin,
0063         float                  tmax,
0064         quad2*                 prd,
0065         unsigned               visibilityMask
0066         )
0067 {
0068     const float rayTime = 0.0f ;
0069     OptixRayFlags rayFlags = OPTIX_RAY_FLAG_DISABLE_ANYHIT ;   // OPTIX_RAY_FLAG_NONE
0070     const unsigned SBToffset = 0u ;
0071     const unsigned SBTstride = 1u ;
0072     const unsigned missSBTIndex = 0u ;
0073     uint32_t p0, p1 ;
0074     packPointer( prd, p0, p1 );  // scuda_pointer.h : pack prd addr from RG program into two uint32_t passed as payload
0075     optixTrace(
0076             handle,
0077             ray_origin,
0078             ray_direction,
0079             tmin,
0080             tmax,
0081             rayTime,
0082             visibilityMask,
0083             rayFlags,
0084             SBToffset,
0085             SBTstride,
0086             missSBTIndex,
0087             p0, p1
0088             );
0089 }
0090
0091
0092 __forceinline__ __device__ uchar4 make_normal_pixel( const float3& normal, float depth )  // pure
0093 {
0094     return make_uchar4(
0095             static_cast<uint8_t>( clamp( normal.x, 0.0f, 1.0f ) *255.0f ),
0096             static_cast<uint8_t>( clamp( normal.y, 0.0f, 1.0f ) *255.0f ),
0097             static_cast<uint8_t>( clamp( normal.z, 0.0f, 1.0f ) *255.0f ),
0098             static_cast<uint8_t>( clamp( depth   , 0.0f, 1.0f ) *255.0f )
0099             );
0100 }
0101
0102
0103
0104 /**
0105 render : non-pure, uses params for viewpoint inputs and pixels output
0106 -----------------------------------------------------------------------
0107
0108 **/
0109
0110 static __forceinline__ __device__ void render( const uint3& idx, const uint3& dim, quad2* prd )
0111 {
0112     float2 d = 2.0f * make_float2(
0113             static_cast<float>(idx.x)/static_cast<float>(dim.x),
0114             static_cast<float>(idx.y)/static_cast<float>(dim.y)
0115             ) - 1.0f;
0116
0117     //const bool yflip = true ;
0118     //if(yflip) d.y = -d.y ;
0119
0120 #ifdef DBG_PIDX
0121     bool dbg = idx.x == dim.x/2 && idx.y == dim.y/2 ;
0122     if(dbg) printf("//render.DBG_PIDX params.eye (%7.3f %7.3f %7.3f)\n", params.eye.x, params.eye.y, params.eye.z);
0123     if(dbg) printf("//render.DBG_PIDX params.U   (%7.3f %7.3f %7.3f)\n", params.U.x, params.U.y, params.U.z);
0124     if(dbg) printf("//render.DBG_PIDX params.V   (%7.3f %7.3f %7.3f)\n", params.V.x, params.V.y, params.V.z);
0125     if(dbg) printf("//render.DBG_PIDX params.W   (%7.3f %7.3f %7.3f)\n", params.W.x, params.W.y, params.W.z);
0126 #endif
0127
0128     const unsigned cameratype = params.cameratype ;
0129     const float3 dxyUV = d.x * params.U + d.y * params.V ;
0130     const float3 origin    = cameratype == 0u ? params.eye                     : params.eye + dxyUV    ;
0131     const float3 direction = cameratype == 0u ? normalize( dxyUV + params.W )  : normalize( params.W ) ;
0132     //                           cameratype 0u:perspective,                    1u:orthographic
0133
0134     trace(
0135         params.handle,
0136         origin,
0137         direction,
0138         params.tmin,
0139         params.tmax,
0140         prd,
0141         params.vizmask
0142     );
0143
0144     const float3* normal = prd->normal();
0145     float3 diddled_normal = normalize(*normal)*0.5f + 0.5f ;
0146     // "diddling" changes range of elements from -1.f:1.f to 0.f:1.f same as  (n+1.f)/2.f
0147     unsigned index = idx.y * params.width + idx.x ;
0148
0149
0150
0151     float eye_z = -prd->distance()*dot(params.WNORM, direction) ;
0152     const float& A = params.ZPROJ.z ;
0153     const float& B = params.ZPROJ.w ;
0154     float zdepth = cameratype == 0u ? -(A + B/eye_z) : A*eye_z + B  ;  // cf SGLM::zdepth1
0155
0156     if( prd->is_boundary_miss() ) zdepth = 0.999f ;
0157
0158     uchar4 pixel = make_normal_pixel( diddled_normal, zdepth );
0159
0160 #ifdef DBG_PIDX
0161     if(dbg) printf("//render.DBG_PIDX pixel (%d %d %d %d) \n", pixel.x, pixel.y, pixel.z, pixel.w);
0162 #endif
0163
0164     params.pixels[index] = pixel ;
0165 }
0166
0167
0168 extern "C" __global__ void __raygen__rg()
0169 {
0170     const uint3 idx = optixGetLaunchIndex();
0171     const uint3 dim = optixGetLaunchDimensions();
0172
0173 #ifdef DBG_PIDX
0174     bool dbg = idx.x == dim.x/2 && idx.y == dim.y/2 ;
0175     if(dbg)  printf("//__raygen__rg.DBG_PIDX idx(%d,%d,%d) dim(%d,%d,%d)\n", idx.x, idx.y, idx.z, dim.x, dim.y, dim.z );
0176 #endif
0177
0178     quad2 prd ;
0179     prd.zero();
0180
0181     render( idx, dim, &prd );
0182 }
0183
0184 /**
0185 __miss__ms
0186 -------------
0187
0188 * missing "normal" is somewhat render specific and this is used for
0189   all raygenmode but Miss should never happen with real simulations
0190 * Miss can happen with simple geometry testing however when shoot
0191   rays from outside the "world"
0192
0193 **/
0194
0195 extern "C" __global__ void __miss__ms()
0196 {
0197     SOPTIX_MissData* ms = reinterpret_cast<SOPTIX_MissData*>( optixGetSbtDataPointer() );
0198     const unsigned ii_id = 0xffffffffu ;
0199     const unsigned gp_bd = 0xffffffffu ;
0200     const float lposcost = 0.f ;
0201     const float lposfphi = 0.f ;
0202
0203     // printf("//__miss__ms ms.bg_color (%7.3f %7.3f %7.3f) \n", ms->bg_color.x, ms->bg_color.x, ms->bg_color.z );
0204
0205     quad2* prd = SOPTIX_getPRD<quad2>();
0206
0207     prd->q0.f.x = ms->bg_color.x ;  // HMM: thats setting the normal, so it will be diddled
0208     prd->q0.f.y = ms->bg_color.y ;
0209     prd->q0.f.z = ms->bg_color.z ;
0210     prd->q0.f.w = 0.f ;
0211
0212     prd->q1.u.x = 0u ;
0213     prd->q1.u.y = 0u ;
0214     prd->q1.u.z = 0u ;
0215     prd->q1.u.w = 0u ;
0216
0217     prd->set_globalPrimIdx_boundary_(gp_bd);
0218     prd->set_iindex_identity_(ii_id);
0219
0220     prd->set_lpos(lposcost, lposfphi);  // __miss__ms.TRIANGLE
0221
0222
0223 }
0224
0225 /**
0226 __closesthit__ch
0227 =================
0228
0229 optixGetInstanceIndex (aka iindex)
0230     0-based index within IAS
0231
0232 optixGetInstanceId (aka identity)
0233     user supplied instanceId,
0234
0235 optixGetPrimitiveIndex (aka prim_idx)
0236     CustomPrimitiveArray: local index of AABB within the GAS,
0237     TriangleArray: local index of triangle (HMM: within one buildInput?)
0238
0239 optixGetRayTmax
0240     In intersection and CH returns the current smallest reported hitT or the tmax passed into rtTrace
0241     if no hit has been reported
0242
0243
0244 optixGetPrimitiveType
0245     returns OPTIX_PRIMITIVE_TYPE_TRIANGLE or OPTIX_PRIMITIVE_TYPE_CUSTOM
0246
0247
0248 In general will need to branch between::
0249
0250     OPTIX_BUILD_INPUT_TYPE_CUSTOM_PRIMITIVES
0251     OPTIX_BUILD_INPUT_TYPE_TRIANGLES
0252
0253 currently just handles triangles.
0254
0255 **/
0256
0257 extern "C" __global__ void __closesthit__ch()
0258 {
0259     //OptixPrimitiveType type = optixGetPrimitiveType();
0260     //printf("//CH type %u \n", type );  hex(9521) = '0x2531'   OPTIX_PRIMITIVE_TYPE_TRIANGLE
0261
0262     const SOPTIX_HitgroupData* hit_group_data = reinterpret_cast<SOPTIX_HitgroupData*>( optixGetSbtDataPointer() );
0263     const SOPTIX_TriMesh& mesh = hit_group_data->mesh ;
0264
0265     //printf("//__closesthit__ch\n");
0266
0267     const unsigned prim_idx = optixGetPrimitiveIndex();
0268     const float2   barys    = optixGetTriangleBarycentrics();
0269
0270     uint3 tri = mesh.indice[ prim_idx ];
0271     const float3 P0 = mesh.vertex[ tri.x ];
0272     const float3 P1 = mesh.vertex[ tri.y ];
0273     const float3 P2 = mesh.vertex[ tri.z ];
0274
0275     const float3 N0 = mesh.normal[ tri.x ];
0276     const float3 N1 = mesh.normal[ tri.y ];
0277     const float3 N2 = mesh.normal[ tri.z ];
0278
0279     const float3 P = ( 1.0f-barys.x-barys.y)*P0 + barys.x*P1 + barys.y*P2;
0280     const float3 Ng = ( 1.0f-barys.x-barys.y)*N0 + barys.x*N1 + barys.y*N2; // guesss
0281     //const float3 Ng = cross( P1-P0, P2-P0 );
0282
0283     const float3 N = normalize( optixTransformNormalFromObjectToWorldSpace( Ng ) );
0284     // HMM: could get normal by bary-weighting vertex normals ?
0285
0286     unsigned iindex = optixGetInstanceIndex() ;
0287     unsigned identity = optixGetInstanceId() ;
0288     unsigned globalPrimIdx = 0u ;
0289     unsigned boundary = 0u ;
0290     // HMM: need to plant boundary in HitGroupData ?
0291     // cf CSGOptiX/Analytic: node->boundary();// all nodes of tree have same boundary
0292
0293     float t = optixGetRayTmax() ;
0294
0295     // cannot get Object frame ray_origin/direction in CH (only IS,AH)
0296     //const float3 ray_origin = optixGetObjectRayOrigin();
0297     //const float3 ray_direction = optixGetObjectRayDirection();
0298     //const float3 lpos = ray_origin + t*ray_direction  ;
0299     // HMM: could use P to give the local position ?
0300
0301     float lposcost = normalize_cost(P); // scuda.h
0302     float lposfphi = normalize_fphi(P);
0303
0304     quad2* prd = SOPTIX_getPRD<quad2>();
0305
0306     prd->q0.f.x = N.x ;
0307     prd->q0.f.y = N.y ;
0308     prd->q0.f.z = N.z ;
0309     prd->q0.f.w = t ;
0310
0311     prd->set_iindex_identity( iindex, identity ) ;
0312     prd->set_globalPrimIdx_boundary(globalPrimIdx, boundary) ;
0313     prd->set_lpos(lposcost, lposfphi);   // __closesthit__ch.TRIANGLE
0314
0315 }
0316
0317 /**
0318 __intersection__is
0319 ====================
0320
0321 With inbuilt triangles there is no role for IS, the intersection
0322 impl is provided by the Driver.
0323
0324 extern "C" __global__ void __intersection__is()
0325 {
0326 }
0327
0328 **/
0329