19 #if defined(__BRANCHED_PATH__) && defined(__VOLUME__)
21 ccl_device_inline void kernel_split_branched_path_volume_indirect_light_init(KernelGlobals *
kg,
24 kernel_split_branched_path_indirect_loop_init(
kg, ray_index);
39 int num_samples =
kernel_data.integrator.volume_samples;
40 float num_samples_inv = 1.0f / num_samples;
42 Ray volume_ray = branched_state->ray;
44 branched_state->isect.t :
47 float step_size = volume_stack_step_size(
kg, branched_state->path_state.volume_stack);
49 for (
int j = branched_state->next_sample; j < num_samples; j++) {
51 *ps = branched_state->path_state;
54 *pray = branched_state->ray;
57 *tp = branched_state->throughput * num_samples_inv;
64 kg, ps, sd, &volume_ray,
L, tp, step_size);
66 # ifdef __VOLUME_SCATTER__
69 kernel_path_volume_connect_light(
kg, sd, emission_sd, *tp, &branched_state->path_state,
L);
72 if (!kernel_path_volume_bounce(
kg, sd, tp, ps, &
L->state, pray)) {
77 branched_state->next_closure = 0;
78 branched_state->next_sample = j + 1;
86 if (branched_state->shared_sample_count < 2) {
87 if (kernel_split_branched_indirect_start_shared(
kg, ray_index)) {
97 branched_state->next_sample = num_samples;
99 branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0);
100 if (branched_state->waiting_on_shared_samples) {
104 kernel_split_branched_path_indirect_loop_end(
kg, ray_index);
108 kernel_volume_shadow(
123 # ifdef __BRANCHED_PATH__
156 kernel_volume_clean_stack(
kg,
state->volume_stack);
160 Ray volume_ray = *ray;
161 volume_ray.
t = (hit) ? isect->t : FLT_MAX;
167 float step_size = volume_stack_step_size(
kg,
state->volume_stack);
172 kg,
state, sd, &volume_ray,
L, throughput, step_size);
174 # ifdef __VOLUME_SCATTER__
177 kernel_path_volume_connect_light(
kg, sd, emission_sd, *throughput,
state,
L);
180 if (kernel_path_volume_bounce(
kg, sd, throughput,
state, &
L->state, ray)) {
190 # ifdef __BRANCHED_PATH__
193 kernel_split_branched_path_volume_indirect_light_init(
kg, ray_index);
195 if (kernel_split_branched_path_volume_indirect_light_iter(
kg, ray_index)) {
203 # ifdef __BRANCHED_PATH__
218 if (kernel_split_branched_path_volume_indirect_light_iter(
kg, ray_index)) {
ccl_device_inline void path_radiance_sum_indirect(PathRadiance *L)
ccl_device_inline void path_radiance_reset_indirect(PathRadiance *L)
#define ccl_global_size(d)
#define ccl_device_inline
#define ccl_device_noinline
#define CCL_NAMESPACE_END
CCL_NAMESPACE_BEGIN ccl_device void kernel_do_volume(KernelGlobals *kg)
ccl_device_inline void path_state_branch(ccl_addr_space PathState *state, int branch, int num_branches)
ccl_device int get_ray_index(KernelGlobals *kg, int thread_index, int queue_number, ccl_global int *queues, int queuesize, int empty_queue)
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 AS_SHADER_DATA(shader_data_tiny_storage)
#define IS_STATE(ray_state, ray_index, state)
#define __BRANCHED_PATH__
#define ADD_RAY_FLAG(ray_state, ray_index, flag)
#define ASSIGN_RAY_STATE(ray_state, ray_index, state)
@ RAY_BRANCHED_VOLUME_INDIRECT
@ RAY_VOLUME_INDIRECT_NEXT_ITER
@ QUEUE_VOLUME_INDIRECT_ITER
@ QUEUE_ACTIVE_AND_REGENERATED_RAYS