Code refactor: remove rng_state buffer and compute hash on the fly.

A little faster on some benchmark scenes, a little slower on others, seems
about performance neutral on average and saves a little memory.

7 files changed, 4 insertions, 15 deletions

diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
index c8938534fe8..f5ebf4ad73f 100644
--- a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
+++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
@@ -18,7 +18,6 @@
 
 void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg,
                                            float *buffer,
-                                           unsigned int *rng_state,
                                            int sample,
                                            int x, int y,
                                            int offset,
@@ -57,7 +56,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)(
         ccl_global void *split_data_buffer,
         int num_elements,
         ccl_global char *ray_state,
-        ccl_global uint *rng_state,
         int start_sample,
         int end_sample,
         int sx, int sy, int sw, int sh, int offset, int stride,
diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
index d4315ee5ec4..3fefc1b7e9c 100644
--- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
+++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
@@ -75,7 +75,6 @@ CCL_NAMESPACE_BEGIN
 
 void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg,
                                            float *buffer,
-                                           unsigned int *rng_state,
                                            int sample,
                                            int x, int y,
                                            int offset,
@@ -88,7 +87,6 @@ void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg,
 	if(kernel_data.integrator.branched) {
 		kernel_branched_path_trace(kg,
 		                           buffer,
-		                           rng_state,
 		                           sample,
 		                           x, y,
 		                           offset,
@@ -97,7 +95,7 @@ void KERNEL_FUNCTION_FULL_NAME(path_trace)(KernelGlobals *kg,
 	else
 #  endif
 	{
-		kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
+		kernel_path_trace(kg, buffer, sample, x, y, offset, stride);
 	}
 #endif /* KERNEL_STUB */
 }
diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu
index 4d100634421..799cd587fcf 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel.cu
@@ -42,7 +42,7 @@ kernel_cuda_path_trace(WorkTile *tile, uint total_work_size)
 		get_work_pixel(tile, work_index, &x, &y, &sample);
 
 		KernelGlobals kg;
-		kernel_path_trace(&kg, tile->buffer, tile->rng_state, sample, x, y, tile->offset, tile->stride);
+		kernel_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
 	}
 }
 
@@ -58,7 +58,7 @@ kernel_cuda_branched_path_trace(WorkTile *tile, uint total_work_size)
 		get_work_pixel(tile, work_index, &x, &y, &sample);
 
 		KernelGlobals kg;
-		kernel_branched_path_trace(&kg, tile->buffer, tile->rng_state, sample, x, y, tile->offset, tile->stride);
+		kernel_branched_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
 	}
 }
 #endif
diff --git a/intern/cycles/kernel/kernels/cuda/kernel_split.cu b/intern/cycles/kernel/kernels/cuda/kernel_split.cu
index e97e87285a5..43b3d0aa0e6 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel_split.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel_split.cu
@@ -60,7 +60,6 @@ kernel_cuda_path_trace_data_init(
         ccl_global void *split_data_buffer,
         int num_elements,
         ccl_global char *ray_state,
-        ccl_global uint *rng_state,
         int start_sample,
         int end_sample,
         int sx, int sy, int sw, int sh, int offset, int stride,
@@ -76,7 +75,6 @@ kernel_cuda_path_trace_data_init(
 	                 split_data_buffer,
 	                 num_elements,
 	                 ray_state,
-	                 rng_state,
 	                 start_sample,
 	                 end_sample,
 	                 sx, sy, sw, sh, offset, stride,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel.cl
index b7108f3d0f8..521b86121ff 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel.cl
@@ -50,7 +50,6 @@
 __kernel void kernel_ocl_path_trace(
 	ccl_constant KernelData *data,
 	ccl_global float *buffer,
-	ccl_global uint *rng_state,
 
 	KERNEL_BUFFER_PARAMS,
 
@@ -68,7 +67,7 @@ __kernel void kernel_ocl_path_trace(
 	int y = sy + ccl_global_id(1);
 
 	if(x < sx + sw && y < sy + sh)
-		kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
+		kernel_path_trace(kg, buffer, sample, x, y, offset, stride);
 }
 
 #else  /* __COMPILE_ONLY_MEGAKERNEL__ */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
index 95b35e40a45..7125348a49f 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
@@ -24,7 +24,6 @@ __kernel void kernel_ocl_path_trace_data_init(
         ccl_global void *split_data_buffer,
         int num_elements,
         ccl_global char *ray_state,
-        ccl_global uint *rng_state,
 		KERNEL_BUFFER_PARAMS,
         int start_sample,
         int end_sample,
@@ -41,7 +40,6 @@ __kernel void kernel_ocl_path_trace_data_init(
 	                 split_data_buffer,
 	                 num_elements,
 	                 ray_state,
-	                 rng_state,
 	                 KERNEL_BUFFER_ARGS,
 	                 start_sample,
 	                 end_sample,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_split_function.h b/intern/cycles/kernel/kernels/opencl/kernel_split_function.h
index 499138b5581..6aa7681cbed 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_split_function.h
+++ b/intern/cycles/kernel/kernels/opencl/kernel_split_function.h
@@ -23,7 +23,6 @@ __kernel void KERNEL_NAME_EVAL(kernel_ocl_path_trace, KERNEL_NAME)(
 
 		ccl_global void *split_data_buffer,
 		ccl_global char *ray_state,
-		ccl_global uint *rng_state,
 
 		KERNEL_BUFFER_PARAMS,
 
@@ -42,7 +41,6 @@ __kernel void KERNEL_NAME_EVAL(kernel_ocl_path_trace, KERNEL_NAME)(
 	if(ccl_local_id(0) + ccl_local_id(1) == 0) {
 		kg->data = data;
 
-		kernel_split_params.tile.rng_state = rng_state;
 		kernel_split_params.queue_index = queue_index;
 		kernel_split_params.use_queues_flag = use_queues_flag;
 		kernel_split_params.work_pools = work_pools;