[Bf-blender-cvs] [e09cb2d] cycles_kernel_split: Prevent un-necessary megakernel compilation inside splitkernel

[varunsundar08]

Commit: e09cb2d248096940bb61e186b8f34cd26976c224
Author: varunsundar08
Date:   Tue Apr 28 12:46:41 2015 +0530
Branches: cycles_kernel_split
https://developer.blender.org/rBe09cb2d248096940bb61e186b8f34cd26976c224

Prevent un-necessary megakernel compilation inside splitkernel

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

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

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

diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index a99c852..ab96810 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -2361,7 +2361,7 @@ public:
 		current_clos_max = clos_max;
 
 		kernel_init_source = "#include \"kernel.cl\" // " + kernel_md5 + "\n";
-		custom_kernel_build_options = "";
+		custom_kernel_build_options = "-D__SPLIT_KERNEL_COMPILE__ ";
 		device_md5 = device_md5_hash(custom_kernel_build_options);
 		clbin = string_printf("cycles_kernel_%s_%s.clbin", device_md5.c_str(), kernel_md5.c_str());
 		if (!load_split_kernel(&cpProgram, kernel_path, "", device_md5, kernel_init_source, clbin, custom_kernel_build_options))
diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl
index ea9dae1..e744e29 100644
--- a/intern/cycles/kernel/kernel.cl
+++ b/intern/cycles/kernel/kernel.cl
@@ -25,6 +25,15 @@
 #include "kernel_path.h"
 #include "kernel_bake.h"
 
+#ifndef __SPLIT_KERNEL_COMPILE__
+/* kernel_ocl_path_trace is megakernel path-trace invoked by OpenCLDevice class (Megakernel's host side);
+ * OpenCLDeviceSplitKernel (Split-kernel's host side) has a set of split kernels in megakernel's place.
+ * This .cl file also contains extra kernels kernel_ocl_shader, kernel_ocl_bake,
+ * kernel_ocl_convert_to_byte, kernel_ocl_convert_to_half_float which are required
+ * by both OpenCLDevice and OpenCLDeviceSplitKernel class.
+ * The macro __SPLIT_KERNEL_COMPILE__ helps in preventing un-necessary kernel_ocl_path_trace
+ * kernel (megakernel) compilation when using OpenCLDeviceSplitKernel.
+ */
 __kernel void kernel_ocl_path_trace(
 	ccl_constant KernelData *data,
 	ccl_global float *buffer,
@@ -51,6 +60,7 @@ __kernel void kernel_ocl_path_trace(
 	if(x < sx + sw && y < sy + sh)
 		kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
+#endif // __SPLIT_KERNEL_COMPILE__
 
 __kernel void kernel_ocl_shader(
 	ccl_constant KernelData *data,