[Bf-blender-cvs] [84ad20a] master: Fix T44833: Can't use ccl_local space in non-kernel functions
Sergey Sharybin
noreply at git.blender.org
Tue May 26 19:56:33 CEST 2015
Commit: 84ad20acef4c0db91c9a850e81c7dc0a57aef42a
Author: Sergey Sharybin
Date: Tue May 26 19:12:49 2015 +0500
Branches: master
https://developer.blender.org/rB84ad20acef4c0db91c9a850e81c7dc0a57aef42a
Fix T44833: Can't use ccl_local space in non-kernel functions
This commit re-shuffles code in split kernel once again and makes it so common
parts which is in the headers is only responsible to making all the work needed
for specified ray index. Getting ray index, checking for it's validity and
enqueuing tasks are now happening in the device specified part of the kernel.
This actually makes sense because enqueuing is indeed device-specified and i.e.
with CUDA we'll want to enqueue kernels from kernel and avoid CPU roundtrip.
TODO:
- Kernel comments are still placed in the common header files, but since queue
related stuff is not passed to those functions those comments might need to
be split as well.
Just currently read them considering that they're also covering the way how
all devices are invoking the common code path.
- Arguments might need to be wrapped into KernelGlobals, so we don't ened to
pass all them around as function arguments.
===================================================================
M intern/cycles/kernel/CMakeLists.txt
M intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
M intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
M intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
M intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
M intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
M intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
M intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
M intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
M intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
M intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
M intern/cycles/kernel/split/kernel_background_buffer_update.h
M intern/cycles/kernel/split/kernel_data_init.h
M intern/cycles/kernel/split/kernel_direct_lighting.h
M intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h
M intern/cycles/kernel/split/kernel_lamp_emission.h
M intern/cycles/kernel/split/kernel_next_iteration_setup.h
D intern/cycles/kernel/split/kernel_queue_enqueue.h
M intern/cycles/kernel/split/kernel_scene_intersect.h
M intern/cycles/kernel/split/kernel_shader_eval.h
M intern/cycles/kernel/split/kernel_shadow_blocked.h
M intern/cycles/kernel/split/kernel_sum_all_radiance.h
===================================================================
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 89dd354..fa8f36b 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -169,7 +169,6 @@ set(SRC_SPLIT_HEADERS
split/kernel_holdout_emission_blurring_pathtermination_ao.h
split/kernel_lamp_emission.h
split/kernel_next_iteration_setup.h
- split/kernel_queue_enqueue.h
split/kernel_scene_intersect.h
split/kernel_shader_eval.h
split/kernel_shadow_blocked.h
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
index 2d1944d..eff77b8 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
@@ -48,34 +48,81 @@ __kernel void kernel_ocl_path_trace_background_buffer_update(
#endif
int parallel_samples) /* Number of samples to be processed in parallel */
{
- kernel_background_buffer_update(globals,
- data,
- shader_data,
- per_sample_output_buffers,
- rng_state,
- rng_coop,
- throughput_coop,
- PathRadiance_coop,
- Ray_coop,
- PathState_coop,
- L_transparent_coop,
- ray_state,
- sw, sh, sx, sy, stride,
- rng_state_offset_x,
- rng_state_offset_y,
- rng_state_stride,
- work_array,
- Queue_data,
- Queue_index,
- queuesize,
- end_sample,
- start_sample,
+ ccl_local unsigned int local_queue_atomics;
+ if(get_local_id(0) == 0 && get_local_id(1) == 0) {
+ local_queue_atomics = 0;
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ int ray_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
+ if(ray_index == 0) {
+ /* We will empty this queue in this kernel. */
+ Queue_index[QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS] = 0;
+ }
+ char enqueue_flag = 0;
+ ray_index = get_ray_index(ray_index,
+ QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS,
+ Queue_data,
+ queuesize,
+ 1);
+
+#ifdef __COMPUTE_DEVICE_GPU__
+ /* If we are executing on a GPU device, we exit all threads that are not
+ * required.
+ *
+ * If we are executing on a CPU device, then we need to keep all threads
+ * active since we have barrier() calls later in the kernel. CPU devices,
+ * expect all threads to execute barrier statement.
+ */
+ if(ray_index == QUEUE_EMPTY_SLOT) {
+ return;
+ }
+#endif
+
+#ifndef __COMPUTE_DEVICE_GPU__
+ if(ray_index != QUEUE_EMPTY_SLOT) {
+#endif
+ enqueue_flag =
+ kernel_background_buffer_update(globals,
+ data,
+ shader_data,
+ per_sample_output_buffers,
+ rng_state,
+ rng_coop,
+ throughput_coop,
+ PathRadiance_coop,
+ Ray_coop,
+ PathState_coop,
+ L_transparent_coop,
+ ray_state,
+ sw, sh, sx, sy, stride,
+ rng_state_offset_x,
+ rng_state_offset_y,
+ rng_state_stride,
+ work_array,
+ end_sample,
+ start_sample,
#ifdef __WORK_STEALING__
- work_pool_wgs,
- num_samples,
+ work_pool_wgs,
+ num_samples,
#endif
#ifdef __KERNEL_DEBUG__
- debugdata_coop,
+ debugdata_coop,
+#endif
+ parallel_samples,
+ ray_index);
+#ifndef __COMPUTE_DEVICE_GPU__
+ }
#endif
- parallel_samples);
+
+ /* 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,
+ queuesize,
+ &local_queue_atomics,
+ Queue_data,
+ Queue_index);
}
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
index 015f087..c327767 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
@@ -17,130 +17,129 @@
#include "split/kernel_data_init.h"
__kernel void kernel_ocl_path_trace_data_init(
- ccl_global char *globals,
- ccl_global char *shader_data_sd, /* Arguments related to ShaderData */
- ccl_global char *shader_data_sd_DL_shadow, /* Arguments related to ShaderData */
+ ccl_global char *globals,
+ ccl_global char *shader_data_sd, /* Arguments related to ShaderData */
+ ccl_global char *shader_data_sd_DL_shadow, /* Arguments related to ShaderData */
- ccl_global float3 *P_sd,
- ccl_global float3 *P_sd_DL_shadow,
+ ccl_global float3 *P_sd,
+ ccl_global float3 *P_sd_DL_shadow,
- ccl_global float3 *N_sd,
- ccl_global float3 *N_sd_DL_shadow,
+ ccl_global float3 *N_sd,
+ ccl_global float3 *N_sd_DL_shadow,
- ccl_global float3 *Ng_sd,
- ccl_global float3 *Ng_sd_DL_shadow,
+ ccl_global float3 *Ng_sd,
+ ccl_global float3 *Ng_sd_DL_shadow,
- ccl_global float3 *I_sd,
- ccl_global float3 *I_sd_DL_shadow,
+ ccl_global float3 *I_sd,
+ ccl_global float3 *I_sd_DL_shadow,
- ccl_global int *shader_sd,
- ccl_global int *shader_sd_DL_shadow,
+ ccl_global int *shader_sd,
+ ccl_global int *shader_sd_DL_shadow,
- ccl_global int *flag_sd,
- ccl_global int *flag_sd_DL_shadow,
+ ccl_global int *flag_sd,
+ ccl_global int *flag_sd_DL_shadow,
- ccl_global int *prim_sd,
- ccl_global int *prim_sd_DL_shadow,
+ ccl_global int *prim_sd,
+ ccl_global int *prim_sd_DL_shadow,
- ccl_global int *type_sd,
- ccl_global int *type_sd_DL_shadow,
+ ccl_global int *type_sd,
+ ccl_global int *type_sd_DL_shadow,
- ccl_global float *u_sd,
- ccl_global float *u_sd_DL_shadow,
+ ccl_global float *u_sd,
+ ccl_global float *u_sd_DL_shadow,
- ccl_global float *v_sd,
- ccl_global float *v_sd_DL_shadow,
+ ccl_global float *v_sd,
+ ccl_global float *v_sd_DL_shadow,
- ccl_global int *object_sd,
- ccl_global int *object_sd_DL_shadow,
+ ccl_global int *object_sd,
+ ccl_global int *object_sd_DL_shadow,
- ccl_global float *time_sd,
- ccl_global float *time_sd_DL_shadow,
+ ccl_global float *time_sd,
+ ccl_global float *time_sd_DL_shadow,
- ccl_global float *ray_length_sd,
- ccl_global float *ray_length_sd_DL_shadow,
+ ccl_global float *ray_length_sd,
+ ccl_global float *ray_length_sd_DL_shadow,
- ccl_global int *ray_depth_sd,
- ccl_global int *ray_depth_sd_DL_shadow,
+ ccl_global int *ray_depth_sd,
+ ccl_global int *ray_depth_sd_DL_shadow,
- ccl_global int *transparent_depth_sd,
- ccl_global int *transparent_depth_sd_DL_shadow,
+ ccl_global int *transparent_depth_sd,
+ ccl_global int *transparent_depth_sd_DL_shadow,
- /* Ray differentials. */
- ccl_global differential3 *dP_sd,
- ccl_global differential3 *dP_sd_DL_shadow,
+ /* Ray differentials. */
+ ccl_global differential3 *dP_sd,
+ ccl_global differential3 *dP_sd_DL_shadow,
- ccl_global differential3 *dI_sd,
- ccl_global differential3 *dI_sd_DL_shadow,
+ ccl_global differential3 *dI_sd,
+ ccl_global differential3 *dI_sd_DL_shadow,
- ccl_global differential *du_sd,
- ccl_global differential *du_sd_DL_shadow,
+ ccl_global differential *du_sd,
+ ccl_global differential *du_sd_DL_shadow,
- ccl_global differential *dv_sd,
- ccl_global differential *dv_sd_DL_shadow,
+ ccl_global differential *dv_sd,
+ ccl_global differential *dv_sd_DL_shadow,
- /* Dp/Du */
- ccl_global float3 *dPdu_sd,
- ccl_global float3 *dPdu_sd_DL_shadow,
+ /* Dp/Du */
+ ccl_global float3 *dPdu_sd,
+ ccl_global float3 *dPdu_sd_DL_shadow,
- ccl_global float3 *dPdv_sd,
- ccl_global float3 *dPdv_sd_DL_shadow,
+ ccl_global float3 *dPdv_sd,
+ ccl_global float3 *dPdv_sd_DL_shadow,
- /* Object motion. */
- ccl_global Transform *ob_tfm_sd,
- ccl_global Transform *ob_tfm_sd_DL_shadow,
+ /* Object motion. */
+ ccl_global Transform *ob_tfm_sd,
+ ccl_global Transform *ob_tfm_sd_DL_shadow,
- ccl_global Transform *ob_itfm_sd,
- ccl_global Transform *ob_itfm_sd_DL_shadow,
+ ccl_global Transform *ob_itfm_sd,
+ ccl_global Transform *ob_itfm_sd_DL_shadow,
- ShaderClosure *closure_sd,
- ShaderClosure *closure_sd_DL_shadow,
+ ShaderClosure *closure_sd,
+ ShaderClosure *closure_sd_DL_shadow,
- ccl_global int *num_closure_sd,
- ccl_global int *num_closure_sd_DL_shadow,
+ ccl_global int *num_closure_sd,
+ ccl_global int *num_closure_sd_DL_shadow,
- ccl_global float *randb_closure_sd,
- ccl_global float *randb_closure_sd_DL_s
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list