48 #ifdef __BRANCHED_PATH__
49 ccl_device_inline void kernel_split_branched_indirect_light_init(KernelGlobals *
kg,
int ray_index)
51 kernel_split_branched_path_indirect_loop_init(
kg, ray_index);
56 ccl_device void kernel_split_branched_transparent_bounce(KernelGlobals *
kg,
int ray_index)
79 if (!path_state_volume_next(
kg,
state)) {
87 ray->t -= sd->ray_length;
89 # ifdef __RAY_DIFFERENTIALS__
91 ray->dD.dx = -sd->dI.dx;
92 ray->dD.dy = -sd->dI.dy;
97 kernel_volume_stack_enter_exit(
kg, sd,
state->volume_stack);
106 *local_queue_atomics = 0;
150 #ifdef __BRANCHED_PATH__
157 #ifdef __BRANCHED_PATH__
160 kernel_split_branched_transparent_bounce(
kg, ray_index);
163 kernel_split_branched_indirect_light_init(
kg, ray_index);
165 if (kernel_split_branched_path_surface_indirect_light_iter(
166 kg, ray_index, 1.0f,
kernel_split_sd(branched_state_sd, ray_index),
true,
true)) {
170 kernel_split_branched_path_indirect_loop_end(
kg, ray_index);
171 kernel_split_branched_transparent_bounce(
kg, ray_index);
186 #ifdef __BRANCHED_PATH__
207 if (kernel_split_branched_path_surface_indirect_light_iter(
208 kg, ray_index, 1.0f,
kernel_split_sd(branched_state_sd, ray_index),
true,
true)) {
212 kernel_split_branched_path_indirect_loop_end(
kg, ray_index);
213 kernel_split_branched_transparent_bounce(
kg, ray_index);
221 *local_queue_atomics = 0;
237 # ifdef __SUBSURFACE__
241 *local_queue_atomics = 0;
ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
ccl_device_inline void path_radiance_sum_indirect(PathRadiance *L)
CCL_NAMESPACE_BEGIN ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, const ShaderData *sd)
ccl_device_inline void path_radiance_reset_indirect(PathRadiance *L)
#define ccl_global_size(d)
ccl_device_inline uint ccl_local_id(uint d)
#define ccl_device_inline
#define CCL_NAMESPACE_END
CCL_NAMESPACE_BEGIN ccl_device void kernel_next_iteration_setup(KernelGlobals *kg, ccl_local_param unsigned int *local_queue_atomics)
ccl_device_inline void path_state_next(KernelGlobals *kg, ccl_addr_space PathState *state, int label)
ccl_device bool kernel_path_surface_bounce(KernelGlobals *kg, ShaderData *sd, ccl_addr_space float3 *throughput, ccl_addr_space PathState *state, PathRadianceState *L_state, ccl_addr_space Ray *ray)
ccl_device int get_ray_index(KernelGlobals *kg, int thread_index, int queue_number, ccl_global int *queues, int queuesize, int empty_queue)
ccl_device void enqueue_ray_index_local(int ray_index, int queue_number, char enqueue_flag, int queuesize, ccl_local_param unsigned int *local_queue_atomics, ccl_global int *Queue_data, ccl_global int *Queue_index)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_split_path_end(KernelGlobals *kg, int ray_index)
#define kernel_split_params
#define kernel_split_sd(sd, ray_index)
#define kernel_split_state
__kernel void ccl_constant KernelData ccl_global void ccl_global char * ray_state
#define IS_FLAG(ray_state, ray_index, flag)
#define IS_STATE(ray_state, ray_index, state)
#define ADD_RAY_FLAG(ray_state, ray_index, flag)
#define ASSIGN_RAY_STATE(ray_state, ray_index, state)
@ RAY_LIGHT_INDIRECT_NEXT_ITER
@ RAY_BRANCHED_LIGHT_INDIRECT
@ RAY_SUBSURFACE_INDIRECT_NEXT_ITER
@ RAY_VOLUME_INDIRECT_NEXT_ITER
@ QUEUE_LIGHT_INDIRECT_ITER
@ QUEUE_SHADOW_RAY_CAST_DL_RAYS
@ QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS
@ QUEUE_SHADOW_RAY_CAST_AO_RAYS
@ QUEUE_VOLUME_INDIRECT_ITER
@ QUEUE_SUBSURFACE_INDIRECT_ITER
@ QUEUE_ACTIVE_AND_REGENERATED_RAYS
bool active
all scheduled work for the GPU.
#define CCL_LOCAL_MEM_FENCE
#define ccl_barrier(flags)
ccl_device_inline bool is_zero(const float2 &a)