diff options
author | Brecht Van Lommel <brecht@blender.org> | 2021-10-18 20:20:09 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2021-10-19 16:10:55 +0300 |
commit | d06828f0b8ebb083de59fd2cb8c5f8fe6af1da22 (patch) | |
tree | 6ed5d1a8a488d1badaf5b8e462f7c1df2e75cbb2 /intern/cycles/kernel/device | |
parent | 943e73b07e26d64c04ccb7d8f656e3818a57cca0 (diff) |
Cycles: avoid intermediate stack array for writing shadow intersections
Helps save one OptiX payload and is a bit more efficient.
Differential Revision: https://developer.blender.org/D12909
Diffstat (limited to 'intern/cycles/kernel/device')
-rw-r--r-- | intern/cycles/kernel/device/optix/kernel.cu | 22 |
1 files changed, 11 insertions, 11 deletions
diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index e97b25d31a2..574f66ab708 100644 --- a/intern/cycles/kernel/device/optix/kernel.cu +++ b/intern/cycles/kernel/device/optix/kernel.cu @@ -225,16 +225,17 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() optixSetPayload_2(num_hits + 1); - Intersection *const isect_array = get_payload_ptr_0<Intersection>(); + const IntegratorShadowState state = optixGetPayload_0(); if (record_index >= max_hits) { /* 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; for (int i = 1; i < max_hits; i++) { - if (isect_array[i].t > max_recorded_t) { - max_recorded_t = isect_array[i].t; + const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); + if (isect_t > max_recorded_t) { + max_recorded_t = isect_t; max_recorded_hit = i; } } @@ -248,13 +249,12 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() record_index = max_recorded_hit; } - Intersection *const isect = isect_array + record_index; - isect->u = u; - isect->v = v; - isect->t = optixGetRayTmax(); - isect->prim = prim; - isect->object = object; - isect->type = type; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = optixGetRayTmax(); + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type; optixIgnoreIntersection(); # endif /* __TRANSPARENT_SHADOWS__ */ |