19 #ifdef __BRANCHED_PATH__
28 # define BRANCHED_STORE(name) branched_state->name = kernel_split_state.name[ray_index];
30 BRANCHED_STORE(path_state);
31 BRANCHED_STORE(throughput);
33 BRANCHED_STORE(isect);
37 for (
int i = 0; i <
kernel_split_sd(branched_state_sd, ray_index)->num_closure; i++) {
42 # undef BRANCHED_STORE
45 branched_state->next_closure = 0;
46 branched_state->next_sample = 0;
56 # define BRANCHED_RESTORE(name) kernel_split_state.name[ray_index] = branched_state->name;
58 BRANCHED_RESTORE(path_state);
59 BRANCHED_RESTORE(throughput);
60 BRANCHED_RESTORE(ray);
61 BRANCHED_RESTORE(isect);
65 for (
int i = 0; i <
kernel_split_sd(branched_state_sd, ray_index)->num_closure; i++) {
70 # undef BRANCHED_RESTORE
90 # define SPLIT_DATA_ENTRY(type, name, num) \
92 kernel_split_state.name[inactive_ray] = kernel_split_state.name[ray_index]; \
95 # undef SPLIT_DATA_ENTRY
126 float num_samples_adjust,
128 bool reset_path_state,
129 bool wait_for_shared)
135 float3 throughput = branched_state->throughput;
138 float sum_sample_weight = 0.0f;
139 # ifdef __DENOISING_FEATURES__
140 if (ps->denoising_feature_weight > 0.0f) {
141 for (
int i = 0; i < sd->num_closure; i++) {
149 sum_sample_weight += sc->sample_weight;
153 sum_sample_weight = 1.0f;
157 for (
int i = branched_state->next_closure; i < sd->num_closure; i++) {
169 num_samples =
kernel_data.integrator.diffuse_samples;
173 num_samples =
kernel_data.integrator.glossy_samples;
175 num_samples =
kernel_data.integrator.transmission_samples;
177 num_samples =
ceil_to_int(num_samples_adjust * num_samples);
179 float num_samples_inv = num_samples_adjust / num_samples;
181 for (
int j = branched_state->next_sample; j < num_samples; j++) {
182 if (reset_path_state) {
183 *ps = branched_state->path_state;
186 ps->rng_hash =
cmj_hash(branched_state->path_state.rng_hash, i);
193 if (!kernel_branched_path_surface_bounce(
194 kg, sd, sc, j, num_samples, tp, ps, &
L->state, bsdf_ray, sum_sample_weight)) {
198 ps->rng_hash = branched_state->path_state.rng_hash;
201 branched_state->next_closure = i;
202 branched_state->next_sample = j + 1;
205 *tp *= num_samples_inv;
207 if (kernel_split_branched_indirect_start_shared(
kg, ray_index)) {
214 branched_state->next_sample = 0;
217 branched_state->next_closure = sd->num_closure;
219 if (wait_for_shared) {
220 branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0);
221 if (branched_state->waiting_on_shared_samples) {
ccl_device_inline void path_radiance_init(KernelGlobals *kg, PathRadiance *L)
ccl_device_inline void path_radiance_copy_indirect(PathRadiance *L, const PathRadiance *L_src)
#define ccl_device_inline
#define ccl_device_noinline
#define CCL_NAMESPACE_END
ccl_device_inline uint cmj_hash(uint i, uint p)
ccl_device int dequeue_ray_index(int queue_number, ccl_global int *queues, int queue_size, ccl_global int *queue_index)
#define kernel_split_params
#define kernel_split_sd(sd, ray_index)
#define SPLIT_DATA_ENTRIES_BRANCHED_SHARED
#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 REMOVE_RAY_FLAG(ray_state, ray_index, flag)
#define ADD_RAY_FLAG(ray_state, ray_index, flag)
@ RAY_BRANCHED_INDIRECT_SHARED
#define CLOSURE_IS_BSDF_TRANSPARENT(type)
#define CLOSURE_IS_BSDF(type)
#define CLOSURE_IS_BSDF_GLOSSY(type)
#define CLOSURE_IS_BSDF_BSSRDF(type)
#define CLOSURE_IS_BSDF_DIFFUSE(type)
@ CLOSURE_BSDF_TRANSPARENT_ID
#define atomic_fetch_and_inc_uint32(p)
ccl_device_inline int ceil_to_int(float f)