[Bf-blender-cvs] [7950b76] cycles_kernel_split: Cycles kernel split : Use ENQUEUE_SPLIT_KERNEL macro to enqueue split kernels

[Jothi_PM]

Commit: 7950b76a5d5476642523c1fee23229ddb27d0a07
Author: Jothi_PM
Date:   Wed May 6 12:12:56 2015 +0530
Branches: cycles_kernel_split
https://developer.blender.org/rB7950b76a5d5476642523c1fee23229ddb27d0a07

Cycles kernel split : Use ENQUEUE_SPLIT_KERNEL macro to enqueue split kernels

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

M	intern/cycles/device/device_opencl.cpp

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

diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index 5d8772a..39770ef 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -60,6 +60,10 @@ CCL_NAMESPACE_BEGIN
  */
 #define DATA_ALLOCATION_MEM_FACTOR 5000000 //5MB
 
+ /* Macro for Enqueuing split kernels */
+#define ENQUEUE_SPLIT_KERNEL(kernelName, globalSize, localSize) opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, kernelName, 2, NULL, globalSize, localSize, 0, NULL, NULL));
+
+
 static cl_device_type opencl_device_type()
 {
 	char *device = getenv("CYCLES_OPENCL_TEST");
@@ -2868,7 +2872,7 @@ public:
 		opencl_assert(clSetKernelArg(ckPathTraceKernel_SumAllRadiance_SPLIT_KERNEL, narg++, sizeof(start_sample), (void*)&start_sample));
 
 		/* Enqueue ckPathTraceKernel_DataInit_SPLIT_KERNEL kernel */
-		opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_DataInit_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
+		ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_DataInit_SPLIT_KERNEL, global_size, local_size);
 		bool activeRaysAvailable = true;
 
 		/* Record number of time host intervention has been made */
@@ -2882,15 +2886,15 @@ public:
 
 			/* Do path-iteration in host [Enqueue Path-iteration kernels] */
 			for(int PathIter = 0; PathIter < PathIteration_times; PathIter++) {
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_SceneIntersect_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_LampEmission_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_QueueEnqueue_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_BG_BufferUpdate_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_Shader_Lighting_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_Holdout_Emission_Blurring_Pathtermination_AO_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, 2, NULL, global_size_shadow_blocked, local_size, 0, NULL, NULL));
-				opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_SetUpNextIteration_SPLIT_KERNEL, 2, NULL, global_size, local_size, 0, NULL, NULL));
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_SceneIntersect_SPLIT_KERNEL, global_size, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_LampEmission_SPLIT_KERNEL, global_size, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_QueueEnqueue_SPLIT_KERNEL, global_size, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_BG_BufferUpdate_SPLIT_KERNEL, global_size, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_Shader_Lighting_SPLIT_KERNEL, global_size, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_Holdout_Emission_Blurring_Pathtermination_AO_SPLIT_KERNEL, global_size, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, global_size, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, global_size_shadow_blocked, local_size);
+				ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_SetUpNextIteration_SPLIT_KERNEL, global_size, local_size);
 			}
 
 			/* Read ray-state into Host memory to decide if we should exit path-iteration in host */
@@ -2925,7 +2929,7 @@ public:
 		size_t sum_all_radiance_global_size[2];
 		sum_all_radiance_global_size[0] = (((d_w - 1) / sum_all_radiance_local_size[0]) + 1) * sum_all_radiance_local_size[0];
 		sum_all_radiance_global_size[1] = (((d_h - 1) / sum_all_radiance_local_size[1]) + 1) * sum_all_radiance_local_size[1];
-		opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel_SumAllRadiance_SPLIT_KERNEL, 2, NULL, sum_all_radiance_global_size, sum_all_radiance_local_size, 0, NULL, NULL));
+		ENQUEUE_SPLIT_KERNEL(ckPathTraceKernel_SumAllRadiance_SPLIT_KERNEL, sum_all_radiance_global_size, sum_all_radiance_local_size);
 
 		if(numHostIntervention == 0) {
 			/* This means that we are executing kernel more than required