89 #define TACHYON_INTERNAL 1 98 #define TACHYON_USE_RAY_STEP 1 99 #define TACHYON_TRANS_USE_INCIDENT 1 100 #define TACHYON_RAY_STEP N*rtLaunch.scene.epsilon*4.0f 101 #define TACHYON_RAY_STEP2 ray_direction*rtLaunch.scene.epsilon*4.0f 104 #define REVERSE_RAY_STEP (scene_epsilon*10.0f) 105 #define REVERSE_RAY_LENGTH 3.0f 109 #define TACHYON_USE_SPHERES_HEARNBAKER 1 115 const uint3 launch_index = optixGetLaunchIndex();
116 const uint3 launch_dim = optixGetLaunchDimensions();
117 const int idx = launch_index.y*launch_dim.x + launch_index.x;
124 return index.y*dim.x + index.x;
136 static __forceinline__ __device__
138 const uint64_t uptr =
static_cast<uint64_t
>( i0 ) << 32 | i1;
139 void* ptr =
reinterpret_cast<void*
>( uptr );
143 static __forceinline__ __device__
145 const uint64_t uptr =
reinterpret_cast<uint64_t
>( ptr );
147 i1 = uptr & 0x00000000ffffffff;
151 static __forceinline__ __device__ T *
getPRD() {
152 const uint32_t p0 = optixGetPayload_0();
153 const uint32_t p1 = optixGetPayload_1();
174 return optixGetPayload_2();
182 struct PerRayData_shadow {
191 return __int_as_float(optixGetPayload_0());
197 optixSetPayload_0(__float_as_int(attenuation));
215 const float3 &cam_pos,
216 float fade_start,
float fade_end,
218 float camdist =
length(hit_point - cam_pos);
222 float fade_len = fade_start - fade_end;
223 alpha *= __saturatef((camdist - fade_start) / fade_len);
229 float3 ray_direction, float3 center,
230 float rad, float2 &tinterval) {
231 float3 V = center - ray_origin;
232 float b =
dot(V, ray_direction);
233 float disc = b*b + rad*rad -
dot(V, V);
238 tinterval.x = b-disc;
239 tinterval.y = b+disc;
248 float3 ray_direction,
249 float &tmin,
float &tmax,
250 const float4 plane) {
252 float dt =
dot(ray_direction, n);
253 float t = (-plane.w -
dot(n, ray_origin))/dt;
254 if(t > tmin && t < tmax) {
262 float3 p = ray_origin + tmin * ray_direction;
280 const int code = optixGetExceptionCode();
281 const uint3 launch_index = optixGetLaunchIndex();
284 case OPTIX_EXCEPTION_CODE_STACK_OVERFLOW:
285 printf(
"TachyonOptiX) Stack overflow, launch idx (%u,%u)\n",
286 launch_index.x, launch_index.y);
289 case OPTIX_EXCEPTION_CODE_TRACE_DEPTH_EXCEEDED:
290 printf(
"TachyonOptiX) Max trace depth exceeded, launch idx (%u,%u)\n",
291 launch_index.x, launch_index.y);
294 case OPTIX_EXCEPTION_CODE_TRAVERSAL_DEPTH_EXCEEDED:
295 printf(
"TachyonOptiX) Max traversal depth exceeded, launch idx (%u,%u)\n",
296 launch_index.x, launch_index.y);
299 case OPTIX_EXCEPTION_CODE_TRAVERSAL_INVALID_MISS_SBT:
300 printf(
"TachyonOptiX) Invalid miss SBT record idx, launch idx (%u,%u)\n",
301 launch_index.x, launch_index.y);
305 case OPTIX_EXCEPTION_CODE_TRAVERSAL_INVALID_HIT_SBT:
306 printf(
"TachyonOptiX) Invalid hit SBT record idx, launch idx (%u,%u)\n",
307 launch_index.x, launch_index.y);
311 #if OPTIX_VERSION >= 70100 312 case OPTIX_EXCEPTION_CODE_BUILTIN_IS_MISMATCH:
313 printf(
"TachyonOptiX) Built-in IS mismatch, launch idx (%u,%u)\n",
314 launch_index.x, launch_index.y);
317 case OPTIX_EXCEPTION_CODE_INVALID_RAY:
318 printf(
"TachyonOptiX) Trace call contains Inf/NaN, launch idx (%u,%u):\n" 319 "TachyonOptiX) @ %s\n",
320 launch_index.x, launch_index.y, optixGetExceptionLineInfo());
324 case OPTIX_EXCEPTION_CODE_CALLABLE_PARAMETER_MISMATCH:
325 printf(
"TachyonOptiX) Callable param mismatch, launch idx (%d,%d)\n",
326 launch_index.x, launch_index.y);
331 case OPTIX_EXCEPTION_CODE_TRAVERSAL_INVALID_TRAVERSABLE:
333 printf(
"TachyonOptiX) Caught exception 0x%X (%d) at launch idx (%u,%u)\n",
334 code, code, launch_index.x, launch_index.y );
401 #if defined(TACHYON_USE_GEOMFLAGS) 402 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
417 if (tmesh.tex2d !=
nullptr) {
418 const int primID = optixGetPrimitiveIndex();
421 if (tmesh.indices == NULL) {
423 index = make_int3(idx3, idx3+1, idx3+2);
425 index = tmesh.indices[primID];
428 const float2 barycentrics = optixGetTriangleBarycentrics();
430 float2 txc0 = tmesh.tex2d[index.x];
431 float2 txc1 = tmesh.tex2d[index.y];
432 float2 txc2 = tmesh.tex2d[index.z];
435 float2 texcoord = (txc0 * (1.0f - barycentrics.x - barycentrics.y) +
436 txc1 * barycentrics.x + txc2 * barycentrics.y);
441 float4 tx = tex2D<float4>(mat.tex, texcoord.x, texcoord.y);
450 attenuation *= (1.0f - opacity);
453 if (attenuation < 0.001f) {
456 #if defined(TACHYON_RAYSTATS) 460 optixIgnoreIntersection();
467 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
471 const uint3 launch_index = optixGetLaunchIndex();
472 if (launch_index.x == 994) {
473 printf(
"AH xy:%d %d mat[%d] diffuse: %g opacity: %g atten: %g\n",
474 launch_index.x, launch_index.y,
482 attenuation *= (1.0f - opacity);
485 if (attenuation < 0.001f) {
488 #if defined(TACHYON_RAYSTATS) 492 optixIgnoreIntersection();
503 const float3 ray_origin = optixGetWorldRayOrigin();
504 const float3 ray_direction = optixGetWorldRayDirection();
505 const float t_hit = optixGetRayTmax();
506 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
510 float3 hit_point = ray_origin + t_hit * ray_direction;
513 float clipalpha = 1.0f;
523 attenuation *= (1.0f - (clipalpha * opacity));
526 if (attenuation < 0.001f) {
529 #if defined(TACHYON_RAYSTATS) 533 optixIgnoreIntersection();
560 #if defined(TACHYON_RAYSTATS) 592 val = __saturatef(val);
598 #
if defined(TACHYON_RAYSTATS)
629 val = __saturatef(val);
635 #
if defined(TACHYON_RAYSTATS)
659 #if defined(TACHYON_OPTIXDENOISER) 670 colrgba4f = clamp_float4(colrgba4f);
694 tonedcol = tonemap_color(colrgba4f,
702 colrgba4f = tonedcol;
710 #if defined(TACHYON_RAYSTATS) 711 static void __inline__ __device__ raystats_clear_incr(
unsigned int idx) {
736 static __device__ __inline__
737 void dof_ray(
const float cam_dof_focal_dist,
const float cam_dof_aperture_rad,
738 const float3 &ray_origin_orig, float3 &ray_origin,
739 const float3 &ray_direction_orig, float3 &ray_direction,
740 unsigned int &randseed,
const float3 &up,
const float3 &right) {
741 float3 focuspoint = ray_origin_orig + ray_direction_orig * cam_dof_focal_dist;
744 ray_origin = ray_origin_orig + dofjxy.x*right + dofjxy.y*up;
745 ray_direction =
normalize(focuspoint - ray_origin);
751 static __device__ __inline__
752 void dof_ray(
const float cam_dof_focal_dist,
const float cam_dof_aperture_rad,
753 const float3 &ray_origin_orig, float3 &ray_origin,
754 const float3 &ray_direction_orig, float3 &ray_direction,
755 float2 &qrnxy,
const float3 &up,
const float3 &right) {
756 float3 focuspoint = ray_origin_orig + ray_direction_orig * cam_dof_focal_dist;
759 ray_origin = ray_origin_orig + dofjxy.x*right + dofjxy.y*up;
760 ray_direction =
normalize(focuspoint - ray_origin);
769 template<
int STEREO_ON,
int DOF_ON>
770 static __device__ __inline__
772 #if defined(TACHYON_TIME_COLORING) 773 clock_t t0 = clock();
776 const uint3 launch_dim = optixGetLaunchDimensions();
777 const uint3 launch_index = optixGetLaunchIndex();
779 #if defined(TACHYON_RAYSTATS) 781 raystats_clear_incr(idx);
792 uint viewport_sz_y, viewport_idx_y;
795 viewport_sz_y = launch_dim.y >> 1;
796 if (launch_index.y >= viewport_sz_y) {
798 viewport_idx_y = launch_index.y - viewport_sz_y;
799 eyepos = cam.
pos + cam.U * cam.stereo_eyesep * 0.5f;
802 viewport_idx_y = launch_index.y;
803 eyepos = cam.pos - cam.U * cam.stereo_eyesep * 0.5f;
807 viewport_sz_y = launch_dim.y;
808 viewport_idx_y = launch_index.y;
816 float2 aspect = make_float2(
float(launch_dim.x) /
float(viewport_sz_y), 1.0f) * cam.zoom;
817 float2 viewportscale = 1.0f / make_float2(launch_dim.x, viewport_sz_y);
818 float2 d = make_float2(launch_index.x, viewport_idx_y) * viewportscale * aspect * 2.f - aspect;
824 float3 ray_origin = eyepos;
829 jxy = jxy * viewportscale * aspect * 2.f + d;
830 float3 ray_direction =
normalize(jxy.x*cam.U + jxy.y*cam.V + cam.W);
834 dof_ray(cam.dof_focal_dist, cam.dof_aperture_rad,
835 eyepos, ray_origin, ray_direction, ray_direction,
836 randseed, cam.V, cam.U);
856 OptixVisibilityMask( 255 ),
857 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
868 #if defined(TACHYON_TIME_COLORING) 869 accumulate_time_coloring(col, t0);
876 tachyon_camera_perspective_general<0, 0>();
880 tachyon_camera_perspective_general<0, 1>();
884 tachyon_camera_perspective_general<1, 0>();
888 tachyon_camera_perspective_general<1, 1>();
897 template<
int STEREO_ON,
int DOF_ON>
898 static __device__ __inline__
900 #if defined(TACHYON_TIME_COLORING) 901 clock_t t0 = clock();
904 const uint3 launch_dim = optixGetLaunchDimensions();
905 const uint3 launch_index = optixGetLaunchIndex();
907 #if defined(TACHYON_RAYSTATS) 909 raystats_clear_incr(idx);
920 uint viewport_sz_y, viewport_idx_y;
921 float3 view_direction;
924 viewport_sz_y = launch_dim.y >> 1;
925 if (launch_index.y >= viewport_sz_y) {
927 viewport_idx_y = launch_index.y - viewport_sz_y;
928 eyepos = cam.
pos + cam.U * cam.stereo_eyesep * 0.5f;
931 viewport_idx_y = launch_index.y;
932 eyepos = cam.pos - cam.U * cam.stereo_eyesep * 0.5f;
934 view_direction =
normalize(cam.pos-eyepos +
normalize(cam.W) * cam.stereo_convergence_dist);
937 viewport_sz_y = launch_dim.y;
938 viewport_idx_y = launch_index.y;
946 float2 aspect = make_float2(
float(launch_dim.x) /
float(viewport_sz_y), 1.0f) * cam.zoom;
947 float2 viewportscale = 1.0f / make_float2(launch_dim.x, viewport_sz_y);
949 float2 d = make_float2(launch_index.x, viewport_idx_y) * viewportscale * aspect * 2.f - aspect;
955 float3 ray_direction = view_direction;
959 jxy = jxy * viewportscale * aspect * 2.f + d;
960 float3 ray_origin = eyepos + jxy.x*cam.U + jxy.y*cam.V;
964 dof_ray(cam.dof_focal_dist, cam.dof_aperture_rad,
965 ray_origin, ray_origin, view_direction, ray_direction,
966 randseed, cam.V, cam.U);
985 OptixVisibilityMask( 255 ),
986 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
997 #if defined(TACHYON_TIME_COLORING) 998 accumulate_time_coloring(col, t0);
1005 tachyon_camera_orthographic_general<0, 0>();
1009 tachyon_camera_orthographic_general<0, 1>();
1013 tachyon_camera_orthographic_general<1, 0>();
1017 tachyon_camera_orthographic_general<1, 1>();
1044 template<
int STEREO_ON,
int DOF_ON>
1045 static __device__ __inline__
1047 #if defined(TACHYON_TIME_COLORING) 1048 clock_t t0 = clock();
1051 const uint3 launch_dim = optixGetLaunchDimensions();
1052 const uint3 launch_index = optixGetLaunchIndex();
1054 #if defined(TACHYON_RAYSTATS) 1056 raystats_clear_incr(idx);
1062 uint facesz = launch_dim.y;
1063 uint face = (launch_index.x / facesz) % 6;
1064 uint2 face_idx = make_uint2(launch_index.x % facesz, launch_index.y);
1075 float3 face_U, face_V, face_W;
1123 viewport_sz_x = launch_dim.x >> 1;
1124 if (launch_index.x >= viewport_sz_x) {
1127 eyeshift = 0.5f * cam.stereo_eyesep;
1131 eyeshift = -0.5f * cam.stereo_eyesep;
1135 viewport_sz_x = launch_dim.x;
1143 float facescale = 1.0f / facesz;
1144 float2 d = make_float2(face_idx.x, face_idx.y) * facescale * 2.f - 1.0f;
1153 jxy = jxy * facescale * 2.f + d;
1154 float3 ray_direction =
normalize(jxy.x*face_U + jxy.y*face_V + face_W);
1156 float3 ray_origin = cam.pos;
1158 ray_origin += eyeshift *
cross(ray_direction, cam.V);
1163 dof_ray(cam.dof_focal_dist, cam.dof_aperture_rad,
1164 ray_origin, ray_origin, ray_direction, ray_direction,
1165 randseed, face_V, face_U);
1184 OptixVisibilityMask( 255 ),
1185 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
1196 #if defined(TACHYON_TIME_COLORING) 1197 accumulate_time_coloring(col, t0);
1205 tachyon_camera_cubemap_general<0, 0>();
1209 tachyon_camera_cubemap_general<0, 1>();
1213 tachyon_camera_cubemap_general<1, 0>();
1217 tachyon_camera_cubemap_general<1, 1>();
1233 template<
int STEREO_ON,
int DOF_ON>
1234 static __device__ __inline__
1236 #if defined(TACHYON_TIME_COLORING) 1237 clock_t t0 = clock();
1240 const uint3 launch_dim = optixGetLaunchDimensions();
1241 const uint3 launch_index = optixGetLaunchIndex();
1243 #if defined(TACHYON_RAYSTATS) 1245 raystats_clear_incr(idx);
1255 uint viewport_sz_y, viewport_idx_y;
1259 viewport_sz_y = launch_dim.y >> 1;
1260 if (launch_index.y >= viewport_sz_y) {
1262 viewport_idx_y = launch_index.y - viewport_sz_y;
1266 viewport_idx_y = launch_index.y;
1267 eyeshift = 0.5f * cam.stereo_eyesep;
1271 viewport_sz_y = launch_dim.y;
1272 viewport_idx_y = launch_index.y;
1280 float thetamax = 0.5 * fov;
1288 float2 viewport_sz = make_float2(launch_dim.x, viewport_sz_y);
1289 float2 radperpix = fov / viewport_sz;
1290 float2 viewport_mid = viewport_sz * 0.5f;
1300 float2 viewport_idx = make_float2(launch_index.x, viewport_idx_y) + jxy;
1303 float2 p = (viewport_idx - viewport_mid) * radperpix;
1304 float theta = hypotf(p.x, p.y);
1308 if (theta < thetamax) {
1309 float3 ray_direction;
1310 float3 ray_origin = cam.pos;
1315 ray_direction = cam.W;
1317 float sintheta, costheta;
1318 sincosf(theta, &sintheta, &costheta);
1319 float rsin = sintheta / theta;
1320 ray_direction = cam.U*rsin*p.x + cam.V*rsin*p.y + cam.W*costheta;
1324 ray_origin += eyeshift *
cross(ray_direction, cam.W);
1328 float rcos = costheta / theta;
1329 float3 ray_up = -cam.U*rcos*p.x -cam.V*rcos*p.y + cam.W*sintheta;
1330 float3 ray_right = cam.U*(p.y/theta) + cam.V*(-p.x/theta);
1331 dof_ray(cam.dof_focal_dist, cam.dof_aperture_rad,
1332 ray_origin, ray_origin, ray_direction, ray_direction,
1333 randseed, ray_up, ray_right);
1353 OptixVisibilityMask( 255 ),
1354 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
1366 #if defined(TACHYON_TIME_COLORING) 1367 accumulate_time_coloring(col, t0);
1375 tachyon_camera_dome_general<0, 0>();
1379 tachyon_camera_dome_general<0, 1>();
1383 tachyon_camera_dome_general<1, 0>();
1387 tachyon_camera_dome_general<1, 1>();
1409 template<
int STEREO_ON,
int DOF_ON>
1410 static __device__ __inline__
1412 #if defined(TACHYON_TIME_COLORING) 1413 clock_t t0 = clock();
1416 const uint3 launch_dim = optixGetLaunchDimensions();
1417 const uint3 launch_index = optixGetLaunchIndex();
1419 #if defined(TACHYON_RAYSTATS) 1421 raystats_clear_incr(idx);
1433 uint viewport_sz_y, viewport_idx_y;
1437 viewport_sz_y = launch_dim.y >> 1;
1438 if (launch_index.y >= viewport_sz_y) {
1440 viewport_idx_y = launch_index.y - viewport_sz_y;
1444 viewport_idx_y = launch_index.y;
1445 eyeshift = 0.5f * cam.stereo_eyesep;
1449 viewport_sz_y = launch_dim.y;
1450 viewport_idx_y = launch_index.y;
1454 float2 viewport_sz = make_float2(launch_dim.x, viewport_sz_y);
1455 float2 radperpix =
M_PIf / viewport_sz * make_float2(2.0f, 1.0f);
1456 float2 viewport_mid = viewport_sz * 0.5f;
1466 float2 viewport_idx = make_float2(launch_index.x, viewport_idx_y) + jxy;
1467 float2 rangle = (viewport_idx - viewport_mid) * radperpix;
1469 float sin_ax, cos_ax, sin_ay, cos_ay;
1470 sincosf(rangle.x, &sin_ax, &cos_ax);
1471 sincosf(rangle.y, &sin_ay, &cos_ay);
1473 float3 ray_direction =
normalize(cos_ay * (cos_ax * cam.W + sin_ax * cam.U) + sin_ay * cam.V);
1475 float3 ray_origin = cam.pos;
1477 ray_origin += eyeshift *
cross(ray_direction, cam.V);
1482 float3 ray_right =
normalize(cos_ay * (-sin_ax * cam.W - cos_ax * cam.U) + sin_ay * cam.V);
1483 float3 ray_up =
cross(ray_direction, ray_right);
1484 dof_ray(cam.dof_focal_dist, cam.dof_aperture_rad,
1485 ray_origin, ray_origin, ray_direction, ray_direction,
1486 randseed, ray_up, ray_right);
1505 OptixVisibilityMask( 255 ),
1506 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
1517 #if defined(TACHYON_TIME_COLORING) 1518 accumulate_time_coloring(col, t0);
1525 tachyon_camera_equirectangular_general<0, 0>();
1529 tachyon_camera_equirectangular_general<0, 1>();
1533 tachyon_camera_equirectangular_general<1, 0>();
1537 tachyon_camera_equirectangular_general<1, 1>();
1549 template<
int STEREO_ON,
int DOF_ON>
1550 static __device__ __inline__
1552 #if defined(TACHYON_TIME_COLORING) 1553 clock_t t0 = clock();
1556 const uint3 launch_dim = optixGetLaunchDimensions();
1557 const uint3 launch_index = optixGetLaunchIndex();
1559 #if defined(TACHYON_RAYSTATS) 1561 raystats_clear_incr(idx);
1573 uint viewport_sz_y, viewport_idx_y;
1577 viewport_sz_y = launch_dim.y >> 1;
1578 if (launch_index.y >= viewport_sz_y) {
1580 viewport_idx_y = launch_index.y - viewport_sz_y;
1584 viewport_idx_y = launch_index.y;
1585 eyeshift = 0.5f * cam.stereo_eyesep;
1589 viewport_sz_y = launch_dim.y;
1590 viewport_idx_y = launch_index.y;
1594 float2 viewport_sz = make_float2(launch_dim.x, viewport_sz_y);
1595 float2 viewport_sz_inv = make_float2(1.0f / launch_dim.x,
1596 1.0f / viewport_sz_y);
1606 float2 viewport_idx = make_float2(launch_index.x, viewport_idx_y) + jxy;
1608 float2 px = viewport_idx * viewport_sz_inv;
1609 px = (px - make_float2(0.5f, 0.5f)) * 2.0f;
1612 float3 ray_direction = OctDecode<1>(px);
1614 float3 ray_origin = cam.pos;
1616 ray_origin += eyeshift *
cross(ray_direction, cam.V);
1622 float3 ray_up =
cross(ray_direction, ray_right);
1623 dof_ray(cam.dof_focal_dist, cam.dof_aperture_rad,
1624 ray_origin, ray_origin, ray_direction, ray_direction,
1625 randseed, ray_up, ray_right);
1644 OptixVisibilityMask( 255 ),
1645 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
1656 #if defined(TACHYON_TIME_COLORING) 1657 accumulate_time_coloring(col, t0);
1664 tachyon_camera_octahedral_general<0, 0>();
1668 tachyon_camera_octahedral_general<0, 1>();
1672 tachyon_camera_octahedral_general<1, 0>();
1676 tachyon_camera_octahedral_general<1, 1>();
1684 template<
int STEREO_ON,
int DOF_ON>
1685 static __device__ __inline__
1687 #if defined(TACHYON_TIME_COLORING) 1688 clock_t t0 = clock();
1691 const uint3 launch_dim = optixGetLaunchDimensions();
1692 const uint3 launch_index = optixGetLaunchIndex();
1694 #if defined(TACHYON_RAYSTATS) 1696 raystats_clear_incr(idx);
1706 uint viewport_sz_x, viewport_idx_x;
1710 viewport_sz_x = launch_dim.x >> 1;
1711 if (launch_index.x >= viewport_sz_x) {
1713 viewport_idx_x = launch_index.x - viewport_sz_x;
1717 viewport_idx_x = launch_index.x;
1718 eyeshift = -0.5f * cam.stereo_eyesep;
1722 viewport_sz_x = launch_dim.x;
1723 viewport_idx_x = launch_index.x;
1730 float2 aspect = make_float2(
float(viewport_sz_x) /
float(launch_dim.y), 1.0f) * cam.zoom;
1731 float2 viewportscale = 1.0f / make_float2(viewport_sz_x, launch_dim.y);
1732 float2 d = make_float2(viewport_idx_x, launch_index.y) * viewportscale * aspect * 2.f - aspect;
1751 float2 cp = make_float2(viewport_sz_x >> 1, launch_dim.y >> 1) * viewportscale * aspect * 2.f - aspect;;
1753 float r2 = dr.x*dr.x + dr.y*dr.y;
1754 float r = 0.24f*r2*r2 + 0.22f*r2 + 1.0f;
1758 unsigned int randseed = tea<4>(idx, subframecount);
1760 float3 eyepos = cam.pos;
1762 eyepos += eyeshift * cam.U;
1765 float3 ray_origin = eyepos;
1774 jxy *= (subframecount > 0 || s > 0);
1776 jxy = jxy * viewportscale * aspect * 2.f + d;
1777 float3 ray_direction =
normalize(jxy.x*cam.U + jxy.y*cam.V + cam.W);
1781 dof_ray(cam.dof_focal_dist, cam.dof_aperture_rad,
1782 eyepos, ray_origin, ray_direction, ray_direction,
1783 randseed, cam.V, cam.U);
1802 OptixVisibilityMask( 255 ),
1803 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
1814 #if defined(TACHYON_TIME_COLORING) 1815 accumulate_time_coloring(col, t0);
1822 tachyon_camera_oculus_rift_general<0, 0>();
1826 tachyon_camera_oculus_rift_general<0, 1>();
1830 tachyon_camera_oculus_rift_general<1, 0>();
1834 tachyon_camera_oculus_rift_general<1, 1>();
1848 printf(
"This should never happen!\n");
1860 static __inline__ __device__
1862 float3 world_shading_normal =
normalize(optixTransformNormalFromObjectToWorldSpace(Nshading));
1863 float3 world_geometric_normal =
normalize(optixTransformNormalFromObjectToWorldSpace(Ngeometric));
1864 const float3 ray_dir = optixGetWorldRayDirection();
1865 return faceforward(world_shading_normal, -ray_dir, world_geometric_normal);
1879 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
1880 const float3 ray_origin = optixGetObjectRayOrigin();
1881 const float3 obj_ray_direction = optixGetObjectRayDirection();
1882 const int primID = optixGetPrimitiveIndex();
1884 float3 base = sbtHG.
cone.
base[primID];
1885 float3 apex = sbtHG.
cone.
apex[primID];
1889 float3 axis = (apex - base);
1890 float3 obase = ray_origin - base;
1891 float3 oapex = ray_origin - apex;
1892 float m0 =
dot(axis, axis);
1893 float m1 =
dot(obase, axis);
1894 float m2 =
dot(obj_ray_direction, axis);
1895 float m3 =
dot(obj_ray_direction, obase);
1896 float m5 =
dot(obase, obase);
1897 float m9 =
dot(oapex, axis);
1901 float rr = baserad - apexrad;
1902 float hy = m0 + rr*rr;
1903 float k2 = m0*m0 - m2*m2*hy;
1904 float k1 = m0*m0*m3 - m1*m2*hy + m0*baserad*(rr*m2*1.0f );
1905 float k0 = m0*m0*m5 - m1*m1*hy + m0*baserad*(rr*m1*2.0f - m0*baserad);
1906 float h = k1*k1 - k2*k0;
1910 float t = (-k1-sqrt(h))/k2;
1911 float y = m1 + t*m2;
1912 if (y < 0.0f || y > m0)
1919 static __host__ __device__ __inline__
1921 const float3 ray_origin = optixGetWorldRayOrigin();
1922 const float3 ray_direction = optixGetWorldRayDirection();
1923 const float t_hit = optixGetRayTmax();
1924 const int primID = optixGetPrimitiveIndex();
1928 float3 base = sbtHG.
cone.
base[primID];
1929 float3 apex = sbtHG.
cone.
apex[primID];
1933 float3 axis = (apex - base);
1934 float3 obase = ray_origin - base;
1935 float3 oapex = ray_origin - apex;
1936 float m0 =
dot(axis, axis);
1937 float m1 =
dot(obase, axis);
1938 float m2 =
dot(ray_direction, axis);
1939 float m3 =
dot(ray_direction, obase);
1940 float m5 =
dot(obase, obase);
1941 float m9 =
dot(oapex, axis);
1945 float rr = baserad - apexrad;
1946 float hy = m0 + rr*rr;
1947 float k2 = m0*m0 - m2*m2*hy;
1948 float k1 = m0*m0*m3 - m1*m2*hy + m0*baserad*(rr*m2*1.0f );
1949 float k0 = m0*m0*m5 - m1*m1*hy + m0*baserad*(rr*m1*2.0f - m0*baserad);
1950 float h = k1*k1 - k2*k0;
1954 float t = (-k1-sqrt(h))/k2;
1955 float y = m1 + t*m2;
1959 float3 hit = t * ray_direction;
1960 float3 Ng =
normalize(m0*(m0*(obase + hit) + rr*axis*baserad) - axis*hy*y);
1972 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
1973 const float3 ray_origin = optixGetObjectRayOrigin();
1974 const float3 obj_ray_direction = optixGetObjectRayDirection();
1975 const int primID = optixGetPrimitiveIndex();
1977 float3 start = sbtHG.
cyl.
start[primID];
1978 float3 end = sbtHG.
cyl.
end[primID];
1979 float radius = sbtHG.
cyl.
radius[primID];
1981 float3 axis = (end - start);
1982 float3 rc = ray_origin - start;
1983 float3 n =
cross(obj_ray_direction, axis);
1984 float lnsq =
dot(n, n);
1990 float invln = rsqrtf(lnsq);
1996 float3 O =
cross(rc, axis);
1997 float t = -
dot(O, n) * invln;
2000 float s =
dot(obj_ray_direction, O);
2001 s = fabs(sqrtf(radius*radius - d*d) / s);
2002 float axlen =
length(axis);
2007 float3 hit = ray_origin + obj_ray_direction * tin;
2008 float3 tmp2 = hit - start;
2009 float tmp =
dot(tmp2, axis_u);
2010 if ((tmp > 0.0f) && (tmp < axlen)) {
2016 hit = ray_origin + obj_ray_direction * tout;
2018 tmp =
dot(tmp2, axis_u);
2019 if ((tmp > 0.0f) && (tmp < axlen)) {
2026 static __host__ __device__ __inline__
2028 const float3 ray_origin = optixGetWorldRayOrigin();
2029 const float3 ray_direction = optixGetWorldRayDirection();
2030 const float t_hit = optixGetRayTmax();
2031 const int primID = optixGetPrimitiveIndex();
2034 float3 start = sbtHG.
cyl.
start[primID];
2035 float3 end = sbtHG.
cyl.
end[primID];
2037 float3 hit = ray_origin + ray_direction * t_hit;
2038 float3 tmp2 = hit - start;
2039 float tmp =
dot(tmp2, axis_u);
2040 float3 Ng =
normalize(hit - (tmp * axis_u + start));
2048 extern "C" __global__
void cylinder_array_color_bounds(
int primIdx,
float result[6]) {
2049 const float3 start = cylinder_buffer[primIdx].start;
2050 const float3 end = start + cylinder_buffer[primIdx].axis;
2051 const float3 rad =
make_float3(cylinder_buffer[primIdx].radius);
2052 optix::Aabb* aabb = (optix::Aabb*)result;
2054 if (rad.x > 0.0f && !isinf(rad.x)) {
2055 aabb->m_min = fminf(start - rad, end - rad);
2056 aabb->m_max =
fmaxf(start + rad, end + rad);
2083 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
2084 const float3 ray_origin = optixGetWorldRayOrigin();
2085 const float3 ray_direction = optixGetWorldRayDirection();
2086 const int primID = optixGetPrimitiveIndex();
2092 if (qmesh.indices == NULL) {
2093 int idx4 = primID*4;
2094 index = make_int4(idx4, idx4+1, idx4+2, idx4+3);
2096 index = qmesh.indices[primID];
2100 const float3 &v00 = qmesh.vertices[index.x];
2101 const float3 &v10 = qmesh.vertices[index.y];
2102 const float3 &v11 = qmesh.vertices[index.z];
2103 const float3 &v01 = qmesh.vertices[index.w];
2105 float3 e01 = v10 - v00;
2106 float3 e03 = v01 - v00;
2107 float3 P =
cross(ray_direction, e03);
2108 float det =
dot(e01, P);
2109 if (
fabsf(det) < quadepsilon)
2113 float inv_det = __frcp_rn(det);
2115 float inv_det = 1.0f / det;
2117 float3 T = ray_origin - v00;
2118 float alpha =
dot(T, P) * inv_det;
2121 #if defined(QUAD_VERTEX_REORDERING) 2125 float3 Q =
cross(T, e01);
2126 float beta =
dot(ray_direction, Q) * inv_det;
2129 #if defined(QUAD_VERTEX_REORDERING) 2134 if ((alpha + beta) > 1.0f) {
2137 float3 e23 = v01 - v11;
2138 float3 e21 = v10 - v11;
2139 float3 P_prime =
cross(ray_direction, e21);
2140 float det_prime =
dot(e23, P_prime);
2141 if (
fabsf(det_prime) < quadepsilon)
2144 float inv_det_prime = __frcp_rn(det_prime);
2146 float inv_det_prime = 1.0f / det_prime;
2148 float3 T_prime = ray_origin - v11;
2149 float alpha_prime =
dot(T_prime, P_prime) * inv_det_prime;
2150 if (alpha_prime < 0.0f)
2152 float3 Q_prime =
cross(T_prime, e23);
2153 float beta_prime =
dot(ray_direction, Q_prime) * inv_det_prime;
2154 if (beta_prime < 0.0f)
2158 float t =
dot(e03, Q) * inv_det;
2165 __float_as_int(alpha),
2166 __float_as_int(beta));
2172 static __host__ __device__ __inline__
2174 float &alpha11,
float &beta11) {
2175 const int primID = optixGetPrimitiveIndex();
2179 const int4 index = qmesh.
indices[primID];
2183 const float3 &v00 = qmesh.vertices[index.x];
2184 const float3 &v10 = qmesh.vertices[index.y];
2185 const float3 &v11 = qmesh.vertices[index.z];
2186 const float3 &v01 = qmesh.vertices[index.w];
2188 float3 e01 = v10 - v00;
2189 float3 e02 = v11 - v00;
2190 float3 e03 = v01 - v00;
2191 float3 n =
cross(e01, e03);
2194 if ((absn.x >= absn.y) && (absn.x >= absn.z)) {
2195 alpha11 = ((e02.y * e03.z) - (e02.z * e03.y)) / n.x;
2196 beta11 = ((e01.y * e02.z) - (e01.z * e02.y)) / n.x;
2197 }
else if ((absn.y >= absn.x) && (absn.y >= absn.z)) {
2198 alpha11 = ((e02.z * e03.x) - (e02.x * e03.z)) / n.y;
2199 beta11 = ((e01.z * e02.x) - (e01.x * e02.z)) / n.y;
2201 alpha11 = ((e02.x * e03.y) - (e02.y * e03.x)) / n.z;
2202 beta11 = ((e01.x * e02.y) - (e01.y * e02.x)) / n.z;
2211 static __host__ __device__ __inline__
2213 const float alpha,
const float beta,
2214 const float &alpha11,
const float &beta11,
2215 float &u,
float &v) {
2217 const float alpha11minus1 = alpha11 - 1.0f;
2218 const float beta11minus1 = beta11 - 1.0f;
2220 if (
fabsf(alpha11minus1) < quadepsilon) {
2223 if (
fabsf(beta11minus1) < quadepsilon) {
2226 v = beta / ((u * beta11minus1) + 1.0f);
2228 }
else if (
fabsf(beta11minus1) < quadepsilon) {
2231 u = alpha / ((v * alpha11minus1) + 1.0f);
2233 float A = -beta11minus1;
2234 float B = (alpha * beta11minus1) - (beta * alpha11minus1) - 1.0f;
2236 float D = (B * B) - (4.0f * A * C);
2237 float Q = -0.5f * (B + (((B < 0.0f) ? -1.0f : 1.0f) * sqrtf(D)));
2239 if ((u < 0.0f) || (u > 1.0f))
2241 v = beta / ((u * beta11minus1) + 1.0f);
2246 static __host__ __device__ __inline__
2248 float3 &shading_normal) {
2249 const int primID = optixGetPrimitiveIndex();
2254 if (qmesh.indices == NULL) {
2255 int idx4 = primID*4;
2256 index = make_int4(idx4, idx4+1, idx4+2, idx4+3);
2258 index = qmesh.indices[primID];
2261 float alpha = __int_as_float(optixGetAttribute_0());
2262 float beta = __int_as_float(optixGetAttribute_1());
2266 float alpha11, beta11;
2277 if (u > 1.0f || v > 1.0f)
2278 printf(
"quad: u:%.2f v:%.2f a:%.2f b:%.2f a11:%.2f b11:%.2f\n",
2279 u, v, alpha, beta, alpha11, beta11);
2287 if (qmesh.normals !=
nullptr) {
2288 const float3 &v00 = qmesh.vertices[index.x];
2289 const float3 &v10 = qmesh.vertices[index.y];
2291 const float3 &v01 = qmesh.vertices[index.w];
2294 const float3& n00 = qmesh.normals[index.x];
2295 const float3& n10 = qmesh.normals[index.y];
2296 const float3& n11 = qmesh.normals[index.z];
2297 const float3& n01 = qmesh.normals[index.w];
2300 Ns =
normalize(((n00 * (1.0f - u)) + (n10 * u)) * (1.0f - v) +
2301 ((n01 * (1.0f - u)) + (n11 * u)) * v);
2303 const float3 &v00 = qmesh.vertices[index.x];
2304 const float3 &v10 = qmesh.vertices[index.y];
2306 const float3 &v01 = qmesh.vertices[index.w];
2312 if (qmesh.vertcolors3f !=
nullptr) {
2313 const float3 c00 = qmesh.vertcolors3f[index.x];
2314 const float3 c10 = qmesh.vertcolors3f[index.y];
2315 const float3 c11 = qmesh.vertcolors3f[index.z];
2316 const float3 c01 = qmesh.vertcolors3f[index.w];
2319 hit_color = ((c00 * (1.0f - u)) + (c10 * u)) * (1.0f - v) +
2320 ((c01 * (1.0f - u)) + (c11 * u)) * v;
2321 }
else if (qmesh.vertcolors4u !=
nullptr) {
2322 const float ci2f = 1.0f / 255.0f;
2323 const float3 c00 = qmesh.vertcolors4u[index.x] * ci2f;
2324 const float3 c10 = qmesh.vertcolors4u[index.y] * ci2f;
2325 const float3 c11 = qmesh.vertcolors4u[index.z] * ci2f;
2326 const float3 c01 = qmesh.vertcolors4u[index.w] * ci2f;
2329 hit_color = ((c00 * (1.0f - u)) + (c10 * u)) * (1.0f - v) +
2330 ((c01 * (1.0f - u)) + (c11 * u)) * v;
2344 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
2345 const float3 obj_ray_origin = optixGetObjectRayOrigin();
2346 const float3 obj_ray_direction = optixGetObjectRayDirection();
2347 const int primID = optixGetPrimitiveIndex();
2349 const float3 center = sbtHG.
ring.
center[primID];
2350 const float3 norm = sbtHG.
ring.
norm[primID];
2351 const float inrad = sbtHG.
ring.
inrad[primID];
2352 const float outrad = sbtHG.
ring.
outrad[primID];
2354 float d = -
dot(center, norm);
2355 float t = -(d +
dot(norm, obj_ray_origin));
2356 float td =
dot(norm, obj_ray_direction);
2360 float3 hit = obj_ray_origin + t * obj_ray_direction;
2361 float rd =
length(hit - center);
2362 if ((rd > inrad) && (rd < outrad)) {
2370 static __host__ __device__ __inline__
2372 const int primID = optixGetPrimitiveIndex();
2375 float3 Ng = sbtHG.
ring.
norm[primID];
2383 extern "C" __global__
void ring_array_color_bounds(
int primIdx,
float result[6]) {
2384 const float3 center = ring_buffer[primIdx].center;
2385 const float3 rad =
make_float3(ring_buffer[primIdx].outrad);
2386 optix::Aabb* aabb = (optix::Aabb*)result;
2388 if (rad.x > 0.0f && !isinf(rad.x)) {
2389 aabb->m_min = center - rad;
2390 aabb->m_max = center + rad;
2400 #if defined(TACHYON_USE_SPHERES_HEARNBAKER) 2409 static __host__ __device__ __inline__
2411 const float3 ray_origin = optixGetObjectRayOrigin();
2412 const float3 obj_ray_direction = optixGetObjectRayDirection();
2417 const float3 ray_direction =
normalize_len(obj_ray_direction, ray_invlen);
2419 float3 deltap = center - ray_origin;
2420 float ddp =
dot(ray_direction, deltap);
2421 float3 remedyTerm = deltap - ddp * ray_direction;
2422 float disc = rad*rad -
dot(remedyTerm, remedyTerm);
2424 float disc_root = sqrtf(disc);
2426 #if 0 && defined(FASTONESIDEDSPHERES) 2427 float t1 = ddp - disc_root;
2431 float t1 = ddp - disc_root;
2435 float t2 = ddp + disc_root;
2447 static __host__ __device__ __inline__
2448 void sphere_intersect_classic(float3 center,
float rad) {
2449 const float3 ray_origin = optixGetObjectRayOrigin();
2450 const float3 obj_ray_direction = optixGetObjectRayDirection();
2455 const float3 ray_direction =
normalize_len(obj_ray_direction, ray_invlen);
2457 float3 deltap = center - ray_origin;
2458 float ddp =
dot(ray_direction, deltap);
2459 float disc = ddp*ddp + rad*rad -
dot(deltap, deltap);
2461 float disc_root = sqrtf(disc);
2463 #if 0 && defined(FASTONESIDEDSPHERES) 2465 float t1 = ddp - disc_root;
2469 float t2 = ddp + disc_root;
2473 float t1 = ddp - disc_root;
2487 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
2488 const int primID = optixGetPrimitiveIndex();
2491 float radius = xyzr.w;
2493 #if defined(TACHYON_USE_SPHERES_HEARNBAKER) 2496 sphere_intersect_classic(center, radius);
2501 static __host__ __device__ __inline__
2503 const float3 ray_origin = optixGetWorldRayOrigin();
2504 const float3 ray_direction = optixGetWorldRayDirection();
2505 const float t_hit = optixGetRayTmax();
2506 const int primID = optixGetPrimitiveIndex();
2511 float radius = xyzr.w;
2512 float3 deltap = center - ray_origin;
2513 float3 Ng = (t_hit * ray_direction - deltap) * (1.0f / radius);
2519 extern "C" __global__
void sphere_array_bounds(
int primIdx,
float result[6]) {
2520 const float3 cen = sphere_buffer[primIdx].center;
2521 const float3 rad =
make_float3(sphere_buffer[primIdx].radius);
2522 optix::Aabb* aabb = (optix::Aabb*)result;
2524 if (rad.x > 0.0f && !isinf(rad.x)) {
2525 aabb->m_min = cen - rad;
2526 aabb->m_max = cen + rad;
2536 extern "C" __global__
void sphere_array_color_bounds(
int primIdx,
float result[6]) {
2537 const float3 cen = sphere_color_buffer[primIdx].center;
2538 const float3 rad =
make_float3(sphere_color_buffer[primIdx].radius);
2539 optix::Aabb* aabb = (optix::Aabb*)result;
2541 if (rad.x > 0.0f && !isinf(rad.x)) {
2542 aabb->m_min = cen - rad;
2543 aabb->m_max = cen + rad;
2555 #if OPTIX_VERSION >= 70100 2556 static __host__ __device__ __inline__
2557 void get_shadevars_curves_linear(
const GeomSBTHG &sbtHG, float3 &hit_color,
2558 float3 &shading_normal) {
2559 const int primID = optixGetPrimitiveIndex();
2567 #if OPTIX_VERSION >= 70400 2568 static __host__ __device__ __inline__
2569 void get_shadevars_curves_catmullrom(
const GeomSBTHG &sbtHG, float3 &hit_color,
2570 float3 &shading_normal) {
2571 const int primID = optixGetPrimitiveIndex();
2584 static __host__ __device__ __inline__
2586 float &hit_alpha, float3 &shading_normal) {
2587 const int primID = optixGetPrimitiveIndex();
2592 if (tmesh.indices == NULL) {
2593 int idx3 = primID*3;
2594 index = make_int3(idx3, idx3+1, idx3+2);
2596 index = tmesh.indices[primID];
2599 const float2 barycentrics = optixGetTriangleBarycentrics();
2603 if (tmesh.packednormals !=
nullptr) {
2608 const float3& n0 =
unpackNormal(tmesh.packednormals[primID].y);
2609 const float3& n1 =
unpackNormal(tmesh.packednormals[primID].z);
2610 const float3& n2 =
unpackNormal(tmesh.packednormals[primID].w);
2614 const float3& n0 =
unpackNormal(tmesh.packednormals[index.x].y);
2615 const float3& n1 =
unpackNormal(tmesh.packednormals[index.y].z);
2616 const float3& n2 =
unpackNormal(tmesh.packednormals[index.z].w);
2620 Ns =
normalize(n0 * (1.0f - barycentrics.x - barycentrics.y) +
2621 n1 * barycentrics.x + n2 * barycentrics.y);
2622 }
else if (tmesh.normals !=
nullptr) {
2623 const float3 &A = tmesh.vertices[index.x];
2624 const float3 &B = tmesh.vertices[index.y];
2625 const float3 &C = tmesh.vertices[index.z];
2628 const float3& n0 = tmesh.normals[index.x];
2629 const float3& n1 = tmesh.normals[index.y];
2630 const float3& n2 = tmesh.normals[index.z];
2633 Ns =
normalize(n0 * (1.0f - barycentrics.x - barycentrics.y) +
2634 n1 * barycentrics.x + n2 * barycentrics.y);
2636 const float3 &A = tmesh.vertices[index.x];
2637 const float3 &B = tmesh.vertices[index.y];
2638 const float3 &C = tmesh.vertices[index.z];
2644 if (tmesh.vertcolors4u !=
nullptr) {
2645 const float ci2f = 1.0f / 255.0f;
2646 const float3 c0 = tmesh.vertcolors4u[index.x] * ci2f;
2647 const float3 c1 = tmesh.vertcolors4u[index.y] * ci2f;
2648 const float3 c2 = tmesh.vertcolors4u[index.z] * ci2f;
2651 hit_color = (c0 * (1.0f - barycentrics.x - barycentrics.y) +
2652 c1 * barycentrics.x + c2 * barycentrics.y);
2653 }
else if (tmesh.vertcolors3f !=
nullptr) {
2654 const float3 c0 = tmesh.vertcolors3f[index.x];
2655 const float3 c1 = tmesh.vertcolors3f[index.y];
2656 const float3 c2 = tmesh.vertcolors3f[index.z];
2659 hit_color = (c0 * (1.0f - barycentrics.x - barycentrics.y) +
2660 c1 * barycentrics.x + c2 * barycentrics.y);
2663 }
else if (tmesh.tex2d !=
nullptr) {
2664 float2 txc0 = tmesh.tex2d[index.x];
2665 float2 txc1 = tmesh.tex2d[index.y];
2666 float2 txc2 = tmesh.tex2d[index.z];
2669 float2 texcoord = (txc0 * (1.0f - barycentrics.x - barycentrics.y) +
2670 txc1 * barycentrics.x + txc2 * barycentrics.y);
2675 float4 tx = tex2D<float4>(mat.tex, texcoord.x, texcoord.y);
2678 }
else if (tmesh.tex3d !=
nullptr) {
2679 float3 txc0 = tmesh.tex3d[index.x];
2680 float3 txc1 = tmesh.tex3d[index.y];
2681 float3 txc2 = tmesh.tex3d[index.z];
2684 float3 texcoord = (txc0 * (1.0f - barycentrics.x - barycentrics.y) +
2685 txc1 * barycentrics.x + txc2 * barycentrics.y);
2690 float4 tx = tex3D<float4>(mat.tex, texcoord.x, texcoord.y, texcoord.z);
2703 __device__ __inline__
void generic_tri_bounds(optix::Aabb *aabb,
2704 float3 v0, float3 v1, float3 v2) {
2708 if (area > 0.0f && !isinf(area)) {
2709 aabb->m_min = fminf(fminf(v0, v1), v2);
2716 aabb->m_min = fminf(fminf(v0, v1), v2);
2725 extern "C" __global__
void ort_tri_intersect(
int primIdx) {
2726 float3 v0 = stri_buffer[primIdx].v0;
2727 float3 v1 = tri_buffer[primIdx].v1;
2728 float3 v2 = tri_buffer[primIdx].v2;
2732 float t, beta, gamma;
2733 if (intersect_triangle(ray, v0, v1, v2, n, t, beta, gamma)) {
2734 if (rtPotentialIntersection(t)) {
2735 shading_normal = geometric_normal =
normalize(n);
2738 prim_color = uniform_color;
2739 rtReportIntersection(0);
2744 extern "C" __global__
void ort_tri_bounds(
int primIdx,
float result[6]) {
2745 float3 v0 = tri_buffer[primIdx].v0;
2746 float3 v1 = tri_buffer[primIdx].v1;
2747 float3 v2 = tri_buffer[primIdx].v2;
2749 optix::Aabb *aabb = (optix::Aabb*)result;
2750 generic_tri_bounds(aabb, v0, v1, v2);
2757 extern "C" __global__
void ort_stri_intersect(
int primIdx) {
2758 float3 v0 = stri_buffer[primIdx].v0;
2759 float3 v1 = stri_buffer[primIdx].v1;
2760 float3 v2 = stri_buffer[primIdx].v2;
2764 float t, beta, gamma;
2765 if (intersect_triangle(ray, v0, v1, v2, n, t, beta, gamma)) {
2766 if (rtPotentialIntersection(t)) {
2767 float3 n0 = stri_buffer[primIdx].n0;
2768 float3 n1 = stri_buffer[primIdx].n1;
2769 float3 n2 = stri_buffer[primIdx].n2;
2770 shading_normal =
normalize(n1*beta + n2*gamma + n0*(1.0f-beta-gamma));
2774 prim_color = uniform_color;
2775 rtReportIntersection(0);
2780 extern "C" __global__
void ort_stri_bounds(
int primIdx,
float result[6]) {
2781 float3 v0 = stri_buffer[primIdx].v0;
2782 float3 v1 = stri_buffer[primIdx].v1;
2783 float3 v2 = stri_buffer[primIdx].v2;
2785 optix::Aabb *aabb = (optix::Aabb*)result;
2786 generic_tri_bounds(aabb, v0, v1, v2);
2793 extern "C" __global__
void ort_vcstri_intersect(
int primIdx) {
2794 float3 v0 = vcstri_buffer[primIdx].v0;
2795 float3 v1 = vcstri_buffer[primIdx].v1;
2796 float3 v2 = vcstri_buffer[primIdx].v2;
2800 float t, beta, gamma;
2801 if (intersect_triangle(ray, v0, v1, v2, n, t, beta, gamma)) {
2802 if (rtPotentialIntersection(t)) {
2803 float3 n0 = vcstri_buffer[primIdx].n0;
2804 float3 n1 = vcstri_buffer[primIdx].n1;
2805 float3 n2 = vcstri_buffer[primIdx].n2;
2806 shading_normal =
normalize(n1*beta + n2*gamma + n0*(1.0f-beta-gamma));
2809 float3 c0 = vcstri_buffer[primIdx].c0;
2810 float3 c1 = vcstri_buffer[primIdx].c1;
2811 float3 c2 = vcstri_buffer[primIdx].c2;
2812 prim_color = c1*beta + c2*gamma + c0*(1.0f-beta-gamma);
2813 rtReportIntersection(0);
2818 extern "C" __global__
void ort_vcstri_bounds(
int primIdx,
float result[6]) {
2819 float3 v0 = vcstri_buffer[primIdx].v0;
2820 float3 v1 = vcstri_buffer[primIdx].v1;
2821 float3 v2 = vcstri_buffer[primIdx].v2;
2823 optix::Aabb *aabb = (optix::Aabb*)result;
2824 generic_tri_bounds(aabb, v0, v1, v2);
2836 static __device__ __forceinline__
float fog_coord(float3 hit_point) {
2839 const float3 ray_direction = optixGetWorldRayDirection();
2840 const float t_hit = optixGetRayTmax();
2847 switch (scene.fog_mode) {
2849 f = (scene.fog_end - r) / (scene.fog_end - scene.fog_start);
2856 v = scene.fog_density * r;
2864 v = scene.fog_density * r;
2872 return __saturatef(f);
2876 static __device__ __forceinline__ float3
fog_color(
float fogmod, float3 hit_col) {
2895 unsigned int randseed = tea<2>(teabits1, teabits1);
2901 float ndotambl =
dot(N, dir);
2904 if (ndotambl < 0.0f) {
2905 ndotambl = -ndotambl;
2909 float3 aoray_origin, aoray_direction;
2911 #ifdef USE_REVERSE_SHADOW_RAYS 2923 aoray_direction = -dir;
2927 #if defined(TACHYON_USE_RAY_STEP) 2932 aoray_direction = dir;
2936 uint32_t p0 = __float_as_int(1.0f);
2944 OptixVisibilityMask( 255 ),
2946 OPTIX_RAY_FLAG_NONE,
2951 OPTIX_RAY_FLAG_DISABLE_ANYHIT
2952 | OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT
2953 | OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT,
2960 inten += ndotambl * __int_as_float(p0);
2969 template<
int SHADOWS_ON>
2972 float3 &N, float3 &L,
2978 float shadow_tmax) {
2979 float inten =
dot(N, L);
2982 float light_attenuation =
static_cast<float>(inten > 0.0f);
2985 float3 shadowray_origin, shadowray_direction;
2986 float tmax=shadow_tmax;
2987 #ifdef USE_REVERSE_SHADOW_RAYS 2999 shadowray_direction = -L
3000 tmax = fminf(tmax, shadow_tmax));
3006 shadowray_direction = L;
3010 uint32_t p0 = __float_as_int(1.0f);
3014 shadowray_direction,
3018 OptixVisibilityMask( 255 ),
3020 OPTIX_RAY_FLAG_NONE,
3025 OPTIX_RAY_FLAG_DISABLE_ANYHIT
3026 | OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT
3027 | OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT,
3033 light_attenuation = __int_as_float(p0);
3035 #if defined(TACHYON_RAYSTATS) 3043 if (!SHADOWS_ON || light_attenuation > 0.0f) {
3044 result += col * p_Kd * inten * light_attenuation;
3047 const float3 ray_direction = optixGetWorldRayDirection();
3048 float3 H =
normalize(L - ray_direction);
3049 float nDh =
dot(N, H);
3051 float power = powf(nDh, p_phong_exp);
3052 phongcol +=
make_float3(p_Ks) * power * light_attenuation;
3081 template<
int CLIP_VIEW_ON,
3088 int TRANSMISSION_ON>
3090 float p_Ka,
float p_Kd,
float p_Ks,
3091 float p_phong_exp,
float p_reflectivity,
3093 float p_outline,
float p_outlinewidth,
3096 const float3 ray_origin = optixGetWorldRayOrigin();
3097 const float3 ray_direction = optixGetWorldRayDirection();
3098 const float t_hit = optixGetRayTmax();
3100 float3 hit_point = ray_origin + t_hit * ray_direction;
3106 float fogmod = 1.0f;
3111 #if defined(TACHYON_RAYSTATS) 3126 if ((p_opacity < 1.0f) && (prd.
transcnt < 1)) {
3130 new_prd.
alpha = 1.0f;
3144 float3 transray_direction = ray_direction;
3145 float3 transray_origin;
3146 #if defined(TACHYON_USE_RAY_STEP) 3147 #if defined(TACHYON_TRANS_USE_INCIDENT) 3156 transray_origin = hit_point;
3170 OptixVisibilityMask( 255 ),
3171 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
3178 #if defined(TACHYON_RAYSTATS) 3188 float3 col = prim_color;
3193 shade_light<SHADOWS_ON>(result, hit_point, N, L, p_Kd, p_Ks, p_phong_exp,
3200 float3 L = Lpos - hit_point;
3203 shade_light<SHADOWS_ON>(result, hit_point, N, L, p_Kd, p_Ks, p_phong_exp,
3204 col, phongcol, shadow_tmax);
3213 shade_light<SHADOWS_ON>(result, hit_point, N, L, p_Kd, p_Ks, p_phong_exp,
3214 col, phongcol, shadow_tmax);
3222 #if defined(TACHYON_RAYSTATS) 3228 if (OUTLINE_ON && p_outline > 0.0f) {
3229 float edgefactor =
dot(N, ray_direction);
3230 edgefactor *= edgefactor;
3231 edgefactor = 1.0f - edgefactor;
3232 edgefactor = 1.0f - powf(edgefactor, (1.0f - p_outlinewidth) * 32.0f);
3233 float outlinefactor = __saturatef((1.0f - p_outline) + (edgefactor * p_outline));
3234 result *= outlinefactor;
3243 if (REFLECTION_ON && p_reflectivity > 0.0f) {
3252 float3 reflray_direction =
reflect(ray_direction, N);
3254 float3 reflray_origin;
3255 #if defined(TACHYON_USE_RAY_STEP) 3258 reflray_origin = hit_point;
3272 OptixVisibilityMask( 255 ),
3273 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
3280 #if defined(TACHYON_RAYSTATS) 3283 result += p_reflectivity * new_prd.
result;
3290 float alpha = p_opacity;
3298 float fade_start = 1.00f;
3299 float fade_end = 0.20f;
3304 float fade_len = fade_start - fade_end;
3305 alpha *= __saturatef((camdist - fade_start) / fade_len);
3316 if ((TRANSMISSION_ON || CLIP_VIEW_ON) && alpha < 0.999f ) {
3319 alpha = 1.0f + cosf(3.1415926f * (1.0f-alpha) *
dot(N, ray_direction));
3320 alpha = alpha*alpha * 0.25f;
3328 new_prd.
alpha = 1.0f;
3333 float3 transray_direction = ray_direction;
3334 float3 transray_origin;
3335 #if defined(TACHYON_USE_RAY_STEP) 3336 #if defined(TACHYON_TRANS_USE_INCIDENT) 3345 transray_origin = hit_point;
3359 OptixVisibilityMask( 255 ),
3360 OPTIX_RAY_FLAG_DISABLE_ANYHIT,
3367 #if defined(TACHYON_RAYSTATS) 3371 result += (1.0f - alpha) * new_prd.
result;
3372 prd.
alpha = alpha + (1.0f - alpha) * new_prd.
alpha;
3377 if (FOG_ON && fogmod < 1.0f) {
3396 const GeomSBTHG &sbtHG = *
reinterpret_cast<const GeomSBTHG*
>(optixGetSbtDataPointer());
3399 float3 shading_normal;
3401 float hit_alpha=1.0f;
3402 int vertexcolorset=0;
3405 #if defined(TACHYON_MERGED_CLOSESTHIT_DISPATCH) 3406 unsigned int hit_kind = optixGetHitKind();
3407 unsigned int hit_prim_type = 0;
3408 #if OPTIX_VERSION >= 70100 3409 hit_prim_type = optixGetPrimitiveType(hit_kind);
3419 unsigned int mergeprimtype = (hit_prim_type << 16) | (0xFE & hit_kind);
3420 switch (mergeprimtype) {
3447 #if OPTIX_VERSION >= 70400 3448 case RT_PRM_CATMULLROM:
3449 get_shadevars_curves_catmullrom(sbtHG, hit_color, shading_normal);
3452 #if OPTIX_VERSION >= 70200 3454 get_shadevars_curves_linear(sbtHG, hit_color, shading_normal);
3460 printf(
"Unrecognized merged prim: %08x\n", mergeprimtype);
3466 #else // !defined(TACHYON_MERGED_CLOSESTHIT_DISPATCH) 3469 unsigned int hit_kind = optixGetHitKind();
3471 #if !defined(TACHYON_FLATTEN_CLOSESTHIT_DISPATCH) 3472 #if OPTIX_VERSION >= 70100 3474 unsigned int hit_prim_type = optixGetPrimitiveType(hit_kind);
3478 if (hit_prim_type == OPTIX_PRIMITIVE_TYPE_TRIANGLE) {
3480 }
else if (hit_prim_type == OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR) {
3481 get_shadevars_curves_linear(sbtHG, hit_color, shading_normal);
3482 #if OPTIX_VERSION >= 70400 3483 }
else if (hit_prim_type == OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM) {
3484 get_shadevars_curves_catmullrom(sbtHG, hit_color, shading_normal);
3488 #endif // TACHYON_FLATTEN_CLOSESTHIT_DISPATCH 3495 #if (OPTIX_VERSION < 70100) || defined(TACHYON_FLATTEN_CLOSESTHIT_DISPATCH) 3499 case OPTIX_HIT_KIND_TRIANGLE_FRONT_FACE:
3500 case OPTIX_HIT_KIND_TRIANGLE_BACK_FACE:
3527 #if defined(TACHYON_FLATTEN_CLOSESTHIT_DISPATCH) 3528 #if OPTIX_VERSION >= 70100 3531 OptixPrimitiveType hit_prim_type = optixGetPrimitiveType(hit_kind);
3534 if (hit_prim_type == OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR) {
3535 get_shadevars_curves_linear(sbtHG, hit_color, shading_normal);
3536 #if OPTIX_VERSION >= 70400 3537 }
else if (hit_prim_type == OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM) {
3538 get_shadevars_curves_catmullrom(sbtHG, hit_color, shading_normal);
3547 #endif // !defined(TACHYON_MERGED_CLOSESTHIT_DISPATCH) 3550 if (!vertexcolorset) {
3552 const int primID = optixGetPrimitiveIndex();
3565 shader_template<0, 0, 1, 1, 1, 1, 1, 1>(hit_color, shading_normal,
3566 mat.ambient, mat.diffuse,
3567 mat.specular, mat.shininess,
3569 mat.opacity * hit_alpha,
3570 mat.outline, mat.outlinewidth,
__global__ void __closesthit__radiance_general()
__host__ __device__ float3 normalize(const float3 &v)
Normalize input vector to unit length.
static __host__ __device__ __inline__ void quad_calc_barycentrics_v11(const GeomSBTHG &sbtHG, float &alpha11, float &beta11)
__global__ void __raygen__camera_dome_master_dof()
static __forceinline__ __device__ void packPointer(void *ptr, uint32_t &i0, uint32_t &i1)
__global__ void __exception__all()
static __inline__ __device__ float3 calc_ffworld_normal(const float3 &Nshading, const float3 &Ngeometric)
__global__ void __raygen__camera_orthographic_stereo()
uchar4 * framebuffer
8-bit unorm RGBA framebuffer
float accum_normalize
precalc 1.0f / subframe_index
#define REVERSE_RAY_LENGTH
static __device__ __inline__ void tachyon_camera_perspective_general()
int headlight_mode
Extra VR camera-located headlight.
static __device__ __inline__ void jitter_offset2f(unsigned int &pval, float2 &xy)
float bg_grad_invrange
miss background gradient inverse range
static __device__ __inline__ void dof_ray(const float cam_dof_focal_dist, const float cam_dof_aperture_rad, const float3 &ray_origin_orig, float3 &ray_origin, const float3 &ray_direction_orig, float3 &ray_direction, unsigned int &randseed, const float3 &up, const float3 &right)
float4 * PosRadius
X,Y,Z,Radius packed for coalescing.
static __forceinline__ __device__ void * unpackPointer(uint32_t i0, uint32_t i1)
float3 bg_grad_updir
miss background gradient up direction
__constant__ tachyonLaunchParams rtLaunch
launch parameters in constant memory, filled by optixLaunch)
int update_colorbuffer
accumulation copyout flag
static __host__ __device__ __inline__ void sphere_intersect_hearn_baker(float3 center, float rad)
static __device__ __inline__ void jitter_disc2f_qrn(float2 &qrnxy, float2 &xy, float radius)
int tonemap_mode
output tone mapping mode
__host__ __device__ float3 normalize_len(const float3 v, float &l)
Normalize input vector to unit length, and return its original length.
float ao_lightscale
2.0f/float(ao_samples)
static __device__ float shade_ambient_occlusion(float3 hit, float3 N, float aoimportance)
static __forceinline__ __device__ float getPayloadShadowAttenuation()
any-hit programs read-modify-update shadow attenuation value carried in ray payload register 0 ...
__global__ void __raygen__camera_perspective()
static __device__ __inline__ void tachyon_camera_octahedral_general()
__global__ void __raygen__camera_equirectangular_stereo_dof()
float ao_maxdist
AO maximum occlusion distance.
__host__ __device__ float4 make_float4(const float3 &a, const float &b)
__host__ __device__ float3 make_float3(const float s)
float3 * prim_color
optional per-primitive color array
float3 W
camera orthonormal W (view) axis
static void __inline__ __device__ accumulate_color(int idx, float4 colrgba4f)
__global__ void __intersection__sphere_array()
__host__ __device__ float3 fmaxf(const float3 &a, const float3 &b)
float3 pos
camera position
__global__ void __intersection__cone_array_color()
__device__ void clip_ray_by_plane(float3 ray_origin, float3 ray_direction, float &tmin, float &tmax, const float4 plane)
static __device__ __forceinline__ float3 fog_color(float fogmod, float3 hit_col)
int max_trans
max transparent surface crossing count
float tonemap_exposure
tone mapping exposure gain parameter
Tachyon ray tracing engine core routines and data structures compiled to PTX for runtime JIT to build...
float ao_direct
AO direct lighting scaling factor.
__global__ void __raygen__camera_oculus_rift()
int materialindex
material index for this array
__global__ void __raygen__camera_oculus_rift_stereo_dof()
__host__ __device__ float3 faceforward(const float3 &n, const float3 &i, const float3 &nref)
Ensure that an interpolated surface normal n faces in the same direction as dictated by a geometric n...
static __host__ __device__ __inline__ void quad_calc_bilinear_coords(const GeomSBTHG &sbtHG, const float alpha, const float beta, const float &alpha11, const float &beta11, float &u, float &v)
Tachyon OptiX global launch parameter structure containing the active camera, framebuffer, materials, and any global scene parameters required for shading.
void jitter_sphere3f(rng_frand_handle *rngh, float *dir)
int fb_clearall
clear/overwrite all FB components
struct tachyonLaunchParams::@4 lights
__host__ __device__ float3 fabsf(const float3 &a)
int fog_mode
fog type (or off)
static __host__ __device__ __inline__ void get_shadevars_ring_array(const GeomSBTHG &sbtHG, float3 &shading_normal)
__global__ void __closesthit__shadow_nop()
__global__ void any_hit_shadow_clip_sphere()
static __forceinline__ __device__ T * getPRD()
__global__ void __raygen__UNKNOWN()
static __host__ __device__ __inline__ void get_shadevars_trimesh(const GeomSBTHG &sbtHG, float3 &hit_color, float &hit_alpha, float3 &shading_normal)
__global__ void __raygen__camera_octahedral_stereo()
__host__ __device__ float3 reflect(const float3 &i, const float3 &n)
calculate reflection direction from incident direction i, and surface normal n.
__global__ void __miss__radiance_solid_bg()
static __device__ __inline__ void jitter_disc2f(unsigned int &pval, float2 &xy, float radius)
__global__ void __intersection__cylinder_array_color()
__global__ void __raygen__camera_equirectangular_stereo()
__global__ void __anyhit__shadow_transmission()
static __host__ __device__ __inline__ float3 unpackNormal(uint packed)
static __forceinline__ __device__ void setPayloadShadowAttenuation(const float attenuation)
any-hit programs read-modify-update shadow attenuation value carried in ray payload register 0 ...
int shadows_enabled
global shadow flag
__global__ void __raygen__camera_dome_master_stereo_dof()
static __device__ __forceinline__ float fog_coord(float3 hit_point)
__global__ void __anyhit__shadow_opaque()
__global__ void __raygen__camera_perspective_stereo_dof()
__global__ void __raygen__camera_dome_master()
float3 * dir_lights
list of directional light directions
float bg_grad_botval
miss background gradient bottom value
struct tachyonLaunchParams::@3 scene
static int __forceinline__ __device__ subframe_count()
__device__ void sphere_fade_and_clip(const float3 &hit_point, const float3 &cam_pos, float fade_start, float fade_end, float &alpha)
enable alpha transparency
__host__ __device__ float dot(const float3 &a, const float3 &b)
static __device__ __inline__ void tachyon_camera_dome_general()
__global__ void __raygen__camera_equirectangular_dof()
__global__ void __raygen__camera_cubemap()
static __host__ __device__ __inline__ void get_shadevars_sphere_array(const GeomSBTHG &sbtHG, float3 &shading_normal)
__global__ void __intersection__ring_array()
#define TACHYON_RAY_STEP2
static __device__ void shader_template(float3 prim_color, float3 N, float p_Ka, float p_Kd, float p_Ks, float p_phong_exp, float p_reflectivity, float p_opacity, float p_outline, float p_outlinewidth, int p_transmode)
material-specific shading property
int2 size
framebuffer size
float3 bg_color_grad_bot
miss background gradient (bottom)
static __host__ __device__ __inline__ void get_shadevars_quadmesh(const GeomSBTHG &sbtHG, float3 &hit_color, float3 &shading_normal)
static __forceinline__ __device__ float4 linear_to_sRGB_approx_20(const float4 &linear)
__global__ void __raygen__camera_cubemap_dof()
any-hit traversal reversal
__global__ void __miss__radiance_gradient_bg_sky_sphere()
__global__ void __anyhit__radiance_nop()
static __forceinline__ __device__ uint32_t getPayloadAAsample()
radiance PRD aasample count is stored in ray payload register 2
__global__ void __raygen__camera_orthographic_stereo_dof()
int ao_samples
number of AO samples per AA ray
__global__ void __intersection__quadmesh()
static __host__ __device__ __inline__ void get_shadevars_cone_array(const GeomSBTHG &sbtHG, float3 &shading_normal)
__global__ void __raygen__camera_orthographic_dof()
int subframe_index
accumulation subframe index
__global__ void __raygen__camera_perspective_dof()
__host__ __device__ float3 cross(const float3 &a, const float3 &b)
calculate the cross product between vectors a and b.
__global__ void __raygen__camera_octahedral()
float bg_grad_noisemag
miss background gradient noise magnitude
int aa_samples
AA samples per launch.
static __device__ __inline__ void shade_light(float3 &result, float3 &hit_point, float3 &N, float3 &L, float p_Kd, float p_Ks, float p_phong_exp, float3 &col, float3 &phongcol, float shadow_tmax)
scene-wide shading property
static __device__ __inline__ void tachyon_camera_equirectangular_general()
OptixTraversableHandle traversable
global OptiX scene traversable handle
__global__ void __raygen__camera_equirectangular()
float3 bg_color_grad_top
miss background gradient (top)
float clipview_end
clipping sphere/plane end coord
static __device__ __inline__ void tachyon_camera_oculus_rift_general()
struct tachyonLaunchParams::@2 frame
__host__ __device__ float length(const float3 &v)
float clipview_start
clipping sphere/plane start coord
static __device__ __inline__ void tachyon_camera_orthographic_general()
float opacity
surface opacity
__device__ void ray_sphere_clip_interval(float3 ray_origin, float3 ray_direction, float3 center, float rad, float2 &tinterval)
int colorspace
output colorspace
struct tachyonLaunchParams::@5 cam
int num_dir_lights
directional light count
float3 * pos_lights
list of positional light positions
#define UINT32_RAND_MAX_INV
__global__ void __miss__shadow_nop()
custom prim quadrilateral
__global__ void __raygen__camera_octahedral_dof()
__global__ void __raygen__camera_dome_master_stereo()
__global__ void __raygen__camera_cubemap_stereo()
rt_material * materials
device memory material array
__global__ void __raygen__camera_oculus_rift_stereo()
__global__ void __raygen__camera_perspective_stereo()
custom prim quadrilateral
float3 bg_color
miss background color
__global__ void __raygen__camera_oculus_rift_dof()
int num_pos_lights
positional light count
float ao_ambient
AO ambient factor.
float3 uniform_color
uniform color for entire sphere array
float4 * accum_buffer
32-bit FP RGBA accumulation buffer
__global__ void __raygen__camera_octahedral_stereo_dof()
float stereo_eyesep
stereo eye separation, in world coords
__global__ void __miss__radiance_gradient_bg_sky_plane()
int clipview_mode
VR clipping view on/off.
static __device__ __inline__ int tachyon1DLaunchIndex(void)
Helper function to return 1-D framebuffer offset computed from the current thread's launch_index...
int max_depth
global max ray tracing recursion depth
enable tex cutout transparency
__global__ void __raygen__camera_orthographic()
static __device__ __inline__ void tachyon_camera_cubemap_general()
float epsilon
global epsilon value
__global__ void __raygen__camera_cubemap_stereo_dof()
static __host__ __device__ __inline__ void get_shadevars_cylinder_array(const GeomSBTHG &sbtHG, float3 &shading_normal)
static __forceinline__ __device__ float4 linear_to_sRGB(const float4 &lin)