Page MenuHome
Differential D12909 Diff 43508 intern/cycles/kernel/device/optix/kernel.cu

Changeset View

Standalone View

View Options

intern/cycles/kernel/device/optix/kernel.cu

Show First 20 Lines • Show All 169 Lines • ▼ Show 20 Lines
} }
extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
{ {
#ifdef __SHADOW_RECORD_ALL__ #ifdef __SHADOW_RECORD_ALL__
int prim = optixGetPrimitiveIndex(); int prim = optixGetPrimitiveIndex();
const uint object = get_object_id(); const uint object = get_object_id();
# ifdef __VISIBILITY_FLAG__ # ifdef __VISIBILITY_FLAG__
const uint visibility = optixGetPayload_4(); const uint visibility = optixGetPayload_3();
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) { if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection(); return optixIgnoreIntersection();
} }
# endif # endif
float u = 0.0f, v = 0.0f; float u = 0.0f, v = 0.0f;
int type = 0; int type = 0;
if (optixIsTriangleHit()) { if (optixIsTriangleHit()) {
Show All 15 Lines else {
if (u == 0.0f || u == 1.0f) { if (u == 0.0f || u == 1.0f) {
return optixIgnoreIntersection(); return optixIgnoreIntersection();
} }
} }
# endif # endif
# ifndef __TRANSPARENT_SHADOWS__ # ifndef __TRANSPARENT_SHADOWS__
/* No transparent shadows support compiled in, make opaque. */ /* No transparent shadows support compiled in, make opaque. */
optixSetPayload_5(true); optixSetPayload_4(true);
return optixTerminateRay(); return optixTerminateRay();
# else # else
const int max_hits = optixGetPayload_3(); const int max_hits = optixGetPayload_2();
/* If no transparent shadows, all light is blocked and we can stop immediately. */ /* If no transparent shadows, all light is blocked and we can stop immediately. */
if (max_hits == 0 || if (max_hits == 0 ||
!(intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) { !(intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) {
optixSetPayload_5(true); optixSetPayload_4(true);
return optixTerminateRay(); return optixTerminateRay();
} }
/* Record transparent intersection. */ /* Record transparent intersection. */
const int num_hits = optixGetPayload_2(); const int num_hits = optixGetPayload_1();
int record_index = num_hits; int record_index = num_hits;
optixSetPayload_2(num_hits + 1); optixSetPayload_1(num_hits + 1);
Intersection *const isect_array = get_payload_ptr_0<Intersection>(); const IntegratorShadowState state = optixGetPayload_0();
if (record_index >= max_hits) { if (record_index >= max_hits) {
/* If maximum number of hits reached, find a hit to replace. */ /* If maximum number of hits reached, find a hit to replace. */
float max_recorded_t = isect_array[0].t; float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t);
int max_recorded_hit = 0; int max_recorded_hit = 0;
for (int i = 1; i < max_hits; i++) { for (int i = 1; i < max_hits; i++) {
if (isect_array[i].t > max_recorded_t) { const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t);
max_recorded_t = isect_array[i].t; if (isect_t > max_recorded_t) {
max_recorded_t = isect_t;
max_recorded_hit = i; max_recorded_hit = i;
} }
} }
if (optixGetRayTmax() >= max_recorded_t) { if (optixGetRayTmax() >= max_recorded_t) {
/* Accept hit, so that OptiX won't consider any more hits beyond the distance of the /* Accept hit, so that OptiX won't consider any more hits beyond the distance of the
* current hit anymore. */ * current hit anymore. */
return; return;
} }
record_index = max_recorded_hit; record_index = max_recorded_hit;
} }
Intersection *const isect = isect_array + record_index; INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u;
isect->u = u; INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v;
isect->v = v; INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = optixGetRayTmax();
isect->t = optixGetRayTmax(); INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim;
isect->prim = prim; INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object;
isect->object = object; INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type;
isect->type = type;
optixIgnoreIntersection(); optixIgnoreIntersection();
# endif /* __TRANSPARENT_SHADOWS__ */ # endif /* __TRANSPARENT_SHADOWS__ */
#endif /* __SHADOW_RECORD_ALL__ */ #endif /* __SHADOW_RECORD_ALL__ */
} }
extern "C" __global__ void __anyhit__kernel_optix_volume_test() extern "C" __global__ void __anyhit__kernel_optix_volume_test()
{ {
▲ Show 20 Lines • Show All 119 Lines • Show Last 20 Lines