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
+}
_______________________________________________
Bf-blender-cvs mailing list
[email protected]
http://lists.blender.org/mailman/listinfo/bf-blender-cvs
