[Bf-blender-cvs] [270b54a] cycles_kernel_split: Cycles kernel split : Refactor: replace address specifiers with ccl_ macros
Kavitha Sampath
noreply at git.blender.org
Wed May 6 09:54:20 CEST 2015
Commit: 270b54adcbf9e2883810b31aaf421c284f8cdc80
Author: Kavitha Sampath
Date: Wed May 6 12:10:39 2015 +0530
Branches: cycles_kernel_split
https://developer.blender.org/rB270b54adcbf9e2883810b31aaf421c284f8cdc80
Cycles kernel split : Refactor: replace address specifiers with ccl_ macros
__local with ccl_local
__global with ccl_global
__private with ccl_private
inline with ccl_device_inline
===================================================================
M intern/cycles/kernel/kernel_Background_BufferUpdate.cl
M intern/cycles/kernel/kernel_DirectLighting.cl
M intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
M intern/cycles/kernel/kernel_LampEmission.cl
M intern/cycles/kernel/kernel_NextIterationSetUp.cl
M intern/cycles/kernel/kernel_QueueEnqueue.cl
M intern/cycles/kernel/kernel_SceneIntersect.cl
M intern/cycles/kernel/kernel_ShaderEval.cl
M intern/cycles/kernel/kernel_ShadowBlocked.cl
M intern/cycles/kernel/kernel_compat_opencl.h
M intern/cycles/kernel/kernel_passes.h
M intern/cycles/kernel/kernel_queues.h
M intern/cycles/kernel/kernel_work_stealing.h
===================================================================
diff --git a/intern/cycles/kernel/kernel_Background_BufferUpdate.cl b/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
index 3f1db1b..7be0e84 100644
--- a/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
+++ b/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
@@ -128,7 +128,7 @@ __kernel void kernel_ocl_path_trace_Background_BufferUpdate_SPLIT_KERNEL(
int parallel_samples /* Number of samples to be processed in parallel */
)
{
- __local unsigned int local_queue_atomics;
+ ccl_local unsigned int local_queue_atomics;
if(get_local_id(0) == 0 && get_local_id(1) == 0) {
local_queue_atomics = 0;
}
diff --git a/intern/cycles/kernel/kernel_DirectLighting.cl b/intern/cycles/kernel/kernel_DirectLighting.cl
index 47dae0a..f9ce761 100644
--- a/intern/cycles/kernel/kernel_DirectLighting.cl
+++ b/intern/cycles/kernel/kernel_DirectLighting.cl
@@ -65,7 +65,7 @@ __kernel void kernel_ocl_path_trace_DirectLighting_SPLIT_KERNEL(
int queuesize /* Size (capacity) of each queue */
)
{
- __local unsigned int local_queue_atomics;
+ ccl_local unsigned int local_queue_atomics;
if(get_local_id(0) == 0 && get_local_id(1) == 0) {
local_queue_atomics = 0;
}
diff --git a/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl b/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
index 91adbf0..591b970 100644
--- a/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
+++ b/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
@@ -98,8 +98,8 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_AO
int parallel_samples /* Number of samples to be processed in parallel */
)
{
- __local unsigned int local_queue_atomics_bg;
- __local unsigned int local_queue_atomics_ao;
+ ccl_local unsigned int local_queue_atomics_bg;
+ ccl_local unsigned int local_queue_atomics_ao;
if(get_local_id(0) == 0 && get_local_id(1) == 0) {
local_queue_atomics_bg = 0;
local_queue_atomics_ao = 0;
diff --git a/intern/cycles/kernel/kernel_LampEmission.cl b/intern/cycles/kernel/kernel_LampEmission.cl
index 4cd4482..38d0984 100644
--- a/intern/cycles/kernel/kernel_LampEmission.cl
+++ b/intern/cycles/kernel/kernel_LampEmission.cl
@@ -67,7 +67,7 @@ __kernel void kernel_ocl_path_trace_LampEmission_SPLIT_KERNEL(
}
/* Fetch use_queues_flag */
- __local char local_use_queues_flag;
+ ccl_local char local_use_queues_flag;
if(get_local_id(0) == 0 && get_local_id(1) == 0) {
local_use_queues_flag = use_queues_flag[0];
}
diff --git a/intern/cycles/kernel/kernel_NextIterationSetUp.cl b/intern/cycles/kernel/kernel_NextIterationSetUp.cl
index c7fa362..c406696 100644
--- a/intern/cycles/kernel/kernel_NextIterationSetUp.cl
+++ b/intern/cycles/kernel/kernel_NextIterationSetUp.cl
@@ -84,7 +84,7 @@ __kernel void kernel_ocl_path_trace_SetupNextIteration_SPLIT_KERNEL(
)
{
- __local unsigned int local_queue_atomics;
+ ccl_local unsigned int local_queue_atomics;
if(get_local_id(0) == 0 && get_local_id(1) == 0) {
local_queue_atomics = 0;
}
diff --git a/intern/cycles/kernel/kernel_QueueEnqueue.cl b/intern/cycles/kernel/kernel_QueueEnqueue.cl
index 3926574..d83290d 100644
--- a/intern/cycles/kernel/kernel_QueueEnqueue.cl
+++ b/intern/cycles/kernel/kernel_QueueEnqueue.cl
@@ -60,7 +60,7 @@ __kernel void kernel_ocl_path_trace_QueueEnqueue_SPLIT_KERNEL(
)
{
/* We have only 2 cases (Hit/Not-Hit) */
- __local unsigned int local_queue_atomics[2];
+ ccl_local unsigned int local_queue_atomics[2];
int lidx = get_local_id(1) * get_local_size(0) + get_local_id(0);
int ray_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
diff --git a/intern/cycles/kernel/kernel_SceneIntersect.cl b/intern/cycles/kernel/kernel_SceneIntersect.cl
index 3db25af..eb37e1d 100644
--- a/intern/cycles/kernel/kernel_SceneIntersect.cl
+++ b/intern/cycles/kernel/kernel_SceneIntersect.cl
@@ -86,7 +86,7 @@ __kernel void kernel_ocl_path_trace_SceneIntersect_SPLIT_KERNEL(
int y = get_global_id(1);
/* Fetch use_queues_flag */
- __local char local_use_queues_flag;
+ ccl_local char local_use_queues_flag;
if(get_local_id(0) == 0 && get_local_id(1) == 0) {
local_use_queues_flag = use_queues_flag[0];
}
diff --git a/intern/cycles/kernel/kernel_ShaderEval.cl b/intern/cycles/kernel/kernel_ShaderEval.cl
index 7b0fd95..ac0f0fe 100644
--- a/intern/cycles/kernel/kernel_ShaderEval.cl
+++ b/intern/cycles/kernel/kernel_ShaderEval.cl
@@ -62,7 +62,7 @@ __kernel void kernel_ocl_path_trace_ShaderEvaluation_SPLIT_KERNEL(
{
int ray_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
/* Enqeueue RAY_TO_REGENERATE rays into QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue */
- __local unsigned int local_queue_atomics;
+ ccl_local unsigned int local_queue_atomics;
if(get_local_id(0) == 0 && get_local_id(1) == 0) {
local_queue_atomics = 0;
}
diff --git a/intern/cycles/kernel/kernel_ShadowBlocked.cl b/intern/cycles/kernel/kernel_ShadowBlocked.cl
index 6f380ae..52fa89d 100644
--- a/intern/cycles/kernel/kernel_ShadowBlocked.cl
+++ b/intern/cycles/kernel/kernel_ShadowBlocked.cl
@@ -74,8 +74,8 @@ __kernel void kernel_ocl_path_trace_ShadowBlocked_DirectLighting_SPLIT_KERNEL(
int lidx = get_local_id(1) * get_local_id(0) + get_local_id(0);
- __local unsigned int ao_queue_length;
- __local unsigned int dl_queue_length;
+ ccl_local unsigned int ao_queue_length;
+ ccl_local unsigned int dl_queue_length;
if(lidx == 0) {
ao_queue_length = Queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS];
dl_queue_length = Queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS];
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index a601b61..29de7dd 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -37,6 +37,8 @@
#define ccl_may_alias
#define ccl_constant __constant
#define ccl_global __global
+#define ccl_local __local
+#define ccl_private __private
#ifdef __SPLIT_KERNEL__
#define ccl_addr_space __global
diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h
index ebcfa83..8b15144 100644
--- a/intern/cycles/kernel/kernel_passes.h
+++ b/intern/cycles/kernel/kernel_passes.h
@@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN
#if defined(__SPLIT_KERNEL__) && defined(__WORK_STEALING__)
/* Utility functions for float atomics */
/* float atomics impl credits : http://suhorukov.blogspot.in/2011/12/opencl-11-atomic-operations-on-floating.html */
-inline void atomic_add_float(volatile __global float *source, const float operand) {
+ccl_device_inline void atomic_add_float(volatile ccl_global float *source, const float operand) {
union {
unsigned int intVal;
float floatVal;
@@ -34,7 +34,7 @@ inline void atomic_add_float(volatile __global float *source, const float operan
prevVal.floatVal = *source;
newVal.floatVal = prevVal.floatVal + operand;
- } while(atomic_cmpxchg((volatile __global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal);
+ } while(atomic_cmpxchg((volatile ccl_global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal);
}
#endif // __SPLIT_KERNEL__ && __WORK_STEALING__
diff --git a/intern/cycles/kernel/kernel_queues.h b/intern/cycles/kernel/kernel_queues.h
index df52dcf..9e65e2b 100644
--- a/intern/cycles/kernel/kernel_queues.h
+++ b/intern/cycles/kernel/kernel_queues.h
@@ -72,7 +72,7 @@ ccl_device void enqueue_ray_index_local(
int queue_number, /* Queue in which to enqueue ray index */
char enqueue_flag, /* True for threads whose ray index has to be enqueued */
int queuesize, /* queue size */
- __local unsigned int *local_queue_atomics, /* To to local queue atomics */
+ ccl_local unsigned int *local_queue_atomics, /* To to local queue atomics */
ccl_global int *Queue_data, /* Queues */
ccl_global int *Queue_index /* To do global queue atomics */
)
@@ -101,7 +101,7 @@ ccl_device void enqueue_ray_index_local(
ccl_device unsigned int get_local_queue_index(
int queue_number, /* Queue in which to enqueue the ray; -1 if no queue */
- __local unsigned int *local_queue_atomics
+ ccl_local unsigned int *local_queue_atomics
)
{
int my_lqidx = atomic_inc(&local_queue_atomics[queue_number]);
@@ -110,8 +110,8 @@ ccl_device unsigned int get_local_queue_index(
ccl_device unsigned int get_global_per_queue_offset(
int queue_number,
- __local unsigned int *local_queue_atomics,
- __global int* global_queue_atomics
+ ccl_local unsigned int *local_queue_atomics,
+ ccl_global int* global_queue_atomics
)
{
unsigned int queue_offset = atomic_add((&global_queue_atomics[queue_number]), local_queue_atomics[queue_number]);
@@ -122,7 +122,7 @@ ccl_device unsigned int get_global_queue_index(
int queue_number,
int queuesize,
unsigned int lqidx,
- __local unsigned int * global_per_queue_offset
+ ccl_local unsigned int * global_per_queue_offset
)
{
int my_gqidx = queuesize * queue_number + lqidx + global_per_queue_offset[queue_number];
diff --git a/intern/cycles/kernel/kernel_work_stealing.h b/intern/cycles/kernel/kernel_work_stealing.h
index 9ffab1b..0231fe7 100644
--- a/intern/cycles/kernel/kernel_work_stealing.h
+++ b/intern/cycles/kernel/kernel_work_stealing.h
@@ -61,8 +61,8 @@ unsigned int get_total_work(unsigned int tile_dim_x,
/* Returns 0 in case there is no next work available */
/* Returns 1 in case work assigned is valid */
-int get_next_work(__global unsigned int *work_pool,
- __private unsigned int *my_work,
+int get_next_work(ccl_global unsigned int *work_pool,
+ ccl_private unsigned int *my_work,
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list