diff options
Diffstat (limited to 'intern/cycles/kernel/split')
27 files changed, 0 insertions, 2946 deletions
diff --git a/intern/cycles/kernel/split/kernel_adaptive_adjust_samples.h b/intern/cycles/kernel/split/kernel_adaptive_adjust_samples.h deleted file mode 100644 index 437a5c9581b..00000000000 --- a/intern/cycles/kernel/split/kernel_adaptive_adjust_samples.h +++ /dev/null @@ -1,43 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_adaptive_adjust_samples(KernelGlobals *kg) -{ - int pixel_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (pixel_index < kernel_split_params.tile.w * kernel_split_params.tile.h) { - int x = kernel_split_params.tile.x + pixel_index % kernel_split_params.tile.w; - int y = kernel_split_params.tile.y + pixel_index / kernel_split_params.tile.w; - int buffer_offset = (kernel_split_params.tile.offset + x + - y * kernel_split_params.tile.stride) * - kernel_data.film.pass_stride; - ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; - int sample = kernel_split_params.tile.start_sample + kernel_split_params.tile.num_samples; - if (buffer[kernel_data.film.pass_sample_count] < 0.0f) { - buffer[kernel_data.film.pass_sample_count] = -buffer[kernel_data.film.pass_sample_count]; - float sample_multiplier = sample / buffer[kernel_data.film.pass_sample_count]; - if (sample_multiplier != 1.0f) { - kernel_adaptive_post_adjust(kg, buffer, sample_multiplier); - } - } - else { - kernel_adaptive_post_adjust(kg, buffer, sample / (sample - 1.0f)); - } - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_adaptive_filter_x.h b/intern/cycles/kernel/split/kernel_adaptive_filter_x.h deleted file mode 100644 index 93f41f7ced4..00000000000 --- a/intern/cycles/kernel/split/kernel_adaptive_filter_x.h +++ /dev/null @@ -1,30 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_adaptive_filter_x(KernelGlobals *kg) -{ - int pixel_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (pixel_index < kernel_split_params.tile.h && - kernel_split_params.tile.start_sample + kernel_split_params.tile.num_samples >= - kernel_data.integrator.adaptive_min_samples) { - int y = kernel_split_params.tile.y + pixel_index; - kernel_do_adaptive_filter_x(kg, y, &kernel_split_params.tile); - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_adaptive_filter_y.h b/intern/cycles/kernel/split/kernel_adaptive_filter_y.h deleted file mode 100644 index eca53d079ec..00000000000 --- a/intern/cycles/kernel/split/kernel_adaptive_filter_y.h +++ /dev/null @@ -1,29 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_adaptive_filter_y(KernelGlobals *kg) -{ - int pixel_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (pixel_index < kernel_split_params.tile.w && - kernel_split_params.tile.start_sample + kernel_split_params.tile.num_samples >= - kernel_data.integrator.adaptive_min_samples) { - int x = kernel_split_params.tile.x + pixel_index; - kernel_do_adaptive_filter_y(kg, x, &kernel_split_params.tile); - } -} -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_adaptive_stopping.h b/intern/cycles/kernel/split/kernel_adaptive_stopping.h deleted file mode 100644 index c8eb1ebd705..00000000000 --- a/intern/cycles/kernel/split/kernel_adaptive_stopping.h +++ /dev/null @@ -1,37 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_adaptive_stopping(KernelGlobals *kg) -{ - int pixel_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (pixel_index < kernel_split_params.tile.w * kernel_split_params.tile.h && - kernel_split_params.tile.start_sample + kernel_split_params.tile.num_samples >= - kernel_data.integrator.adaptive_min_samples) { - int x = kernel_split_params.tile.x + pixel_index % kernel_split_params.tile.w; - int y = kernel_split_params.tile.y + pixel_index / kernel_split_params.tile.w; - int buffer_offset = (kernel_split_params.tile.offset + x + - y * kernel_split_params.tile.stride) * - kernel_data.film.pass_stride; - ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; - kernel_do_adaptive_stopping(kg, - buffer, - kernel_split_params.tile.start_sample + - kernel_split_params.tile.num_samples - 1); - } -} -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_branched.h b/intern/cycles/kernel/split/kernel_branched.h deleted file mode 100644 index 45f5037d321..00000000000 --- a/intern/cycles/kernel/split/kernel_branched.h +++ /dev/null @@ -1,231 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -#ifdef __BRANCHED_PATH__ - -/* sets up the various state needed to do an indirect loop */ -ccl_device_inline void kernel_split_branched_path_indirect_loop_init(KernelGlobals *kg, - int ray_index) -{ - SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index]; - - /* save a copy of the state to restore later */ -# define BRANCHED_STORE(name) branched_state->name = kernel_split_state.name[ray_index]; - - BRANCHED_STORE(path_state); - BRANCHED_STORE(throughput); - BRANCHED_STORE(ray); - BRANCHED_STORE(isect); - BRANCHED_STORE(ray_state); - - *kernel_split_sd(branched_state_sd, ray_index) = *kernel_split_sd(sd, ray_index); - for (int i = 0; i < kernel_split_sd(branched_state_sd, ray_index)->num_closure; i++) { - kernel_split_sd(branched_state_sd, ray_index)->closure[i] = - kernel_split_sd(sd, ray_index)->closure[i]; - } - -# undef BRANCHED_STORE - - /* Set loop counters to initial position. */ - branched_state->next_closure = 0; - branched_state->next_sample = 0; -} - -/* ends an indirect loop and restores the previous state */ -ccl_device_inline void kernel_split_branched_path_indirect_loop_end(KernelGlobals *kg, - int ray_index) -{ - SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index]; - - /* restore state */ -# define BRANCHED_RESTORE(name) kernel_split_state.name[ray_index] = branched_state->name; - - BRANCHED_RESTORE(path_state); - BRANCHED_RESTORE(throughput); - BRANCHED_RESTORE(ray); - BRANCHED_RESTORE(isect); - BRANCHED_RESTORE(ray_state); - - *kernel_split_sd(sd, ray_index) = *kernel_split_sd(branched_state_sd, ray_index); - for (int i = 0; i < kernel_split_sd(branched_state_sd, ray_index)->num_closure; i++) { - kernel_split_sd(sd, ray_index)->closure[i] = - kernel_split_sd(branched_state_sd, ray_index)->closure[i]; - } - -# undef BRANCHED_RESTORE - - /* leave indirect loop */ - REMOVE_RAY_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_INDIRECT); -} - -ccl_device_inline bool kernel_split_branched_indirect_start_shared(KernelGlobals *kg, - int ray_index) -{ - ccl_global char *ray_state = kernel_split_state.ray_state; - - int inactive_ray = dequeue_ray_index(QUEUE_INACTIVE_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - kernel_split_params.queue_index); - - if (!IS_STATE(ray_state, inactive_ray, RAY_INACTIVE)) { - return false; - } - -# define SPLIT_DATA_ENTRY(type, name, num) \ - if (num) { \ - kernel_split_state.name[inactive_ray] = kernel_split_state.name[ray_index]; \ - } - SPLIT_DATA_ENTRIES_BRANCHED_SHARED -# undef SPLIT_DATA_ENTRY - - *kernel_split_sd(sd, inactive_ray) = *kernel_split_sd(sd, ray_index); - for (int i = 0; i < kernel_split_sd(sd, ray_index)->num_closure; i++) { - kernel_split_sd(sd, inactive_ray)->closure[i] = kernel_split_sd(sd, ray_index)->closure[i]; - } - - kernel_split_state.branched_state[inactive_ray].shared_sample_count = 0; - kernel_split_state.branched_state[inactive_ray].original_ray = ray_index; - kernel_split_state.branched_state[inactive_ray].waiting_on_shared_samples = false; - - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - PathRadiance *inactive_L = &kernel_split_state.path_radiance[inactive_ray]; - - path_radiance_init(kg, inactive_L); - path_radiance_copy_indirect(inactive_L, L); - - ray_state[inactive_ray] = RAY_REGENERATED; - ADD_RAY_FLAG(ray_state, inactive_ray, RAY_BRANCHED_INDIRECT_SHARED); - ADD_RAY_FLAG(ray_state, inactive_ray, IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT)); - - atomic_fetch_and_inc_uint32( - (ccl_global uint *)&kernel_split_state.branched_state[ray_index].shared_sample_count); - - return true; -} - -/* bounce off surface and integrate indirect light */ -ccl_device_noinline bool kernel_split_branched_path_surface_indirect_light_iter( - KernelGlobals *kg, - int ray_index, - float num_samples_adjust, - ShaderData *saved_sd, - bool reset_path_state, - bool wait_for_shared) -{ - SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index]; - - ShaderData *sd = saved_sd; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - float3 throughput = branched_state->throughput; - ccl_global PathState *ps = &kernel_split_state.path_state[ray_index]; - - float sum_sample_weight = 0.0f; -# ifdef __DENOISING_FEATURES__ - if (ps->denoising_feature_weight > 0.0f) { - for (int i = 0; i < sd->num_closure; i++) { - const ShaderClosure *sc = &sd->closure[i]; - - /* transparency is not handled here, but in outer loop */ - if (!CLOSURE_IS_BSDF(sc->type) || CLOSURE_IS_BSDF_TRANSPARENT(sc->type)) { - continue; - } - - sum_sample_weight += sc->sample_weight; - } - } - else { - sum_sample_weight = 1.0f; - } -# endif /* __DENOISING_FEATURES__ */ - - for (int i = branched_state->next_closure; i < sd->num_closure; i++) { - const ShaderClosure *sc = &sd->closure[i]; - - if (!CLOSURE_IS_BSDF(sc->type)) - continue; - /* transparency is not handled here, but in outer loop */ - if (sc->type == CLOSURE_BSDF_TRANSPARENT_ID) - continue; - - int num_samples; - - if (CLOSURE_IS_BSDF_DIFFUSE(sc->type)) - num_samples = kernel_data.integrator.diffuse_samples; - else if (CLOSURE_IS_BSDF_BSSRDF(sc->type)) - num_samples = 1; - else if (CLOSURE_IS_BSDF_GLOSSY(sc->type)) - num_samples = kernel_data.integrator.glossy_samples; - else - num_samples = kernel_data.integrator.transmission_samples; - - num_samples = ceil_to_int(num_samples_adjust * num_samples); - - float num_samples_inv = num_samples_adjust / num_samples; - - for (int j = branched_state->next_sample; j < num_samples; j++) { - if (reset_path_state) { - *ps = branched_state->path_state; - } - - ps->rng_hash = cmj_hash(branched_state->path_state.rng_hash, i); - - ccl_global float3 *tp = &kernel_split_state.throughput[ray_index]; - *tp = throughput; - - ccl_global Ray *bsdf_ray = &kernel_split_state.ray[ray_index]; - - if (!kernel_branched_path_surface_bounce( - kg, sd, sc, j, num_samples, tp, ps, &L->state, bsdf_ray, sum_sample_weight)) { - continue; - } - - ps->rng_hash = branched_state->path_state.rng_hash; - - /* update state for next iteration */ - branched_state->next_closure = i; - branched_state->next_sample = j + 1; - - /* start the indirect path */ - *tp *= num_samples_inv; - - if (kernel_split_branched_indirect_start_shared(kg, ray_index)) { - continue; - } - - return true; - } - - branched_state->next_sample = 0; - } - - branched_state->next_closure = sd->num_closure; - - if (wait_for_shared) { - branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0); - if (branched_state->waiting_on_shared_samples) { - return true; - } - } - - return false; -} - -#endif /* __BRANCHED_PATH__ */ - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_buffer_update.h b/intern/cycles/kernel/split/kernel_buffer_update.h deleted file mode 100644 index b96feca582f..00000000000 --- a/intern/cycles/kernel/split/kernel_buffer_update.h +++ /dev/null @@ -1,154 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel takes care of rays that hit the background (sceneintersect - * kernel), and for the rays of state RAY_UPDATE_BUFFER it updates the ray's - * accumulated radiance in the output buffer. This kernel also takes care of - * rays that have been determined to-be-regenerated. - * - * We will empty QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue in this kernel. - * - * Typically all rays that are in state RAY_HIT_BACKGROUND, RAY_UPDATE_BUFFER - * will be eventually set to RAY_TO_REGENERATE state in this kernel. - * Finally all rays of ray_state RAY_TO_REGENERATE will be regenerated and put - * in queue QUEUE_ACTIVE_AND_REGENERATED_RAYS. - * - * State of queues when this kernel is called: - * At entry, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE rays. - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with - * RAY_UPDATE_BUFFER, RAY_HIT_BACKGROUND, RAY_TO_REGENERATE rays. - * At exit, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE and - * RAY_REGENERATED rays. - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be empty. - */ -ccl_device void kernel_buffer_update(KernelGlobals *kg, - ccl_local_param unsigned int *local_queue_atomics) -{ - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (ray_index == 0) { - /* We will empty this queue in this kernel. */ - kernel_split_params.queue_index[QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS] = 0; - } - char enqueue_flag = 0; - ray_index = get_ray_index(kg, - ray_index, - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - - if (ray_index != QUEUE_EMPTY_SLOT) { - ccl_global char *ray_state = kernel_split_state.ray_state; - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index]; - bool ray_was_updated = false; - - if (IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER)) { - ray_was_updated = true; - uint sample = state->sample; - uint buffer_offset = kernel_split_state.buffer_offset[ray_index]; - ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; - - /* accumulate result in output buffer */ - kernel_write_result(kg, buffer, sample, L); - - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_TO_REGENERATE); - } - - if (kernel_data.film.cryptomatte_passes) { - /* Make sure no thread is writing to the buffers. */ - ccl_barrier(CCL_LOCAL_MEM_FENCE); - if (ray_was_updated && state->sample - 1 == kernel_data.integrator.aa_samples) { - uint buffer_offset = kernel_split_state.buffer_offset[ray_index]; - ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; - ccl_global float *cryptomatte_buffer = buffer + kernel_data.film.pass_cryptomatte; - kernel_sort_id_slots(cryptomatte_buffer, 2 * kernel_data.film.cryptomatte_depth); - } - } - - if (IS_STATE(ray_state, ray_index, RAY_TO_REGENERATE)) { - /* We have completed current work; So get next work */ - ccl_global uint *work_pools = kernel_split_params.work_pools; - uint total_work_size = kernel_split_params.total_work_size; - uint work_index; - - if (!get_next_work(kg, work_pools, total_work_size, ray_index, &work_index)) { - /* If work is invalid, this means no more work is available and the thread may exit */ - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_INACTIVE); - } - - if (IS_STATE(ray_state, ray_index, RAY_TO_REGENERATE)) { - ccl_global WorkTile *tile = &kernel_split_params.tile; - uint x, y, sample; - get_work_pixel(tile, work_index, &x, &y, &sample); - - /* Store buffer offset for writing to passes. */ - uint buffer_offset = (tile->offset + x + y * tile->stride) * kernel_data.film.pass_stride; - kernel_split_state.buffer_offset[ray_index] = buffer_offset; - - /* Initialize random numbers and ray. */ - uint rng_hash; - kernel_path_trace_setup(kg, sample, x, y, &rng_hash, ray); - - if (ray->t != 0.0f) { - /* Initialize throughput, path radiance, Ray, PathState; - * These rays proceed with path-iteration. - */ - *throughput = make_float3(1.0f, 1.0f, 1.0f); - path_radiance_init(kg, L); - path_state_init(kg, - AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]), - state, - rng_hash, - sample, - ray); -#ifdef __SUBSURFACE__ - kernel_path_subsurface_init_indirect(&kernel_split_state.ss_rays[ray_index]); -#endif - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - enqueue_flag = 1; - } - else { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_TO_REGENERATE); - } - } - } - } - - /* Enqueue RAY_REGENERATED rays into QUEUE_ACTIVE_AND_REGENERATED_RAYS; - * These rays will be made active during next SceneIntersectkernel. - */ - enqueue_ray_index_local(ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - enqueue_flag, - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h deleted file mode 100644 index 2f83a10316d..00000000000 --- a/intern/cycles/kernel/split/kernel_data_init.h +++ /dev/null @@ -1,115 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel Initializes structures needed in path-iteration kernels. - * - * Note on Queues: - * All slots in queues are initialized to queue empty slot; - * The number of elements in the queues is initialized to 0; - */ - -#ifndef __KERNEL_CPU__ -ccl_device void kernel_data_init( -#else -void KERNEL_FUNCTION_FULL_NAME(data_init)( -#endif - KernelGlobals *kg, - ccl_constant KernelData *data, - ccl_global void *split_data_buffer, - int num_elements, - ccl_global char *ray_state, - -#ifdef __KERNEL_OPENCL__ - KERNEL_BUFFER_PARAMS, -#endif - - int start_sample, - int end_sample, - int sx, - int sy, - int sw, - int sh, - int offset, - int stride, - ccl_global int *Queue_index, /* Tracks the number of elements in queues */ - int queuesize, /* size (capacity) of the queue */ - ccl_global char *use_queues_flag, /* flag to decide if scene-intersect kernel should use queues - to fetch ray index */ - ccl_global unsigned int *work_pools, /* Work pool for each work group */ - unsigned int num_samples, - ccl_global float *buffer) -{ -#ifdef KERNEL_STUB - STUB_ASSERT(KERNEL_ARCH, data_init); -#else - -# ifdef __KERNEL_OPENCL__ - kg->data = data; -# endif - - kernel_split_params.tile.x = sx; - kernel_split_params.tile.y = sy; - kernel_split_params.tile.w = sw; - kernel_split_params.tile.h = sh; - - kernel_split_params.tile.start_sample = start_sample; - kernel_split_params.tile.num_samples = num_samples; - - kernel_split_params.tile.offset = offset; - kernel_split_params.tile.stride = stride; - - kernel_split_params.tile.buffer = buffer; - - kernel_split_params.total_work_size = sw * sh * num_samples; - - kernel_split_params.work_pools = work_pools; - - kernel_split_params.queue_index = Queue_index; - kernel_split_params.queue_size = queuesize; - kernel_split_params.use_queues_flag = use_queues_flag; - - split_data_init(kg, &kernel_split_state, num_elements, split_data_buffer, ray_state); - -# ifdef __KERNEL_OPENCL__ - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); -# endif - - int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - - /* Initialize queue data and queue index. */ - if (thread_index < queuesize) { - for (int i = 0; i < NUM_QUEUES; i++) { - kernel_split_state.queue_data[i * queuesize + thread_index] = QUEUE_EMPTY_SLOT; - } - } - - if (thread_index == 0) { - for (int i = 0; i < NUM_QUEUES; i++) { - Queue_index[i] = 0; - } - - /* The scene-intersect kernel should not use the queues very first time. - * since the queue would be empty. - */ - *use_queues_flag = 0; - } -#endif /* KERENL_STUB */ -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_direct_lighting.h b/intern/cycles/kernel/split/kernel_direct_lighting.h deleted file mode 100644 index 3be2b35812f..00000000000 --- a/intern/cycles/kernel/split/kernel_direct_lighting.h +++ /dev/null @@ -1,152 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel takes care of direct lighting logic. - * However, the "shadow ray cast" part of direct lighting is handled - * in the next kernel. - * - * This kernels determines the rays for which a shadow_blocked() function - * associated with direct lighting should be executed. Those rays for which - * a shadow_blocked() function for direct-lighting must be executed, are - * marked with flag RAY_SHADOW_RAY_CAST_DL and enqueued into the queue - * QUEUE_SHADOW_RAY_CAST_DL_RAYS - * - * Note on Queues: - * This kernel only reads from the QUEUE_ACTIVE_AND_REGENERATED_RAYS queue - * and processes only the rays of state RAY_ACTIVE; If a ray needs to execute - * the corresponding shadow_blocked part, after direct lighting, the ray is - * marked with RAY_SHADOW_RAY_CAST_DL flag. - * - * State of queues when this kernel is called: - * - State of queues QUEUE_ACTIVE_AND_REGENERATED_RAYS and - * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be same before and after this - * kernel call. - * - QUEUE_SHADOW_RAY_CAST_DL_RAYS queue will be filled with rays for which a - * shadow_blocked function must be executed, after this kernel call - * Before this kernel call the QUEUE_SHADOW_RAY_CAST_DL_RAYS will be empty. - */ -ccl_device void kernel_direct_lighting(KernelGlobals *kg, - ccl_local_param unsigned int *local_queue_atomics) -{ - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - char enqueue_flag = 0; - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - - if (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) { - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - - /* direct lighting */ -#ifdef __EMISSION__ - bool flag = (kernel_data.integrator.use_direct_light && (sd->flag & SD_BSDF_HAS_EVAL)); - -# ifdef __BRANCHED_PATH__ - if (flag && kernel_data.integrator.branched) { - flag = false; - enqueue_flag = 1; - } -# endif /* __BRANCHED_PATH__ */ - -# ifdef __SHADOW_TRICKS__ - if (flag && state->flag & PATH_RAY_SHADOW_CATCHER) { - flag = false; - enqueue_flag = 1; - } -# endif /* __SHADOW_TRICKS__ */ - - if (flag) { - /* Sample illumination from lights to find path contribution. */ - float light_u, light_v; - path_state_rng_2D(kg, state, PRNG_LIGHT_U, &light_u, &light_v); - float terminate = path_state_rng_light_termination(kg, state); - - LightSample ls; - if (light_sample(kg, -1, light_u, light_v, sd->time, sd->P, state->bounce, &ls)) { - Ray light_ray; - light_ray.time = sd->time; - - BsdfEval L_light; - bool is_lamp; - if (direct_emission(kg, - sd, - AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]), - &ls, - state, - &light_ray, - &L_light, - &is_lamp, - terminate)) { - /* Write intermediate data to global memory to access from - * the next kernel. - */ - kernel_split_state.light_ray[ray_index] = light_ray; - kernel_split_state.bsdf_eval[ray_index] = L_light; - kernel_split_state.is_lamp[ray_index] = is_lamp; - /* Mark ray state for next shadow kernel. */ - enqueue_flag = 1; - } - } - } -#endif /* __EMISSION__ */ - } - -#ifdef __EMISSION__ - /* Enqueue RAY_SHADOW_RAY_CAST_DL rays. */ - enqueue_ray_index_local(ray_index, - QUEUE_SHADOW_RAY_CAST_DL_RAYS, - enqueue_flag, - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); -#endif - -#ifdef __BRANCHED_PATH__ - /* Enqueue RAY_LIGHT_INDIRECT_NEXT_ITER rays - * this is the last kernel before next_iteration_setup that uses local atomics so we do this here - */ - ccl_barrier(CCL_LOCAL_MEM_FENCE); - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - enqueue_ray_index_local( - ray_index, - QUEUE_LIGHT_INDIRECT_ITER, - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_LIGHT_INDIRECT_NEXT_ITER), - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); - -#endif /* __BRANCHED_PATH__ */ -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_do_volume.h b/intern/cycles/kernel/split/kernel_do_volume.h deleted file mode 100644 index 1775e870f07..00000000000 --- a/intern/cycles/kernel/split/kernel_do_volume.h +++ /dev/null @@ -1,227 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -#if defined(__BRANCHED_PATH__) && defined(__VOLUME__) - -ccl_device_inline void kernel_split_branched_path_volume_indirect_light_init(KernelGlobals *kg, - int ray_index) -{ - kernel_split_branched_path_indirect_loop_init(kg, ray_index); - - ADD_RAY_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_VOLUME_INDIRECT); -} - -ccl_device_noinline bool kernel_split_branched_path_volume_indirect_light_iter(KernelGlobals *kg, - int ray_index) -{ - SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index]; - - ShaderData *sd = kernel_split_sd(sd, ray_index); - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]); - - /* GPU: no decoupled ray marching, scatter probabilistically. */ - int num_samples = kernel_data.integrator.volume_samples; - float num_samples_inv = 1.0f / num_samples; - - Ray volume_ray = branched_state->ray; - volume_ray.t = (!IS_STATE(&branched_state->ray_state, 0, RAY_HIT_BACKGROUND)) ? - branched_state->isect.t : - FLT_MAX; - - float step_size = volume_stack_step_size(kg, branched_state->path_state.volume_stack); - - for (int j = branched_state->next_sample; j < num_samples; j++) { - ccl_global PathState *ps = &kernel_split_state.path_state[ray_index]; - *ps = branched_state->path_state; - - ccl_global Ray *pray = &kernel_split_state.ray[ray_index]; - *pray = branched_state->ray; - - ccl_global float3 *tp = &kernel_split_state.throughput[ray_index]; - *tp = branched_state->throughput * num_samples_inv; - - /* branch RNG state */ - path_state_branch(ps, j, num_samples); - - /* integrate along volume segment with distance sampling */ - VolumeIntegrateResult result = kernel_volume_integrate( - kg, ps, sd, &volume_ray, L, tp, step_size); - -# ifdef __VOLUME_SCATTER__ - if (result == VOLUME_PATH_SCATTERED) { - /* direct lighting */ - kernel_path_volume_connect_light(kg, sd, emission_sd, *tp, &branched_state->path_state, L); - - /* indirect light bounce */ - if (!kernel_path_volume_bounce(kg, sd, tp, ps, &L->state, pray)) { - continue; - } - - /* start the indirect path */ - branched_state->next_closure = 0; - branched_state->next_sample = j + 1; - - /* Attempting to share too many samples is slow for volumes as it causes us to - * loop here more and have many calls to kernel_volume_integrate which evaluates - * shaders. The many expensive shader evaluations cause the work load to become - * unbalanced and many threads to become idle in this kernel. Limiting the - * number of shared samples here helps quite a lot. - */ - if (branched_state->shared_sample_count < 2) { - if (kernel_split_branched_indirect_start_shared(kg, ray_index)) { - continue; - } - } - - return true; - } -# endif - } - - branched_state->next_sample = num_samples; - - branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0); - if (branched_state->waiting_on_shared_samples) { - return true; - } - - kernel_split_branched_path_indirect_loop_end(kg, ray_index); - - /* todo: avoid this calculation using decoupled ray marching */ - float3 throughput = kernel_split_state.throughput[ray_index]; - kernel_volume_shadow( - kg, emission_sd, &kernel_split_state.path_state[ray_index], &volume_ray, &throughput); - kernel_split_state.throughput[ray_index] = throughput; - - return false; -} - -#endif /* __BRANCHED_PATH__ && __VOLUME__ */ - -ccl_device void kernel_do_volume(KernelGlobals *kg) -{ -#ifdef __VOLUME__ - /* We will empty this queue in this kernel. */ - if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) { - kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0; -# ifdef __BRANCHED_PATH__ - kernel_split_params.queue_index[QUEUE_VOLUME_INDIRECT_ITER] = 0; -# endif /* __BRANCHED_PATH__ */ - } - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - - if (*kernel_split_params.use_queues_flag) { - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - } - - ccl_global char *ray_state = kernel_split_state.ray_state; - - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE) || - IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) { - ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index]; - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - ccl_global Intersection *isect = &kernel_split_state.isect[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]); - - bool hit = !IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND); - - /* Sanitize volume stack. */ - if (!hit) { - kernel_volume_clean_stack(kg, state->volume_stack); - } - /* volume attenuation, emission, scatter */ - if (state->volume_stack[0].shader != SHADER_NONE) { - Ray volume_ray = *ray; - volume_ray.t = (hit) ? isect->t : FLT_MAX; - -# ifdef __BRANCHED_PATH__ - if (!kernel_data.integrator.branched || - IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT)) { -# endif /* __BRANCHED_PATH__ */ - float step_size = volume_stack_step_size(kg, state->volume_stack); - - { - /* integrate along volume segment with distance sampling */ - VolumeIntegrateResult result = kernel_volume_integrate( - kg, state, sd, &volume_ray, L, throughput, step_size); - -# ifdef __VOLUME_SCATTER__ - if (result == VOLUME_PATH_SCATTERED) { - /* direct lighting */ - kernel_path_volume_connect_light(kg, sd, emission_sd, *throughput, state, L); - - /* indirect light bounce */ - if (kernel_path_volume_bounce(kg, sd, throughput, state, &L->state, ray)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - else { - kernel_split_path_end(kg, ray_index); - } - } -# endif /* __VOLUME_SCATTER__ */ - } - -# ifdef __BRANCHED_PATH__ - } - else { - kernel_split_branched_path_volume_indirect_light_init(kg, ray_index); - - if (kernel_split_branched_path_volume_indirect_light_iter(kg, ray_index)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - } -# endif /* __BRANCHED_PATH__ */ - } - } - -# ifdef __BRANCHED_PATH__ - /* iter loop */ - ray_index = get_ray_index(kg, - ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0), - QUEUE_VOLUME_INDIRECT_ITER, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - - if (IS_STATE(ray_state, ray_index, RAY_VOLUME_INDIRECT_NEXT_ITER)) { - /* for render passes, sum and reset indirect light pass variables - * for the next samples */ - path_radiance_sum_indirect(&kernel_split_state.path_radiance[ray_index]); - path_radiance_reset_indirect(&kernel_split_state.path_radiance[ray_index]); - - if (kernel_split_branched_path_volume_indirect_light_iter(kg, ray_index)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - } -# endif /* __BRANCHED_PATH__ */ - -#endif /* __VOLUME__ */ -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_enqueue_inactive.h b/intern/cycles/kernel/split/kernel_enqueue_inactive.h deleted file mode 100644 index 745313f89f1..00000000000 --- a/intern/cycles/kernel/split/kernel_enqueue_inactive.h +++ /dev/null @@ -1,46 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_enqueue_inactive(KernelGlobals *kg, - ccl_local_param unsigned int *local_queue_atomics) -{ -#ifdef __BRANCHED_PATH__ - /* Enqueue RAY_INACTIVE rays into QUEUE_INACTIVE_RAYS queue. */ - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - - char enqueue_flag = 0; - if (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_INACTIVE)) { - enqueue_flag = 1; - } - - enqueue_ray_index_local(ray_index, - QUEUE_INACTIVE_RAYS, - enqueue_flag, - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); -#endif /* __BRANCHED_PATH__ */ -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h b/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h deleted file mode 100644 index 61722840b0b..00000000000 --- a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h +++ /dev/null @@ -1,149 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel takes care of the logic to process "material of type holdout", - * indirect primitive emission, bsdf blurring, probabilistic path termination - * and AO. - * - * This kernels determines the rays for which a shadow_blocked() function - * associated with AO should be executed. Those rays for which a - * shadow_blocked() function for AO must be executed are marked with flag - * RAY_SHADOW_RAY_CAST_ao and enqueued into the queue - * QUEUE_SHADOW_RAY_CAST_AO_RAYS - * - * Ray state of rays that are terminated in this kernel are changed to RAY_UPDATE_BUFFER - * - * Note on Queues: - * This kernel fetches rays from the queue QUEUE_ACTIVE_AND_REGENERATED_RAYS - * and processes only the rays of state RAY_ACTIVE. - * There are different points in this kernel where a ray may terminate and - * reach RAY_UPDATE_BUFFER state. These rays are enqueued into - * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue. These rays will still be present - * in QUEUE_ACTIVE_AND_REGENERATED_RAYS queue, but since their ray-state has - * been changed to RAY_UPDATE_BUFFER, there is no problem. - * - * State of queues when this kernel is called: - * At entry, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE and - * RAY_REGENERATED rays - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with - * RAY_TO_REGENERATE rays. - * - QUEUE_SHADOW_RAY_CAST_AO_RAYS will be empty. - * At exit, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE, - * RAY_REGENERATED and RAY_UPDATE_BUFFER rays. - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with - * RAY_TO_REGENERATE and RAY_UPDATE_BUFFER rays. - * - QUEUE_SHADOW_RAY_CAST_AO_RAYS will be filled with rays marked with - * flag RAY_SHADOW_RAY_CAST_AO - */ - -ccl_device void kernel_holdout_emission_blurring_pathtermination_ao( - KernelGlobals *kg, ccl_local_param BackgroundAOLocals *locals) -{ - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - locals->queue_atomics_bg = 0; - locals->queue_atomics_ao = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - -#ifdef __AO__ - char enqueue_flag = 0; -#endif - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - - if (ray_index != QUEUE_EMPTY_SLOT) { - ccl_global PathState *state = 0x0; - float3 throughput; - - ccl_global char *ray_state = kernel_split_state.ray_state; - ShaderData *sd = kernel_split_sd(sd, ray_index); - - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - uint buffer_offset = kernel_split_state.buffer_offset[ray_index]; - ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; - - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]); - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - - throughput = kernel_split_state.throughput[ray_index]; - state = &kernel_split_state.path_state[ray_index]; - - if (!kernel_path_shader_apply(kg, sd, state, ray, throughput, emission_sd, L, buffer)) { - kernel_split_path_end(kg, ray_index); - } - } - - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - /* Path termination. this is a strange place to put the termination, it's - * mainly due to the mixed in MIS that we use. gives too many unneeded - * shader evaluations, only need emission if we are going to terminate. - */ - float probability = path_state_continuation_probability(kg, state, throughput); - - if (probability == 0.0f) { - kernel_split_path_end(kg, ray_index); - } - else if (probability < 1.0f) { - float terminate = path_state_rng_1D(kg, state, PRNG_TERMINATE); - if (terminate >= probability) { - kernel_split_path_end(kg, ray_index); - } - else { - kernel_split_state.throughput[ray_index] = throughput / probability; - } - } - -#ifdef __DENOISING_FEATURES__ - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - kernel_update_denoising_features(kg, sd, state, L); - } -#endif - } - -#ifdef __AO__ - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - /* ambient occlusion */ - if (kernel_data.integrator.use_ambient_occlusion) { - enqueue_flag = 1; - } - } -#endif /* __AO__ */ - } - -#ifdef __AO__ - /* Enqueue to-shadow-ray-cast rays. */ - enqueue_ray_index_local(ray_index, - QUEUE_SHADOW_RAY_CAST_AO_RAYS, - enqueue_flag, - kernel_split_params.queue_size, - &locals->queue_atomics_ao, - kernel_split_state.queue_data, - kernel_split_params.queue_index); -#endif -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_indirect_background.h b/intern/cycles/kernel/split/kernel_indirect_background.h deleted file mode 100644 index 6d500650cc0..00000000000 --- a/intern/cycles/kernel/split/kernel_indirect_background.h +++ /dev/null @@ -1,69 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_indirect_background(KernelGlobals *kg) -{ - ccl_global char *ray_state = kernel_split_state.ray_state; - - int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - int ray_index; - - if (kernel_data.integrator.ao_bounces != INT_MAX) { - ray_index = get_ray_index(kg, - thread_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - - if (ray_index != QUEUE_EMPTY_SLOT) { - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - if (path_state_ao_bounce(kg, state)) { - kernel_split_path_end(kg, ray_index); - } - } - } - } - - ray_index = get_ray_index(kg, - thread_index, - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - - if (ray_index == QUEUE_EMPTY_SLOT) { - return; - } - - if (IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) { - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - float3 throughput = kernel_split_state.throughput[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - uint buffer_offset = kernel_split_state.buffer_offset[ray_index]; - ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; - - kernel_path_background(kg, state, ray, throughput, sd, buffer, L); - kernel_split_path_end(kg, ray_index); - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_indirect_subsurface.h b/intern/cycles/kernel/split/kernel_indirect_subsurface.h deleted file mode 100644 index 3f48f8d6f56..00000000000 --- a/intern/cycles/kernel/split/kernel_indirect_subsurface.h +++ /dev/null @@ -1,67 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_indirect_subsurface(KernelGlobals *kg) -{ - int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (thread_index == 0) { - /* We will empty both queues in this kernel. */ - kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0; - kernel_split_params.queue_index[QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS] = 0; - } - - int ray_index; - get_ray_index(kg, - thread_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - ray_index = get_ray_index(kg, - thread_index, - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - -#ifdef __SUBSURFACE__ - if (ray_index == QUEUE_EMPTY_SLOT) { - return; - } - - ccl_global char *ray_state = kernel_split_state.ray_state; - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index]; - - if (IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER)) { - ccl_addr_space SubsurfaceIndirectRays *ss_indirect = &kernel_split_state.ss_rays[ray_index]; - - /* Trace indirect subsurface rays by restarting the loop. this uses less - * stack memory than invoking kernel_path_indirect. - */ - if (ss_indirect->num_rays) { - kernel_path_subsurface_setup_indirect(kg, ss_indirect, state, ray, L, throughput); - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - } -#endif /* __SUBSURFACE__ */ -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_lamp_emission.h b/intern/cycles/kernel/split/kernel_lamp_emission.h deleted file mode 100644 index 7ecb099208d..00000000000 --- a/intern/cycles/kernel/split/kernel_lamp_emission.h +++ /dev/null @@ -1,67 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel operates on QUEUE_ACTIVE_AND_REGENERATED_RAYS. - * It processes rays of state RAY_ACTIVE and RAY_HIT_BACKGROUND. - * We will empty QUEUE_ACTIVE_AND_REGENERATED_RAYS queue in this kernel. - */ -ccl_device void kernel_lamp_emission(KernelGlobals *kg) -{ -#ifndef __VOLUME__ - /* We will empty this queue in this kernel. */ - if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) { - kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0; - } -#endif - /* Fetch use_queues_flag. */ - char local_use_queues_flag = *kernel_split_params.use_queues_flag; - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (local_use_queues_flag) { - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, -#ifndef __VOLUME__ - 1 -#else - 0 -#endif - ); - if (ray_index == QUEUE_EMPTY_SLOT) { - return; - } - } - - if (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE) || - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_HIT_BACKGROUND)) { - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - - float3 throughput = kernel_split_state.throughput[ray_index]; - Ray ray = kernel_split_state.ray[ray_index]; - ccl_global Intersection *isect = &kernel_split_state.isect[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - - kernel_path_lamp_emission(kg, state, &ray, throughput, isect, sd, L); - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_next_iteration_setup.h b/intern/cycles/kernel/split/kernel_next_iteration_setup.h deleted file mode 100644 index 320f6a414bf..00000000000 --- a/intern/cycles/kernel/split/kernel_next_iteration_setup.h +++ /dev/null @@ -1,258 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/*This kernel takes care of setting up ray for the next iteration of - * path-iteration and accumulating radiance corresponding to AO and - * direct-lighting - * - * Ray state of rays that are terminated in this kernel are changed - * to RAY_UPDATE_BUFFER. - * - * Note on queues: - * This kernel fetches rays from the queue QUEUE_ACTIVE_AND_REGENERATED_RAYS - * and processes only the rays of state RAY_ACTIVE. - * There are different points in this kernel where a ray may terminate and - * reach RAY_UPDATE_BUFF state. These rays are enqueued into - * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue. These rays will still be present - * in QUEUE_ACTIVE_AND_REGENERATED_RAYS queue, but since their ray-state has - * been changed to RAY_UPDATE_BUFF, there is no problem. - * - * State of queues when this kernel is called: - * At entry, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE, - * RAY_REGENERATED, RAY_UPDATE_BUFFER rays. - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with - * RAY_TO_REGENERATE and RAY_UPDATE_BUFFER rays. - * At exit, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE, - * RAY_REGENERATED and more RAY_UPDATE_BUFFER rays. - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with - * RAY_TO_REGENERATE and more RAY_UPDATE_BUFFER rays. - */ - -#ifdef __BRANCHED_PATH__ -ccl_device_inline void kernel_split_branched_indirect_light_init(KernelGlobals *kg, int ray_index) -{ - kernel_split_branched_path_indirect_loop_init(kg, ray_index); - - ADD_RAY_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_LIGHT_INDIRECT); -} - -ccl_device void kernel_split_branched_transparent_bounce(KernelGlobals *kg, int ray_index) -{ - ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - -# ifdef __VOLUME__ - if (!(sd->flag & SD_HAS_ONLY_VOLUME)) { -# endif - /* continue in case of transparency */ - *throughput *= shader_bsdf_transparency(kg, sd); - - if (is_zero(*throughput)) { - kernel_split_path_end(kg, ray_index); - return; - } - - /* Update Path State */ - path_state_next(kg, state, LABEL_TRANSPARENT); -# ifdef __VOLUME__ - } - else { - if (!path_state_volume_next(kg, state)) { - kernel_split_path_end(kg, ray_index); - return; - } - } -# endif - - ray->P = ray_offset(sd->P, -sd->Ng); - ray->t -= sd->ray_length; /* clipping works through transparent */ - -# ifdef __RAY_DIFFERENTIALS__ - ray->dP = sd->dP; - ray->dD.dx = -sd->dI.dx; - ray->dD.dy = -sd->dI.dy; -# endif /* __RAY_DIFFERENTIALS__ */ - -# ifdef __VOLUME__ - /* enter/exit volume */ - kernel_volume_stack_enter_exit(kg, sd, state->volume_stack); -# endif /* __VOLUME__ */ -} -#endif /* __BRANCHED_PATH__ */ - -ccl_device void kernel_next_iteration_setup(KernelGlobals *kg, - ccl_local_param unsigned int *local_queue_atomics) -{ - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) { - /* If we are here, then it means that scene-intersect kernel - * has already been executed at least once. From the next time, - * scene-intersect kernel may operate on queues to fetch ray index - */ - *kernel_split_params.use_queues_flag = 1; - - /* Mark queue indices of QUEUE_SHADOW_RAY_CAST_AO_RAYS and - * QUEUE_SHADOW_RAY_CAST_DL_RAYS queues that were made empty during the - * previous kernel. - */ - kernel_split_params.queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS] = 0; - kernel_split_params.queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS] = 0; - } - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - - ccl_global char *ray_state = kernel_split_state.ray_state; - -#ifdef __VOLUME__ - /* Reactivate only volume rays here, most surface work was skipped. */ - if (IS_STATE(ray_state, ray_index, RAY_HAS_ONLY_VOLUME)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_ACTIVE); - } -#endif - - bool active = IS_STATE(ray_state, ray_index, RAY_ACTIVE); - if (active) { - ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index]; - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - -#ifdef __BRANCHED_PATH__ - if (!kernel_data.integrator.branched || IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT)) { -#endif - /* Compute direct lighting and next bounce. */ - if (!kernel_path_surface_bounce(kg, sd, throughput, state, &L->state, ray)) { - kernel_split_path_end(kg, ray_index); - } -#ifdef __BRANCHED_PATH__ - } - else if (sd->flag & SD_HAS_ONLY_VOLUME) { - kernel_split_branched_transparent_bounce(kg, ray_index); - } - else { - kernel_split_branched_indirect_light_init(kg, ray_index); - - if (kernel_split_branched_path_surface_indirect_light_iter( - kg, ray_index, 1.0f, kernel_split_sd(branched_state_sd, ray_index), true, true)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - else { - kernel_split_branched_path_indirect_loop_end(kg, ray_index); - kernel_split_branched_transparent_bounce(kg, ray_index); - } - } -#endif /* __BRANCHED_PATH__ */ - } - - /* Enqueue RAY_UPDATE_BUFFER rays. */ - enqueue_ray_index_local(ray_index, - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS, - IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER) && active, - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); - -#ifdef __BRANCHED_PATH__ - /* iter loop */ - if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) { - kernel_split_params.queue_index[QUEUE_LIGHT_INDIRECT_ITER] = 0; - } - - ray_index = get_ray_index(kg, - ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0), - QUEUE_LIGHT_INDIRECT_ITER, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - - if (IS_STATE(ray_state, ray_index, RAY_LIGHT_INDIRECT_NEXT_ITER)) { - /* for render passes, sum and reset indirect light pass variables - * for the next samples */ - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - - path_radiance_sum_indirect(L); - path_radiance_reset_indirect(L); - - if (kernel_split_branched_path_surface_indirect_light_iter( - kg, ray_index, 1.0f, kernel_split_sd(branched_state_sd, ray_index), true, true)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - else { - kernel_split_branched_path_indirect_loop_end(kg, ray_index); - kernel_split_branched_transparent_bounce(kg, ray_index); - } - } - -# ifdef __VOLUME__ - /* Enqueue RAY_VOLUME_INDIRECT_NEXT_ITER rays */ - ccl_barrier(CCL_LOCAL_MEM_FENCE); - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - enqueue_ray_index_local( - ray_index, - QUEUE_VOLUME_INDIRECT_ITER, - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_VOLUME_INDIRECT_NEXT_ITER), - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); - -# endif /* __VOLUME__ */ - -# ifdef __SUBSURFACE__ - /* Enqueue RAY_SUBSURFACE_INDIRECT_NEXT_ITER rays */ - ccl_barrier(CCL_LOCAL_MEM_FENCE); - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - enqueue_ray_index_local( - ray_index, - QUEUE_SUBSURFACE_INDIRECT_ITER, - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_SUBSURFACE_INDIRECT_NEXT_ITER), - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); -# endif /* __SUBSURFACE__ */ -#endif /* __BRANCHED_PATH__ */ -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_path_init.h b/intern/cycles/kernel/split/kernel_path_init.h deleted file mode 100644 index c686f46a0cd..00000000000 --- a/intern/cycles/kernel/split/kernel_path_init.h +++ /dev/null @@ -1,78 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel initializes structures needed in path-iteration kernels. - * This is the first kernel in ray-tracing logic. - * - * Ray state of rays outside the tile-boundary will be marked RAY_INACTIVE - */ -ccl_device void kernel_path_init(KernelGlobals *kg) -{ - int ray_index = ccl_global_id(0) + ccl_global_id(1) * ccl_global_size(0); - - /* This is the first assignment to ray_state; - * So we don't use ASSIGN_RAY_STATE macro. - */ - kernel_split_state.ray_state[ray_index] = RAY_ACTIVE; - - /* Get work. */ - ccl_global uint *work_pools = kernel_split_params.work_pools; - uint total_work_size = kernel_split_params.total_work_size; - uint work_index; - - if (!get_next_work(kg, work_pools, total_work_size, ray_index, &work_index)) { - /* No more work, mark ray as inactive */ - kernel_split_state.ray_state[ray_index] = RAY_INACTIVE; - - return; - } - - ccl_global WorkTile *tile = &kernel_split_params.tile; - uint x, y, sample; - get_work_pixel(tile, work_index, &x, &y, &sample); - - /* Store buffer offset for writing to passes. */ - uint buffer_offset = (tile->offset + x + y * tile->stride) * kernel_data.film.pass_stride; - kernel_split_state.buffer_offset[ray_index] = buffer_offset; - - /* Initialize random numbers and ray. */ - uint rng_hash; - kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &kernel_split_state.ray[ray_index]); - - if (kernel_split_state.ray[ray_index].t != 0.0f) { - /* Initialize throughput, path radiance, Ray, PathState; - * These rays proceed with path-iteration. - */ - kernel_split_state.throughput[ray_index] = make_float3(1.0f, 1.0f, 1.0f); - path_radiance_init(kg, &kernel_split_state.path_radiance[ray_index]); - path_state_init(kg, - AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]), - &kernel_split_state.path_state[ray_index], - rng_hash, - sample, - &kernel_split_state.ray[ray_index]); -#ifdef __SUBSURFACE__ - kernel_path_subsurface_init_indirect(&kernel_split_state.ss_rays[ray_index]); -#endif - } - else { - ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_TO_REGENERATE); - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_queue_enqueue.h b/intern/cycles/kernel/split/kernel_queue_enqueue.h deleted file mode 100644 index 2db87f7a671..00000000000 --- a/intern/cycles/kernel/split/kernel_queue_enqueue.h +++ /dev/null @@ -1,87 +0,0 @@ -/* - * Copyright 2011-2016 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel enqueues rays of different ray state into their - * appropriate queues: - * - * 1. Rays that have been determined to hit the background from the - * "kernel_scene_intersect" kernel are enqueued in - * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS; - * 2. Rays that have been determined to be actively participating in pat - * -iteration will be enqueued into QUEUE_ACTIVE_AND_REGENERATED_RAYS. - * - * State of queue during other times this kernel is called: - * At entry, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be empty. - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will contain RAY_TO_REGENERATE - * and RAY_UPDATE_BUFFER rays. - * At exit, - * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE rays. - * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with - * RAY_TO_REGENERATE, RAY_UPDATE_BUFFER, RAY_HIT_BACKGROUND rays. - */ -ccl_device void kernel_queue_enqueue(KernelGlobals *kg, ccl_local_param QueueEnqueueLocals *locals) -{ - /* We have only 2 cases (Hit/Not-Hit) */ - int lidx = ccl_local_id(1) * ccl_local_size(0) + ccl_local_id(0); - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - - if (lidx == 0) { - locals->queue_atomics[0] = 0; - locals->queue_atomics[1] = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int queue_number = -1; - - if (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_HIT_BACKGROUND) || - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_UPDATE_BUFFER) || - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_TO_REGENERATE)) { - queue_number = QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS; - } - else if (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE) || - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_HAS_ONLY_VOLUME) || - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED)) { - queue_number = QUEUE_ACTIVE_AND_REGENERATED_RAYS; - } - - unsigned int my_lqidx; - if (queue_number != -1) { - my_lqidx = get_local_queue_index(queue_number, locals->queue_atomics); - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - if (lidx == 0) { - locals->queue_atomics[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = get_global_per_queue_offset( - QUEUE_ACTIVE_AND_REGENERATED_RAYS, locals->queue_atomics, kernel_split_params.queue_index); - locals->queue_atomics[QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS] = get_global_per_queue_offset( - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS, - locals->queue_atomics, - kernel_split_params.queue_index); - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - unsigned int my_gqidx; - if (queue_number != -1) { - my_gqidx = get_global_queue_index( - queue_number, kernel_split_params.queue_size, my_lqidx, locals->queue_atomics); - kernel_split_state.queue_data[my_gqidx] = ray_index; - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_scene_intersect.h b/intern/cycles/kernel/split/kernel_scene_intersect.h deleted file mode 100644 index 9ac95aafd2f..00000000000 --- a/intern/cycles/kernel/split/kernel_scene_intersect.h +++ /dev/null @@ -1,83 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel takes care of scene_intersect function. - * - * This kernel changes the ray_state of RAY_REGENERATED rays to RAY_ACTIVE. - * This kernel processes rays of ray state RAY_ACTIVE - * This kernel determines the rays that have hit the background and changes - * their ray state to RAY_HIT_BACKGROUND. - */ -ccl_device void kernel_scene_intersect(KernelGlobals *kg) -{ - /* Fetch use_queues_flag */ - char local_use_queues_flag = *kernel_split_params.use_queues_flag; - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (local_use_queues_flag) { - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - - if (ray_index == QUEUE_EMPTY_SLOT) { - return; - } - } - - /* All regenerated rays become active here */ - if (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_REGENERATED)) { -#ifdef __BRANCHED_PATH__ - if (kernel_split_state.branched_state[ray_index].waiting_on_shared_samples) { - kernel_split_path_end(kg, ray_index); - } - else -#endif /* __BRANCHED_PATH__ */ - { - ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE); - } - } - - if (!IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) { - return; - } - - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - Ray ray = kernel_split_state.ray[ray_index]; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - - Intersection isect; - const int last_object = state->bounce > 0 ? - intersection_get_object(kg, &kernel_split_state.isect[ray_index]) : - OBJECT_NONE; - bool hit = kernel_path_scene_intersect(kg, state, &ray, &isect, L, last_object); - kernel_split_state.isect[ray_index] = isect; - - if (!hit) { - /* Change the state of rays that hit the background; - * These rays undergo special processing in the - * background_bufferUpdate kernel. - */ - ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_HIT_BACKGROUND); - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_shader_eval.h b/intern/cycles/kernel/split/kernel_shader_eval.h deleted file mode 100644 index c760a2b2049..00000000000 --- a/intern/cycles/kernel/split/kernel_shader_eval.h +++ /dev/null @@ -1,69 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel evaluates ShaderData structure from the values computed - * by the previous kernels. - */ -ccl_device void kernel_shader_eval(KernelGlobals *kg) -{ - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - /* Sorting on cuda split is not implemented */ -#ifdef __KERNEL_CUDA__ - int queue_index = kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS]; -#else - int queue_index = kernel_split_params.queue_index[QUEUE_SHADER_SORTED_RAYS]; -#endif - if (ray_index >= queue_index) { - return; - } - ray_index = get_ray_index(kg, - ray_index, -#ifdef __KERNEL_CUDA__ - QUEUE_ACTIVE_AND_REGENERATED_RAYS, -#else - QUEUE_SHADER_SORTED_RAYS, -#endif - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - - if (ray_index == QUEUE_EMPTY_SLOT) { - return; - } - - ccl_global char *ray_state = kernel_split_state.ray_state; - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - uint buffer_offset = kernel_split_state.buffer_offset[ray_index]; - ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset; - - shader_eval_surface(kg, kernel_split_sd(sd, ray_index), state, buffer, state->flag); -#ifdef __BRANCHED_PATH__ - if (kernel_data.integrator.branched) { - shader_merge_closures(kernel_split_sd(sd, ray_index)); - } - else -#endif - { - shader_prepare_closures(kernel_split_sd(sd, ray_index), state); - } - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_shader_setup.h b/intern/cycles/kernel/split/kernel_shader_setup.h deleted file mode 100644 index 551836d1653..00000000000 --- a/intern/cycles/kernel/split/kernel_shader_setup.h +++ /dev/null @@ -1,74 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* This kernel sets up the ShaderData structure from the values computed - * by the previous kernels. - * - * It also identifies the rays of state RAY_TO_REGENERATE and enqueues them - * in QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue. - */ -ccl_device void kernel_shader_setup(KernelGlobals *kg, - ccl_local_param unsigned int *local_queue_atomics) -{ - /* Enqueue RAY_TO_REGENERATE rays into QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue. */ - if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) { - *local_queue_atomics = 0; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - int queue_index = kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS]; - if (ray_index < queue_index) { - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 0); - } - else { - ray_index = QUEUE_EMPTY_SLOT; - } - - char enqueue_flag = (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_TO_REGENERATE)) ? 1 : - 0; - enqueue_ray_index_local(ray_index, - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS, - enqueue_flag, - kernel_split_params.queue_size, - local_queue_atomics, - kernel_split_state.queue_data, - kernel_split_params.queue_index); - - /* Continue on with shader evaluation. */ - if (IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE)) { - Intersection isect = kernel_split_state.isect[ray_index]; - Ray ray = kernel_split_state.ray[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - - shader_setup_from_ray(kg, sd, &isect, &ray); - -#ifdef __VOLUME__ - if (sd->flag & SD_HAS_ONLY_VOLUME) { - ASSIGN_RAY_STATE(kernel_split_state.ray_state, ray_index, RAY_HAS_ONLY_VOLUME); - } -#endif - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_shader_sort.h b/intern/cycles/kernel/split/kernel_shader_sort.h deleted file mode 100644 index 95d33a42014..00000000000 --- a/intern/cycles/kernel/split/kernel_shader_sort.h +++ /dev/null @@ -1,97 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -ccl_device void kernel_shader_sort(KernelGlobals *kg, ccl_local_param ShaderSortLocals *locals) -{ -#ifndef __KERNEL_CUDA__ - int tid = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - uint qsize = kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS]; - if (tid == 0) { - kernel_split_params.queue_index[QUEUE_SHADER_SORTED_RAYS] = qsize; - } - - uint offset = (tid / SHADER_SORT_LOCAL_SIZE) * SHADER_SORT_BLOCK_SIZE; - if (offset >= qsize) { - return; - } - - int lid = ccl_local_id(1) * ccl_local_size(0) + ccl_local_id(0); - uint input = QUEUE_ACTIVE_AND_REGENERATED_RAYS * (kernel_split_params.queue_size); - uint output = QUEUE_SHADER_SORTED_RAYS * (kernel_split_params.queue_size); - ccl_local uint *local_value = &locals->local_value[0]; - ccl_local ushort *local_index = &locals->local_index[0]; - - /* copy to local memory */ - for (uint i = 0; i < SHADER_SORT_BLOCK_SIZE; i += SHADER_SORT_LOCAL_SIZE) { - uint idx = offset + i + lid; - uint add = input + idx; - uint value = (~0); - if (idx < qsize) { - int ray_index = kernel_split_state.queue_data[add]; - bool valid = (ray_index != QUEUE_EMPTY_SLOT) && - IS_STATE(kernel_split_state.ray_state, ray_index, RAY_ACTIVE); - if (valid) { - value = kernel_split_sd(sd, ray_index)->shader & SHADER_MASK; - } - } - local_value[i + lid] = value; - local_index[i + lid] = i + lid; - } - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - /* skip sorting for cpu split kernel */ -# ifdef __KERNEL_OPENCL__ - - /* bitonic sort */ - for (uint length = 1; length < SHADER_SORT_BLOCK_SIZE; length <<= 1) { - for (uint inc = length; inc > 0; inc >>= 1) { - for (uint ii = 0; ii < SHADER_SORT_BLOCK_SIZE; ii += SHADER_SORT_LOCAL_SIZE) { - uint i = lid + ii; - bool direction = ((i & (length << 1)) != 0); - uint j = i ^ inc; - ushort ioff = local_index[i]; - ushort joff = local_index[j]; - uint iKey = local_value[ioff]; - uint jKey = local_value[joff]; - bool smaller = (jKey < iKey) || (jKey == iKey && j < i); - bool swap = smaller ^ (j < i) ^ direction; - ccl_barrier(CCL_LOCAL_MEM_FENCE); - local_index[i] = (swap) ? joff : ioff; - local_index[j] = (swap) ? ioff : joff; - ccl_barrier(CCL_LOCAL_MEM_FENCE); - } - } - } -# endif /* __KERNEL_OPENCL__ */ - - /* copy to destination */ - for (uint i = 0; i < SHADER_SORT_BLOCK_SIZE; i += SHADER_SORT_LOCAL_SIZE) { - uint idx = offset + i + lid; - uint lidx = local_index[i + lid]; - uint outi = output + idx; - uint ini = input + offset + lidx; - uint value = local_value[lidx]; - if (idx < qsize) { - kernel_split_state.queue_data[outi] = (value == (~0)) ? QUEUE_EMPTY_SLOT : - kernel_split_state.queue_data[ini]; - } - } -#endif /* __KERNEL_CUDA__ */ -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_shadow_blocked_ao.h b/intern/cycles/kernel/split/kernel_shadow_blocked_ao.h deleted file mode 100644 index 5d772fc597b..00000000000 --- a/intern/cycles/kernel/split/kernel_shadow_blocked_ao.h +++ /dev/null @@ -1,59 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* Shadow ray cast for AO. */ -ccl_device void kernel_shadow_blocked_ao(KernelGlobals *kg) -{ - unsigned int ao_queue_length = kernel_split_params.queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS]; - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int ray_index = QUEUE_EMPTY_SLOT; - int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (thread_index < ao_queue_length) { - ray_index = get_ray_index(kg, - thread_index, - QUEUE_SHADOW_RAY_CAST_AO_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - } - - if (ray_index == QUEUE_EMPTY_SLOT) { - return; - } - - ShaderData *sd = kernel_split_sd(sd, ray_index); - ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]); - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - float3 throughput = kernel_split_state.throughput[ray_index]; - -#ifdef __BRANCHED_PATH__ - if (!kernel_data.integrator.branched || - IS_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_INDIRECT)) { -#endif - kernel_path_ao(kg, sd, emission_sd, L, state, throughput, shader_bsdf_alpha(kg, sd)); -#ifdef __BRANCHED_PATH__ - } - else { - kernel_branched_path_ao(kg, sd, emission_sd, L, state, throughput); - } -#endif -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_shadow_blocked_dl.h b/intern/cycles/kernel/split/kernel_shadow_blocked_dl.h deleted file mode 100644 index 5e46d300bca..00000000000 --- a/intern/cycles/kernel/split/kernel_shadow_blocked_dl.h +++ /dev/null @@ -1,98 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -/* Shadow ray cast for direct visible light. */ -ccl_device void kernel_shadow_blocked_dl(KernelGlobals *kg) -{ - unsigned int dl_queue_length = kernel_split_params.queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS]; - ccl_barrier(CCL_LOCAL_MEM_FENCE); - - int ray_index = QUEUE_EMPTY_SLOT; - int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (thread_index < dl_queue_length) { - ray_index = get_ray_index(kg, - thread_index, - QUEUE_SHADOW_RAY_CAST_DL_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - } - -#ifdef __BRANCHED_PATH__ - /* TODO(mai): move this somewhere else? */ - if (thread_index == 0) { - /* Clear QUEUE_INACTIVE_RAYS before next kernel. */ - kernel_split_params.queue_index[QUEUE_INACTIVE_RAYS] = 0; - } -#endif /* __BRANCHED_PATH__ */ - - if (ray_index == QUEUE_EMPTY_SLOT) - return; - - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - Ray ray = kernel_split_state.light_ray[ray_index]; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - float3 throughput = kernel_split_state.throughput[ray_index]; - - BsdfEval L_light = kernel_split_state.bsdf_eval[ray_index]; - ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]); - bool is_lamp = kernel_split_state.is_lamp[ray_index]; - -#if defined(__BRANCHED_PATH__) || defined(__SHADOW_TRICKS__) - bool use_branched = false; - int all = 0; - - if (state->flag & PATH_RAY_SHADOW_CATCHER) { - use_branched = true; - all = 1; - } -# if defined(__BRANCHED_PATH__) - else if (kernel_data.integrator.branched) { - use_branched = true; - - if (IS_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_INDIRECT)) { - all = (kernel_data.integrator.sample_all_lights_indirect); - } - else { - all = (kernel_data.integrator.sample_all_lights_direct); - } - } -# endif /* __BRANCHED_PATH__ */ - - if (use_branched) { - kernel_branched_path_surface_connect_light( - kg, sd, emission_sd, state, throughput, 1.0f, L, all); - } - else -#endif /* defined(__BRANCHED_PATH__) || defined(__SHADOW_TRICKS__)*/ - { - /* trace shadow ray */ - float3 shadow; - - if (!shadow_blocked(kg, sd, emission_sd, state, &ray, &shadow)) { - /* accumulate */ - path_radiance_accum_light(kg, L, state, throughput, &L_light, shadow, 1.0f, is_lamp); - } - else { - path_radiance_accum_total_light(L, state, throughput, &L_light); - } - } -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/split/kernel_split_common.h b/intern/cycles/kernel/split/kernel_split_common.h deleted file mode 100644 index 5114f2b03e5..00000000000 --- a/intern/cycles/kernel/split/kernel_split_common.h +++ /dev/null @@ -1,106 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#ifndef __KERNEL_SPLIT_H__ -#define __KERNEL_SPLIT_H__ - -// clang-format off -#include "kernel/kernel_math.h" -#include "kernel/kernel_types.h" - -#include "kernel/split/kernel_split_data.h" - -#include "kernel/kernel_globals.h" -#include "kernel/kernel_color.h" - -#ifdef __OSL__ -# include "kernel/osl/osl_shader.h" -#endif - -#ifdef __KERNEL_OPENCL__ -# include "kernel/kernels/opencl/kernel_opencl_image.h" -#endif -#ifdef __KERNEL_CUDA__ -# include "kernel/kernels/cuda/kernel_cuda_image.h" -#endif -#ifdef __KERNEL_CPU__ -# include "kernel/kernels/cpu/kernel_cpu_image.h" -#endif - -#include "util/util_atomic.h" - -#include "kernel/kernel_path.h" -#ifdef __BRANCHED_PATH__ -# include "kernel/kernel_path_branched.h" -#endif - -#include "kernel/kernel_queues.h" -#include "kernel/kernel_work_stealing.h" - -#ifdef __BRANCHED_PATH__ -# include "kernel/split/kernel_branched.h" -#endif -// clang-format on - -CCL_NAMESPACE_BEGIN - -ccl_device_inline void kernel_split_path_end(KernelGlobals *kg, int ray_index) -{ - ccl_global char *ray_state = kernel_split_state.ray_state; - -#ifdef __BRANCHED_PATH__ -# ifdef __SUBSURFACE__ - ccl_addr_space SubsurfaceIndirectRays *ss_indirect = &kernel_split_state.ss_rays[ray_index]; - - if (ss_indirect->num_rays) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER); - } - else -# endif /* __SUBSURFACE__ */ - if (IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT_SHARED)) { - int orig_ray = kernel_split_state.branched_state[ray_index].original_ray; - - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - PathRadiance *orig_ray_L = &kernel_split_state.path_radiance[orig_ray]; - - path_radiance_sum_indirect(L); - path_radiance_accum_sample(orig_ray_L, L); - - atomic_fetch_and_dec_uint32( - (ccl_global uint *)&kernel_split_state.branched_state[orig_ray].shared_sample_count); - - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_INACTIVE); - } - else if (IS_FLAG(ray_state, ray_index, RAY_BRANCHED_LIGHT_INDIRECT)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_LIGHT_INDIRECT_NEXT_ITER); - } - else if (IS_FLAG(ray_state, ray_index, RAY_BRANCHED_VOLUME_INDIRECT)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_VOLUME_INDIRECT_NEXT_ITER); - } - else if (IS_FLAG(ray_state, ray_index, RAY_BRANCHED_SUBSURFACE_INDIRECT)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_SUBSURFACE_INDIRECT_NEXT_ITER); - } - else { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER); - } -#else - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER); -#endif -} - -CCL_NAMESPACE_END - -#endif /* __KERNEL_SPLIT_H__ */ diff --git a/intern/cycles/kernel/split/kernel_split_data.h b/intern/cycles/kernel/split/kernel_split_data.h deleted file mode 100644 index decc537b39b..00000000000 --- a/intern/cycles/kernel/split/kernel_split_data.h +++ /dev/null @@ -1,77 +0,0 @@ -/* - * Copyright 2011-2016 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#ifndef __KERNEL_SPLIT_DATA_H__ -#define __KERNEL_SPLIT_DATA_H__ - -#include "kernel/split/kernel_split_data_types.h" - -#include "kernel/kernel_globals.h" - -CCL_NAMESPACE_BEGIN - -ccl_device_inline uint64_t split_data_buffer_size(KernelGlobals *kg, size_t num_elements) -{ - (void)kg; /* Unused on CPU. */ - - uint64_t size = 0; -#define SPLIT_DATA_ENTRY(type, name, num) +align_up(num_elements *num * sizeof(type), 16) - size = size SPLIT_DATA_ENTRIES; -#undef SPLIT_DATA_ENTRY - - uint64_t closure_size = sizeof(ShaderClosure) * (kernel_data.integrator.max_closures - 1); - -#ifdef __BRANCHED_PATH__ - size += align_up(num_elements * (sizeof(ShaderData) + closure_size), 16); -#endif - - size += align_up(num_elements * (sizeof(ShaderData) + closure_size), 16); - - return size; -} - -ccl_device_inline void split_data_init(KernelGlobals *kg, - ccl_global SplitData *split_data, - size_t num_elements, - ccl_global void *data, - ccl_global char *ray_state) -{ - (void)kg; /* Unused on CPU. */ - - ccl_global char *p = (ccl_global char *)data; - -#define SPLIT_DATA_ENTRY(type, name, num) \ - split_data->name = (type *)p; \ - p += align_up(num_elements * num * sizeof(type), 16); - SPLIT_DATA_ENTRIES; -#undef SPLIT_DATA_ENTRY - - uint64_t closure_size = sizeof(ShaderClosure) * (kernel_data.integrator.max_closures - 1); - -#ifdef __BRANCHED_PATH__ - split_data->_branched_state_sd = (ShaderData *)p; - p += align_up(num_elements * (sizeof(ShaderData) + closure_size), 16); -#endif - - split_data->_sd = (ShaderData *)p; - p += align_up(num_elements * (sizeof(ShaderData) + closure_size), 16); - - split_data->ray_state = ray_state; -} - -CCL_NAMESPACE_END - -#endif /* __KERNEL_SPLIT_DATA_H__ */ diff --git a/intern/cycles/kernel/split/kernel_split_data_types.h b/intern/cycles/kernel/split/kernel_split_data_types.h deleted file mode 100644 index 06bdce9947d..00000000000 --- a/intern/cycles/kernel/split/kernel_split_data_types.h +++ /dev/null @@ -1,180 +0,0 @@ -/* - * Copyright 2011-2016 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#ifndef __KERNEL_SPLIT_DATA_TYPES_H__ -#define __KERNEL_SPLIT_DATA_TYPES_H__ - -CCL_NAMESPACE_BEGIN - -/* parameters used by the split kernels, we use a single struct to avoid passing these to each - * kernel */ - -typedef struct SplitParams { - WorkTile tile; - uint total_work_size; - - ccl_global unsigned int *work_pools; - - ccl_global int *queue_index; - int queue_size; - ccl_global char *use_queues_flag; - - /* Place for storing sd->flag. AMD GPU OpenCL compiler workaround */ - int dummy_sd_flag; -} SplitParams; - -/* Global memory variables [porting]; These memory is used for - * co-operation between different kernels; Data written by one - * kernel will be available to another kernel via this global - * memory. - */ - -/* SPLIT_DATA_ENTRY(type, name, num) */ - -#ifdef __BRANCHED_PATH__ - -typedef ccl_global struct SplitBranchedState { - /* various state that must be kept and restored after an indirect loop */ - PathState path_state; - float3 throughput; - Ray ray; - - Intersection isect; - - char ray_state; - - /* indirect loop state */ - int next_closure; - int next_sample; - -# ifdef __SUBSURFACE__ - int ss_next_closure; - int ss_next_sample; - int next_hit; - int num_hits; - - uint lcg_state; - LocalIntersection ss_isect; -# endif /* __SUBSURFACE__ */ - - int shared_sample_count; /* number of branched samples shared with other threads */ - int original_ray; /* index of original ray when sharing branched samples */ - bool waiting_on_shared_samples; -} SplitBranchedState; - -# define SPLIT_DATA_BRANCHED_ENTRIES \ - SPLIT_DATA_ENTRY(SplitBranchedState, branched_state, 1) \ - SPLIT_DATA_ENTRY(ShaderData, _branched_state_sd, 0) -#else -# define SPLIT_DATA_BRANCHED_ENTRIES -#endif /* __BRANCHED_PATH__ */ - -#ifdef __SUBSURFACE__ -# define SPLIT_DATA_SUBSURFACE_ENTRIES \ - SPLIT_DATA_ENTRY(ccl_global SubsurfaceIndirectRays, ss_rays, 1) -#else -# define SPLIT_DATA_SUBSURFACE_ENTRIES -#endif /* __SUBSURFACE__ */ - -#ifdef __VOLUME__ -# define SPLIT_DATA_VOLUME_ENTRIES SPLIT_DATA_ENTRY(ccl_global PathState, state_shadow, 1) -#else -# define SPLIT_DATA_VOLUME_ENTRIES -#endif /* __VOLUME__ */ - -#define SPLIT_DATA_ENTRIES \ - SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \ - SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \ - SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \ - SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \ - SPLIT_DATA_ENTRY(ccl_global Intersection, isect, 1) \ - SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \ - SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \ - SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \ - SPLIT_DATA_ENTRY( \ - ccl_global int, queue_data, (NUM_QUEUES * 2)) /* TODO(mai): this is too large? */ \ - SPLIT_DATA_ENTRY(ccl_global uint, buffer_offset, 1) \ - SPLIT_DATA_ENTRY(ShaderDataTinyStorage, sd_DL_shadow, 1) \ - SPLIT_DATA_SUBSURFACE_ENTRIES \ - SPLIT_DATA_VOLUME_ENTRIES \ - SPLIT_DATA_BRANCHED_ENTRIES \ - SPLIT_DATA_ENTRY(ShaderData, _sd, 0) - -/* Entries to be copied to inactive rays when sharing branched samples - * (TODO: which are actually needed?) */ -#define SPLIT_DATA_ENTRIES_BRANCHED_SHARED \ - SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \ - SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \ - SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \ - SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \ - SPLIT_DATA_ENTRY(ccl_global Intersection, isect, 1) \ - SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \ - SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \ - SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \ - SPLIT_DATA_ENTRY(ShaderDataTinyStorage, sd_DL_shadow, 1) \ - SPLIT_DATA_SUBSURFACE_ENTRIES \ - SPLIT_DATA_VOLUME_ENTRIES \ - SPLIT_DATA_BRANCHED_ENTRIES \ - SPLIT_DATA_ENTRY(ShaderData, _sd, 0) - -/* struct that holds pointers to data in the shared state buffer */ -typedef struct SplitData { -#define SPLIT_DATA_ENTRY(type, name, num) type *name; - SPLIT_DATA_ENTRIES -#undef SPLIT_DATA_ENTRY - - /* this is actually in a separate buffer from the rest of the split state data (so it can be read - * back from the host easily) but is still used the same as the other data so we have it here in - * this struct as well - */ - ccl_global char *ray_state; -} SplitData; - -#ifndef __KERNEL_CUDA__ -# define kernel_split_state (kg->split_data) -# define kernel_split_params (kg->split_param_data) -#else -__device__ SplitData __split_data; -# define kernel_split_state (__split_data) -__device__ SplitParams __split_param_data; -# define kernel_split_params (__split_param_data) -#endif /* __KERNEL_CUDA__ */ - -#define kernel_split_sd(sd, ray_index) \ - ((ShaderData *)(((ccl_global char *)kernel_split_state._##sd) + \ - (sizeof(ShaderData) + \ - sizeof(ShaderClosure) * (kernel_data.integrator.max_closures - 1)) * \ - (ray_index))) - -/* Local storage for queue_enqueue kernel. */ -typedef struct QueueEnqueueLocals { - uint queue_atomics[2]; -} QueueEnqueueLocals; - -/* Local storage for holdout_emission_blurring_pathtermination_ao kernel. */ -typedef struct BackgroundAOLocals { - uint queue_atomics_bg; - uint queue_atomics_ao; -} BackgroundAOLocals; - -typedef struct ShaderSortLocals { - uint local_value[SHADER_SORT_BLOCK_SIZE]; - ushort local_index[SHADER_SORT_BLOCK_SIZE]; -} ShaderSortLocals; - -CCL_NAMESPACE_END - -#endif /* __KERNEL_SPLIT_DATA_TYPES_H__ */ diff --git a/intern/cycles/kernel/split/kernel_subsurface_scatter.h b/intern/cycles/kernel/split/kernel_subsurface_scatter.h deleted file mode 100644 index ba06ae3bc53..00000000000 --- a/intern/cycles/kernel/split/kernel_subsurface_scatter.h +++ /dev/null @@ -1,264 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -CCL_NAMESPACE_BEGIN - -#if defined(__BRANCHED_PATH__) && defined(__SUBSURFACE__) - -ccl_device_inline void kernel_split_branched_path_subsurface_indirect_light_init(KernelGlobals *kg, - int ray_index) -{ - kernel_split_branched_path_indirect_loop_init(kg, ray_index); - - SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index]; - - branched_state->ss_next_closure = 0; - branched_state->ss_next_sample = 0; - - branched_state->num_hits = 0; - branched_state->next_hit = 0; - - ADD_RAY_FLAG(kernel_split_state.ray_state, ray_index, RAY_BRANCHED_SUBSURFACE_INDIRECT); -} - -ccl_device_noinline bool kernel_split_branched_path_subsurface_indirect_light_iter( - KernelGlobals *kg, int ray_index) -{ - SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index]; - - ShaderData *sd = kernel_split_sd(branched_state_sd, ray_index); - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]); - - for (int i = branched_state->ss_next_closure; i < sd->num_closure; i++) { - ShaderClosure *sc = &sd->closure[i]; - - if (!CLOSURE_IS_BSSRDF(sc->type)) - continue; - - /* Closure memory will be overwritten, so read required variables now. */ - Bssrdf *bssrdf = (Bssrdf *)sc; - ClosureType bssrdf_type = sc->type; - float bssrdf_roughness = bssrdf->roughness; - - /* set up random number generator */ - if (branched_state->ss_next_sample == 0 && branched_state->next_hit == 0 && - branched_state->next_closure == 0 && branched_state->next_sample == 0) { - branched_state->lcg_state = lcg_state_init_addrspace(&branched_state->path_state, - 0x68bc21eb); - } - int num_samples = kernel_data.integrator.subsurface_samples * 3; - float num_samples_inv = 1.0f / num_samples; - uint bssrdf_rng_hash = cmj_hash(branched_state->path_state.rng_hash, i); - - /* do subsurface scatter step with copy of shader data, this will - * replace the BSSRDF with a diffuse BSDF closure */ - for (int j = branched_state->ss_next_sample; j < num_samples; j++) { - ccl_global PathState *hit_state = &kernel_split_state.path_state[ray_index]; - *hit_state = branched_state->path_state; - hit_state->rng_hash = bssrdf_rng_hash; - path_state_branch(hit_state, j, num_samples); - - ccl_global LocalIntersection *ss_isect = &branched_state->ss_isect; - float bssrdf_u, bssrdf_v; - path_branched_rng_2D( - kg, bssrdf_rng_hash, hit_state, j, num_samples, PRNG_BSDF_U, &bssrdf_u, &bssrdf_v); - - /* intersection is expensive so avoid doing multiple times for the same input */ - if (branched_state->next_hit == 0 && branched_state->next_closure == 0 && - branched_state->next_sample == 0) { - uint lcg_state = branched_state->lcg_state; - LocalIntersection ss_isect_private; - - branched_state->num_hits = subsurface_scatter_multi_intersect( - kg, &ss_isect_private, sd, hit_state, sc, &lcg_state, bssrdf_u, bssrdf_v, true); - - branched_state->lcg_state = lcg_state; - *ss_isect = ss_isect_private; - } - - hit_state->rng_offset += PRNG_BOUNCE_NUM; - -# ifdef __VOLUME__ - Ray volume_ray = branched_state->ray; - bool need_update_volume_stack = kernel_data.integrator.use_volumes && - sd->object_flag & SD_OBJECT_INTERSECTS_VOLUME; -# endif /* __VOLUME__ */ - - /* compute lighting with the BSDF closure */ - for (int hit = branched_state->next_hit; hit < branched_state->num_hits; hit++) { - ShaderData *bssrdf_sd = kernel_split_sd(sd, ray_index); - *bssrdf_sd = *sd; /* note: copy happens each iteration of inner loop, this is - * important as the indirect path will write into bssrdf_sd */ - - LocalIntersection ss_isect_private = *ss_isect; - subsurface_scatter_multi_setup( - kg, &ss_isect_private, hit, bssrdf_sd, hit_state, bssrdf_type, bssrdf_roughness); - *ss_isect = ss_isect_private; - -# ifdef __VOLUME__ - if (need_update_volume_stack) { - /* Setup ray from previous surface point to the new one. */ - float3 P = ray_offset(bssrdf_sd->P, -bssrdf_sd->Ng); - volume_ray.D = normalize_len(P - volume_ray.P, &volume_ray.t); - - for (int k = 0; k < VOLUME_STACK_SIZE; k++) { - hit_state->volume_stack[k] = branched_state->path_state.volume_stack[k]; - } - - kernel_volume_stack_update_for_subsurface( - kg, emission_sd, &volume_ray, hit_state->volume_stack); - } -# endif /* __VOLUME__ */ - -# ifdef __EMISSION__ - if (branched_state->next_closure == 0 && branched_state->next_sample == 0) { - /* direct light */ - if (kernel_data.integrator.use_direct_light) { - int all = (kernel_data.integrator.sample_all_lights_direct) || - (hit_state->flag & PATH_RAY_SHADOW_CATCHER); - kernel_branched_path_surface_connect_light(kg, - bssrdf_sd, - emission_sd, - hit_state, - branched_state->throughput, - num_samples_inv, - L, - all); - } - } -# endif /* __EMISSION__ */ - - /* indirect light */ - if (kernel_split_branched_path_surface_indirect_light_iter( - kg, ray_index, num_samples_inv, bssrdf_sd, false, false)) { - branched_state->ss_next_closure = i; - branched_state->ss_next_sample = j; - branched_state->next_hit = hit; - - return true; - } - - branched_state->next_closure = 0; - } - - branched_state->next_hit = 0; - } - - branched_state->ss_next_sample = 0; - } - - branched_state->ss_next_closure = sd->num_closure; - - branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0); - if (branched_state->waiting_on_shared_samples) { - return true; - } - - kernel_split_branched_path_indirect_loop_end(kg, ray_index); - - return false; -} - -#endif /* __BRANCHED_PATH__ && __SUBSURFACE__ */ - -ccl_device void kernel_subsurface_scatter(KernelGlobals *kg) -{ - int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - if (thread_index == 0) { - /* We will empty both queues in this kernel. */ - kernel_split_params.queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0; - kernel_split_params.queue_index[QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS] = 0; - } - - int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0); - ray_index = get_ray_index(kg, - ray_index, - QUEUE_ACTIVE_AND_REGENERATED_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - get_ray_index(kg, - thread_index, - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - -#ifdef __SUBSURFACE__ - ccl_global char *ray_state = kernel_split_state.ray_state; - - if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - ccl_global PathState *state = &kernel_split_state.path_state[ray_index]; - PathRadiance *L = &kernel_split_state.path_radiance[ray_index]; - ccl_global Ray *ray = &kernel_split_state.ray[ray_index]; - ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index]; - ccl_global SubsurfaceIndirectRays *ss_indirect = &kernel_split_state.ss_rays[ray_index]; - ShaderData *sd = kernel_split_sd(sd, ray_index); - ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]); - - if (sd->flag & SD_BSSRDF) { - -# ifdef __BRANCHED_PATH__ - if (!kernel_data.integrator.branched || - IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT)) { -# endif - if (kernel_path_subsurface_scatter( - kg, sd, emission_sd, L, state, ray, throughput, ss_indirect)) { - kernel_split_path_end(kg, ray_index); - } -# ifdef __BRANCHED_PATH__ - } - else { - kernel_split_branched_path_subsurface_indirect_light_init(kg, ray_index); - - if (kernel_split_branched_path_subsurface_indirect_light_iter(kg, ray_index)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - } -# endif - } - } - -# ifdef __BRANCHED_PATH__ - if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) { - kernel_split_params.queue_index[QUEUE_SUBSURFACE_INDIRECT_ITER] = 0; - } - - /* iter loop */ - ray_index = get_ray_index(kg, - ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0), - QUEUE_SUBSURFACE_INDIRECT_ITER, - kernel_split_state.queue_data, - kernel_split_params.queue_size, - 1); - - if (IS_STATE(ray_state, ray_index, RAY_SUBSURFACE_INDIRECT_NEXT_ITER)) { - /* for render passes, sum and reset indirect light pass variables - * for the next samples */ - path_radiance_sum_indirect(&kernel_split_state.path_radiance[ray_index]); - path_radiance_reset_indirect(&kernel_split_state.path_radiance[ray_index]); - - if (kernel_split_branched_path_subsurface_indirect_light_iter(kg, ray_index)) { - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_REGENERATED); - } - } -# endif /* __BRANCHED_PATH__ */ - -#endif /* __SUBSURFACE__ */ -} - -CCL_NAMESPACE_END |