[Bf-blender-cvs] [6e72430] cycles_kernel_split: Cycles kernel split: Fix for progressive refine

[Lenny Wang]

Commit: 6e7243004e35fb381c9d74db713810c690fd2e88
Author: Lenny Wang
Date:   Sat Apr 4 13:40:09 2015 +0500
Branches: cycles_kernel_split
https://developer.blender.org/rB6e7243004e35fb381c9d74db713810c690fd2e88

Cycles kernel split: Fix for progressive refine

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

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

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

diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index 1748ea4..0d3b8bb 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -3259,6 +3259,7 @@ One possible tile size is %zux%zu \n", tile_max_x - local_size[0] , tile_max_y -
 		opencl_assert(clSetKernelArg(ckPathTraceKernel_SumAllRadiance_SPLIT_KERNEL, narg++, sizeof(d_w), (void *)&d_w));
 		opencl_assert(clSetKernelArg(ckPathTraceKernel_SumAllRadiance_SPLIT_KERNEL, narg++, sizeof(d_h), (void *)&d_h));
 		opencl_assert(clSetKernelArg(ckPathTraceKernel_SumAllRadiance_SPLIT_KERNEL, narg++, sizeof(d_stride), (void *)&d_stride));
+		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));
diff --git a/intern/cycles/kernel/kernel_SumAllRadiance.cl b/intern/cycles/kernel/kernel_SumAllRadiance.cl
index 372e595..440bc6f 100644
--- a/intern/cycles/kernel/kernel_SumAllRadiance.cl
+++ b/intern/cycles/kernel/kernel_SumAllRadiance.cl
@@ -29,7 +29,8 @@ __kernel void kernel_ocl_path_trace_SumAllRadiance_SPLIT_KERNEL(
 	ccl_constant KernelData *data,               /* To get pass_stride to offet into buffer */
 	ccl_global float *buffer,                    /* Output buffer of RenderTile */
 	ccl_global float *per_sample_output_buffer,  /* Radiance contributed by all samples */
-	int parallel_samples, int sw, int sh, int stride)
+	int parallel_samples, int sw, int sh, int stride,
+	int start_sample)
 {
 	int x = get_global_id(0);
 	int y = get_global_id(1);
@@ -46,11 +47,10 @@ __kernel void kernel_ocl_path_trace_SumAllRadiance_SPLIT_KERNEL(
 
 		for(sample_iterator = 0; sample_iterator < parallel_samples; sample_iterator++) {
 			for(pass_stride_iterator = 0; pass_stride_iterator < num_floats; pass_stride_iterator++) {
-				*(buffer + pass_stride_iterator) = (sample_iterator == 0) ? *(per_sample_output_buffer + pass_stride_iterator)
+				*(buffer + pass_stride_iterator) = (start_sample == 0 && sample_iterator == 0) ? *(per_sample_output_buffer + pass_stride_iterator)
 				: *(buffer + pass_stride_iterator) + *(per_sample_output_buffer + pass_stride_iterator);
 			}
 			per_sample_output_buffer += sample_stride;
 		}
-
 	}
-}
\ No newline at end of file
+}