From d06828f0b8ebb083de59fd2cb8c5f8fe6af1da22 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Mon, 18 Oct 2021 19:20:09 +0200 Subject: 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 --- intern/cycles/kernel/device/optix/kernel.cu | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) (limited to 'intern/cycles/kernel/device') 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(); + 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__ */ -- cgit v1.2.3