[Bf-blender-cvs] [f6c6dd4] master: Cycles: Remove meaningless ifdef checks for features in device_opencl

[Sergey Sharybin]

Commit: f6c6dd44de766725e4409b8cae357817661b91fe
Author: Sergey Sharybin
Date:   Thu May 14 23:44:19 2015 +0500
Branches: master
https://developer.blender.org/rBf6c6dd44de766725e4409b8cae357817661b91fe

Cycles: Remove meaningless ifdef checks for features in device_opencl

This file was actually checking for features enabled on CPU and surely all
of them were enabled, so removing them does not cause any difference.

ideally we'll need to do runtime feature detection and just pass some stuff
as NULL to the kernel, or maybe also have variadic kernel entry points which
is also possible quite easily.

===================================================================

M	intern/cycles/device/device_opencl.cpp
M	intern/cycles/kernel/kernel_data_init.cl

===================================================================

diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index f2ac5fc..9153aa3 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -1569,16 +1569,17 @@ public:
 	cl_mem ray_depth_sd_DL_shadow;
 	cl_mem transparent_depth_sd;
 	cl_mem transparent_depth_sd_DL_shadow;
-#ifdef __RAY_DIFFERENTIALS__
+
+	/* Ray differentials. */
 	cl_mem dP_sd, dI_sd;
 	cl_mem dP_sd_DL_shadow, dI_sd_DL_shadow;
 	cl_mem du_sd, dv_sd;
 	cl_mem du_sd_DL_shadow, dv_sd_DL_shadow;
-#endif
-#ifdef __DPDU__
+
+	/* Dp/Du */
 	cl_mem dPdu_sd, dPdv_sd;
 	cl_mem dPdu_sd_DL_shadow, dPdv_sd_DL_shadow;
-#endif
+
 	cl_mem closure_sd;
 	cl_mem closure_sd_DL_shadow;
 	cl_mem num_closure_sd;
@@ -1719,7 +1720,8 @@ public:
 		ray_depth_sd_DL_shadow = NULL;
 		transparent_depth_sd = NULL;
 		transparent_depth_sd_DL_shadow = NULL;
-#ifdef __RAY_DIFFERENTIALS__
+
+		/* Ray differentials. */
 		dP_sd = NULL;
 		dI_sd = NULL;
 		dP_sd_DL_shadow = NULL;
@@ -1728,13 +1730,13 @@ public:
 		dv_sd = NULL;
 		du_sd_DL_shadow = NULL;
 		dv_sd_DL_shadow = NULL;
-#endif
-#ifdef __DPDU__
+
+		/* Dp/Du */
 		dPdu_sd = NULL;
 		dPdv_sd = NULL;
 		dPdu_sd_DL_shadow = NULL;
 		dPdv_sd_DL_shadow = NULL;
-#endif
+
 		closure_sd = NULL;
 		closure_sd_DL_shadow = NULL;
 		num_closure_sd = NULL;
@@ -2088,7 +2090,8 @@ public:
 		release_mem_object_safe(ray_depth_sd_DL_shadow);
 		release_mem_object_safe(transparent_depth_sd);
 		release_mem_object_safe(transparent_depth_sd_DL_shadow);
-#ifdef __RAY_DIFFERENTIALS__
+
+		/* Ray differentials. */
 		release_mem_object_safe(dP_sd);
 		release_mem_object_safe(dP_sd_DL_shadow);
 		release_mem_object_safe(dI_sd);
@@ -2097,13 +2100,13 @@ public:
 		release_mem_object_safe(du_sd_DL_shadow);
 		release_mem_object_safe(dv_sd);
 		release_mem_object_safe(dv_sd_DL_shadow);
-#endif
-#ifdef __DPDU__
+
+		/* Dp/Du */
 		release_mem_object_safe(dPdu_sd);
 		release_mem_object_safe(dPdu_sd_DL_shadow);
 		release_mem_object_safe(dPdv_sd);
 		release_mem_object_safe(dPdv_sd_DL_shadow);
-#endif
+
 		release_mem_object_safe(closure_sd);
 		release_mem_object_safe(closure_sd_DL_shadow);
 		release_mem_object_safe(num_closure_sd);
@@ -2283,7 +2286,7 @@ public:
 			transparent_depth_sd = mem_alloc(num_global_elements * sizeof(int));
 			transparent_depth_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(int));
 
-#ifdef __RAY_DIFFERENTIALS__
+			/* Ray differentials. */
 			dP_sd = mem_alloc(num_global_elements * sizeof(differential3));
 			dP_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(differential3));
 			dI_sd = mem_alloc(num_global_elements * sizeof(differential3));
@@ -2292,14 +2295,13 @@ public:
 			du_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(differential));
 			dv_sd = mem_alloc(num_global_elements * sizeof(differential));
 			dv_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(differential));
-#endif
 
-#ifdef __DPDU__
+			/* Dp/Du */
 			dPdu_sd = mem_alloc(num_global_elements * sizeof(float3));
 			dPdu_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(float3));
 			dPdv_sd = mem_alloc(num_global_elements * sizeof(float3));
 			dPdv_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(float3));
-#endif
+
 			closure_sd = mem_alloc(num_global_elements * ShaderClosure_size);
 			closure_sd_DL_shadow = mem_alloc(num_global_elements * 2 * ShaderClosure_size);
 			num_closure_sd = mem_alloc(num_global_elements * sizeof(int));
@@ -2388,7 +2390,7 @@ public:
 		start_arg_index +=
 			kernel_set_args(ckPathTraceKernel_data_init,
 			                start_arg_index,
-#ifdef __RAY_DIFFERENTIALS__
+			                /* Ray differentials. */
 			                dP_sd,
 			                dP_sd_DL_shadow,
 			                dI_sd,
@@ -2397,13 +2399,13 @@ public:
 			                du_sd_DL_shadow,
 			                dv_sd,
 			                dv_sd_DL_shadow,
-#endif
-#ifdef __DPDU__
+
+			                /* Dp/Du */
 			                dPdu_sd,
 			                dPdu_sd_DL_shadow,
 			                dPdv_sd,
 			                dPdv_sd_DL_shadow,
-#endif
+
 			                closure_sd,
 			                closure_sd_DL_shadow,
 			                num_closure_sd,
diff --git a/intern/cycles/kernel/kernel_data_init.cl b/intern/cycles/kernel/kernel_data_init.cl
index dbf9e62..7982573 100644
--- a/intern/cycles/kernel/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernel_data_init.cl
@@ -100,7 +100,8 @@ __kernel void kernel_ocl_path_trace_data_initialization(
 
 	ccl_global int *transparent_depth_sd,
 	ccl_global int *transparent_depth_sd_DL_shadow,
-	#ifdef __RAY_DIFFERENTIALS__
+
+	/* Ray differentials. */
 	ccl_global differential3 *dP_sd,
 	ccl_global differential3 *dP_sd_DL_shadow,
 
@@ -112,14 +113,14 @@ __kernel void kernel_ocl_path_trace_data_initialization(
 
 	ccl_global differential *dv_sd,
 	ccl_global differential *dv_sd_DL_shadow,
-	#endif
-	#ifdef __DPDU__
+
+	/* 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,
-	#endif
+
 	ShaderClosure *closure_sd,
 	ShaderClosure *closure_sd_DL_shadow,