Merge branch 'blender2.7'

2 files changed, 2 insertions, 62 deletions

diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 0a2acd3f669..7332346a787 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -36,8 +36,8 @@ set(SRC_CUDA_KERNELS
 )
 
 set(SRC_OPENCL_KERNELS
-	kernels/opencl/kernel.cl
 	kernels/opencl/kernel_bake.cl
+	kernels/opencl/kernel_base.cl
 	kernels/opencl/kernel_displace.cl
 	kernels/opencl/kernel_background.cl
 	kernels/opencl/kernel_state_buffer_size.cl
diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel_base.cl
index aa837e2ae87..1c2d89e8a92 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_base.cl
@@ -14,72 +14,14 @@
  * limitations under the License.
  */
 
-/* OpenCL kernel entry points - unfinished */
+/* OpenCL base kernels entry points */
 
 #include "kernel/kernel_compat_opencl.h"
-#include "kernel/kernel_math.h"
 #include "kernel/kernel_types.h"
 #include "kernel/kernel_globals.h"
-#include "kernel/kernel_color.h"
-#include "kernel/kernels/opencl/kernel_opencl_image.h"
 
 #include "kernel/kernel_film.h"
 
-#if defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__)
-#  include "kernel/kernel_path.h"
-#  include "kernel/kernel_path_branched.h"
-#else  /* __COMPILE_ONLY_MEGAKERNEL__ */
-/* Include only actually used headers for the case
- * when path tracing kernels are not needed.
- */
-#  include "kernel/kernel_random.h"
-#  include "kernel/kernel_differential.h"
-#  include "kernel/kernel_montecarlo.h"
-#  include "kernel/kernel_projection.h"
-#  include "kernel/geom/geom.h"
-#  include "kernel/bvh/bvh.h"
-
-#  include "kernel/kernel_accumulate.h"
-#  include "kernel/kernel_camera.h"
-#  include "kernel/kernel_shader.h"
-#endif  /* defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__) */
-
-#include "kernel/kernel_bake.h"
-
-#ifdef __COMPILE_ONLY_MEGAKERNEL__
-
-__kernel void kernel_ocl_path_trace(
-	ccl_constant KernelData *data,
-	ccl_global float *buffer,
-
-	KERNEL_BUFFER_PARAMS,
-
-	int sample,
-	int sx, int sy, int sw, int sh, int offset, int stride)
-{
-	KernelGlobals kglobals, *kg = &kglobals;
-
-	kg->data = data;
-
-	kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
-	kernel_set_buffer_info(kg);
-
-	int x = sx + ccl_global_id(0);
-	int y = sy + ccl_global_id(1);
-	bool thread_is_active = x < sx + sw && y < sy + sh;
-	if(thread_is_active) {
-		kernel_path_trace(kg, buffer, sample, x, y, offset, stride);
-	}
-	if(kernel_data.film.cryptomatte_passes) {
-		/* Make sure no thread is writing to the buffers. */
-		ccl_barrier(CCL_LOCAL_MEM_FENCE);
-		if(thread_is_active) {
-			kernel_cryptomatte_post(kg, buffer, sample, x, y, offset, stride);
-		}
-	}
-}
-
-#else  /* __COMPILE_ONLY_MEGAKERNEL__ */
 
 __kernel void kernel_ocl_convert_to_byte(
 	ccl_constant KernelData *data,
@@ -144,5 +86,3 @@ __kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, u
 		}
 	}
 }
-
-#endif  /* __COMPILE_ONLY_MEGAKERNEL__ */