diff options
25 files changed, 21 insertions, 93 deletions
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp index 6a1106328fb..72330b02a28 100644 --- a/intern/cycles/device/device_cpu.cpp +++ b/intern/cycles/device/device_cpu.cpp @@ -171,7 +171,7 @@ public: DeviceRequestedFeatures requested_features; - KernelFunctions<void(*)(KernelGlobals *, float *, unsigned int *, int, int, int, int, int)> path_trace_kernel; + KernelFunctions<void(*)(KernelGlobals *, float *, int, int, int, int, int)> path_trace_kernel; KernelFunctions<void(*)(KernelGlobals *, uchar4 *, float *, float, int, int, int, int)> convert_to_half_float_kernel; KernelFunctions<void(*)(KernelGlobals *, uchar4 *, float *, float, int, int, int, int)> convert_to_byte_kernel; KernelFunctions<void(*)(KernelGlobals *, uint4 *, float4 *, float*, int, int, int, int, int)> shader_kernel; @@ -192,7 +192,7 @@ public: KernelFunctions<void(*)(int, int, int, int, int, float*, int*, float*, float3*, int*, int)> filter_finalize_kernel; KernelFunctions<void(*)(KernelGlobals *, ccl_constant KernelData*, ccl_global void*, int, ccl_global char*, - ccl_global uint*, int, int, int, int, int, int, int, int, ccl_global int*, int, + int, int, int, int, int, int, int, int, ccl_global int*, int, ccl_global char*, ccl_global unsigned int*, unsigned int, ccl_global float*)> data_init_kernel; unordered_map<string, KernelFunctions<void(*)(KernelGlobals*, KernelData*)> > split_kernels; @@ -617,7 +617,6 @@ public: void path_trace(DeviceTask &task, RenderTile &tile, KernelGlobals *kg) { float *render_buffer = (float*)tile.buffer; - uint *rng_state = (uint*)tile.rng_state; int start_sample = tile.start_sample; int end_sample = tile.start_sample + tile.num_samples; @@ -629,7 +628,7 @@ public: for(int y = tile.y; y < tile.y + tile.h; y++) { for(int x = tile.x; x < tile.x + tile.w; x++) { - path_trace_kernel()(kg, render_buffer, rng_state, + path_trace_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride); } } @@ -913,7 +912,6 @@ bool CPUSplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim, (void*)split_data.device_pointer, num_global_elements, (char*)ray_state.device_pointer, - (uint*)rtile.rng_state, rtile.start_sample, rtile.start_sample + rtile.num_samples, rtile.x, diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp index 7ee74e9a512..e5464dcf34e 100644 --- a/intern/cycles/device/device_cuda.cpp +++ b/intern/cycles/device/device_cuda.cpp @@ -1322,7 +1322,6 @@ public: wtile->start_sample = sample; wtile->num_samples = 1; wtile->buffer = (float*)cuda_device_ptr(rtile.buffer); - wtile->rng_state = (uint*)cuda_device_ptr(rtile.rng_state); mem_alloc("work_tiles", work_tiles, MEM_READ_ONLY); mem_copy_to(work_tiles); @@ -1945,7 +1944,6 @@ bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim CUdeviceptr d_use_queues_flag = device->cuda_device_ptr(use_queues_flag.device_pointer); CUdeviceptr d_work_pool_wgs = device->cuda_device_ptr(work_pool_wgs.device_pointer); - CUdeviceptr d_rng_state = device->cuda_device_ptr(rtile.rng_state); CUdeviceptr d_buffer = device->cuda_device_ptr(rtile.buffer); int end_sample = rtile.start_sample + rtile.num_samples; @@ -1955,7 +1953,6 @@ bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim CUdeviceptr* split_data_buffer; int* num_elements; CUdeviceptr* ray_state; - CUdeviceptr* rng_state; int* start_sample; int* end_sample; int* sx; @@ -1976,7 +1973,6 @@ bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim &d_split_data, &num_global_elements, &d_ray_state, - &d_rng_state, &rtile.start_sample, &end_sample, &rtile.x, diff --git a/intern/cycles/device/device_multi.cpp b/intern/cycles/device/device_multi.cpp index bc505b676fc..164ed50bdf6 100644 --- a/intern/cycles/device/device_multi.cpp +++ b/intern/cycles/device/device_multi.cpp @@ -281,7 +281,6 @@ public: foreach(SubDevice& sub, devices) { if(sub.device == sub_device) { if(tile.buffer) tile.buffer = sub.ptr_map[tile.buffer]; - if(tile.rng_state) tile.rng_state = sub.ptr_map[tile.rng_state]; } } } diff --git a/intern/cycles/device/device_network.cpp b/intern/cycles/device/device_network.cpp index 571ba9465ca..4ff8647f66b 100644 --- a/intern/cycles/device/device_network.cpp +++ b/intern/cycles/device/device_network.cpp @@ -737,7 +737,6 @@ protected: tile = entry.tile; if(tile.buffer) tile.buffer = ptr_map[tile.buffer]; - if(tile.rng_state) tile.rng_state = ptr_map[tile.rng_state]; result = true; break; @@ -769,7 +768,6 @@ protected: thread_scoped_lock acquire_lock(acquire_mutex); if(tile.buffer) tile.buffer = ptr_imap[tile.buffer]; - if(tile.rng_state) tile.rng_state = ptr_imap[tile.rng_state]; { thread_scoped_lock lock(rpc_lock); diff --git a/intern/cycles/device/device_network.h b/intern/cycles/device/device_network.h index a5d24c66018..7bfebaf5aec 100644 --- a/intern/cycles/device/device_network.h +++ b/intern/cycles/device/device_network.h @@ -142,7 +142,7 @@ public: archive & tile.x & tile.y & tile.w & tile.h; archive & tile.start_sample & tile.num_samples & tile.sample; archive & tile.resolution & tile.offset & tile.stride; - archive & tile.buffer & tile.rng_state; + archive & tile.buffer; } void write() @@ -303,7 +303,7 @@ public: *archive & tile.x & tile.y & tile.w & tile.h; *archive & tile.start_sample & tile.num_samples & tile.sample; *archive & tile.resolution & tile.offset & tile.stride; - *archive & tile.buffer & tile.rng_state; + *archive & tile.buffer; tile.buffers = NULL; } diff --git a/intern/cycles/device/opencl/opencl_mega.cpp b/intern/cycles/device/opencl/opencl_mega.cpp index ec47fdafa3d..f4555eaba4f 100644 --- a/intern/cycles/device/opencl/opencl_mega.cpp +++ b/intern/cycles/device/opencl/opencl_mega.cpp @@ -62,7 +62,6 @@ public: /* Cast arguments to cl types. */ cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer); cl_mem d_buffer = CL_MEM_PTR(rtile.buffer); - cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state); cl_int d_x = rtile.x; cl_int d_y = rtile.y; cl_int d_w = rtile.w; @@ -79,8 +78,7 @@ public: kernel_set_args(ckPathTraceKernel, 0, d_data, - d_buffer, - d_rng_state); + d_buffer); set_kernel_arg_buffers(ckPathTraceKernel, &start_arg_index); diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp index 16a96213100..976cc9df46d 100644 --- a/intern/cycles/device/opencl/opencl_split.cpp +++ b/intern/cycles/device/opencl/opencl_split.cpp @@ -192,7 +192,6 @@ struct CachedSplitMemory { int id; device_memory *split_data; device_memory *ray_state; - device_ptr *rng_state; device_memory *queue_index; device_memory *use_queues_flag; device_memory *work_pools; @@ -225,8 +224,7 @@ public: kg, data, *cached_memory.split_data, - *cached_memory.ray_state, - *cached_memory.rng_state); + *cached_memory.ray_state); device->set_kernel_arg_buffers(program(), &start_arg_index); @@ -356,8 +354,7 @@ public: kernel_data, split_data, num_global_elements, - ray_state, - rtile.rng_state); + ray_state); device->set_kernel_arg_buffers(device->program_data_init(), &start_arg_index); @@ -401,7 +398,6 @@ public: cached_memory.split_data = &split_data; cached_memory.ray_state = &ray_state; - cached_memory.rng_state = &rtile.rng_state; cached_memory.queue_index = &queue_index; cached_memory.use_queues_flag = &use_queues_flag; cached_memory.work_pools = &work_pool_wgs; diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index 793fede0deb..341ada63ec3 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -672,21 +672,20 @@ ccl_device_forceinline void kernel_path_integrate( } ccl_device void kernel_path_trace(KernelGlobals *kg, - ccl_global float *buffer, ccl_global uint *rng_state, + ccl_global float *buffer, int sample, int x, int y, int offset, int stride) { /* buffer offset */ int index = offset + x + y*stride; int pass_stride = kernel_data.film.pass_stride; - rng_state += index; buffer += index*pass_stride; /* Initialize random numbers and sample ray. */ uint rng_hash; Ray ray; - kernel_path_trace_setup(kg, rng_state, sample, x, y, &rng_hash, &ray); + kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &ray); if(ray.t == 0.0f) { kernel_write_result(kg, buffer, sample, NULL); diff --git a/intern/cycles/kernel/kernel_path_branched.h b/intern/cycles/kernel/kernel_path_branched.h index 6e0ec22d581..70d73a20b97 100644 --- a/intern/cycles/kernel/kernel_path_branched.h +++ b/intern/cycles/kernel/kernel_path_branched.h @@ -538,21 +538,20 @@ ccl_device void kernel_branched_path_integrate(KernelGlobals *kg, } ccl_device void kernel_branched_path_trace(KernelGlobals *kg, - ccl_global float *buffer, ccl_global uint *rng_state, + ccl_global float *buffer, int sample, int x, int y, int offset, int stride) { /* buffer offset */ int index = offset + x + y*stride; int pass_stride = kernel_data.film.pass_stride; - rng_state += index; buffer += index*pass_stride; /* initialize random numbers and ray */ uint rng_hash; Ray ray; - kernel_path_trace_setup(kg, rng_state, sample, x, y, &rng_hash, &ray); + kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &ray); /* integrate */ PathRadiance L; diff --git a/intern/cycles/kernel/kernel_path_common.h b/intern/cycles/kernel/kernel_path_common.h index 54dd278a185..d83fd474cde 100644 --- a/intern/cycles/kernel/kernel_path_common.h +++ b/intern/cycles/kernel/kernel_path_common.h @@ -19,7 +19,6 @@ CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_path_trace_setup(KernelGlobals *kg, - ccl_global uint *rng_state, int sample, int x, int y, uint *rng_hash, @@ -30,11 +29,7 @@ ccl_device_inline void kernel_path_trace_setup(KernelGlobals *kg, int num_samples = kernel_data.integrator.aa_samples; - if(sample == kernel_data.integrator.start_sample) { - *rng_state = hash_int_2d(x, y); - } - - path_rng_init(kg, rng_state, sample, num_samples, rng_hash, x, y, &filter_u, &filter_v); + path_rng_init(kg, sample, num_samples, rng_hash, x, y, &filter_u, &filter_v); /* sample camera ray */ diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h index 11798d87cb5..e7a6134b8eb 100644 --- a/intern/cycles/kernel/kernel_random.h +++ b/intern/cycles/kernel/kernel_random.h @@ -15,6 +15,7 @@ */ #include "kernel/kernel_jitter.h" +#include "util/util_hash.h" CCL_NAMESPACE_BEGIN @@ -115,14 +116,13 @@ ccl_device_forceinline void path_rng_2D(KernelGlobals *kg, } ccl_device_inline void path_rng_init(KernelGlobals *kg, - ccl_global uint *rng_state, int sample, int num_samples, uint *rng_hash, int x, int y, float *fx, float *fy) { /* load state */ - *rng_hash = *rng_state; + *rng_hash = hash_int_2d(x, y); *rng_hash ^= kernel_data.integrator.seed; #ifdef __DEBUG_CORRELATION__ diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index bf3a2881666..19c77c1ed4f 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -1460,7 +1460,6 @@ typedef struct WorkTile { uint stride; ccl_global float *buffer; - ccl_global uint *rng_state; } WorkTile; CCL_NAMESPACE_END 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; diff --git a/intern/cycles/kernel/split/kernel_buffer_update.h b/intern/cycles/kernel/split/kernel_buffer_update.h index e8547767480..5e9db821f7b 100644 --- a/intern/cycles/kernel/split/kernel_buffer_update.h +++ b/intern/cycles/kernel/split/kernel_buffer_update.h @@ -108,10 +108,6 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg, uint x, y, sample; get_work_pixel(tile, work_index, &x, &y, &sample); - /* Remap rng_state to current pixel. */ - ccl_global uint *rng_state = kernel_split_params.tile.rng_state; - rng_state += tile->offset + x + y*tile->stride; - /* Store buffer offset for writing to passes. */ uint buffer_offset = (tile->offset + x + y*tile->stride) * kernel_data.film.pass_stride; ccl_global float *buffer = tile->buffer + buffer_offset; @@ -119,7 +115,7 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg, /* Initialize random numbers and ray. */ uint rng_hash; - kernel_path_trace_setup(kg, rng_state, sample, x, y, &rng_hash, ray); + kernel_path_trace_setup(kg, sample, x, y, &rng_hash, ray); if(ray->t != 0.0f) { /* Initialize throughput, path radiance, Ray, PathState; diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h index 2da3ca47466..f4df949fdb1 100644 --- a/intern/cycles/kernel/split/kernel_data_init.h +++ b/intern/cycles/kernel/split/kernel_data_init.h @@ -49,7 +49,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, #ifdef __KERNEL_OPENCL__ KERNEL_BUFFER_PARAMS, @@ -84,7 +83,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( kernel_split_params.tile.offset = offset; kernel_split_params.tile.stride = stride; - kernel_split_params.tile.rng_state = rng_state; kernel_split_params.tile.buffer = buffer; kernel_split_params.total_work_size = sw * sh * num_samples; @@ -122,7 +120,7 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( *use_queues_flag = 0; } - /* zero the tiles pixels and initialize rng_state if this is the first sample */ + /* zero the tiles pixels if this is the first sample */ if(start_sample == 0) { int pass_stride = kernel_data.film.pass_stride; @@ -130,9 +128,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( for(int y = sy; y < sy + sh; y++) { int index = offset + y * stride; memset(buffer + (sx + index) * pass_stride, 0, sizeof(float) * pass_stride * sw); - for(int x = sx; x < sx + sw; x++) { - rng_state[index + x] = hash_int_2d(x, y); - } } #else parallel_for(kg, i, sw * sh * pass_stride) { @@ -146,14 +141,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)( *(buffer + index) = 0.0f; } - - parallel_for(kg, i, sw * sh) { - int x = sx + i % sw; - int y = sy + i / sw; - - int index = (offset + x + y*stride); - *(rng_state + index) = hash_int_2d(x, y); - } #endif } diff --git a/intern/cycles/kernel/split/kernel_path_init.h b/intern/cycles/kernel/split/kernel_path_init.h index 701d39403ad..1bd641b031d 100644 --- a/intern/cycles/kernel/split/kernel_path_init.h +++ b/intern/cycles/kernel/split/kernel_path_init.h @@ -45,10 +45,6 @@ ccl_device void kernel_path_init(KernelGlobals *kg) { uint x, y, sample; get_work_pixel(tile, work_index, &x, &y, &sample); - /* Remap rng_state and buffer to current pixel. */ - ccl_global uint *rng_state = kernel_split_params.tile.rng_state; - rng_state += tile->offset + x + y*tile->stride; - /* Store buffer offset for writing to passes. */ uint buffer_offset = (tile->offset + x + y*tile->stride) * kernel_data.film.pass_stride; ccl_global float *buffer = tile->buffer + buffer_offset; @@ -57,7 +53,6 @@ ccl_device void kernel_path_init(KernelGlobals *kg) { /* Initialize random numbers and ray. */ uint rng_hash; kernel_path_trace_setup(kg, - rng_state, sample, x, y, &rng_hash, diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp index cf402c3f214..cf66e5385c9 100644 --- a/intern/cycles/render/buffers.cpp +++ b/intern/cycles/render/buffers.cpp @@ -108,7 +108,6 @@ RenderTile::RenderTile() stride = 0; buffer = 0; - rng_state = 0; buffers = NULL; } @@ -131,11 +130,6 @@ void RenderBuffers::device_free() device->mem_free(buffer); buffer.clear(); } - - if(rng_state.device_pointer) { - device->mem_free(rng_state); - rng_state.clear(); - } } void RenderBuffers::reset(Device *device, BufferParams& params_) @@ -149,11 +143,6 @@ void RenderBuffers::reset(Device *device, BufferParams& params_) buffer.resize(params.width*params.height*params.get_passes_size()); device->mem_alloc("render_buffer", buffer, MEM_READ_WRITE); device->mem_zero(buffer); - - /* allocate rng state */ - rng_state.resize(params.width, params.height); - - device->mem_alloc("rng_state", rng_state, MEM_READ_WRITE); } bool RenderBuffers::copy_from_device(Device *from_device) diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h index e56556c8abe..552dabe178a 100644 --- a/intern/cycles/render/buffers.h +++ b/intern/cycles/render/buffers.h @@ -74,8 +74,6 @@ public: /* float buffer */ device_vector<float> buffer; - /* random number generator state */ - device_vector<uint> rng_state; Device *device; @@ -149,7 +147,6 @@ public: int tile_index; device_ptr buffer; - device_ptr rng_state; RenderBuffers *buffers; diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp index f68efe38add..69cfb5c3e6e 100644 --- a/intern/cycles/render/session.cpp +++ b/intern/cycles/render/session.cpp @@ -384,7 +384,6 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile) tile_manager.state.buffer.get_offset_stride(rtile.offset, rtile.stride); rtile.buffer = buffers->buffer.device_pointer; - rtile.rng_state = buffers->rng_state.device_pointer; rtile.buffers = buffers; tile->buffers = buffers; @@ -442,7 +441,6 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile) tile->buffers->params.get_offset_stride(rtile.offset, rtile.stride); rtile.buffer = tile->buffers->buffer.device_pointer; - rtile.rng_state = tile->buffers->rng_state.device_pointer; rtile.buffers = tile->buffers; rtile.sample = 0; |